llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> CIR emitted function addresses in address space 0, but some targets put functions in other, like address space 4 on spirv64-amd-amdhsa --- Full diff: https://github.com/llvm/llvm-project/pull/233755.diff 7 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenExpr.cpp (+4-2) - (modified) clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp (+9-1) - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+3-1) - (modified) clang/lib/CIR/Dialect/IR/CIRDialect.cpp (+10-1) - (modified) clang/test/CIR/CodeGen/address-space.c (+7) - (modified) clang/test/CIR/CodeGen/no-proto-then-def.c (+7) - (modified) clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip (+9) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index dfd6fa725e77fe..84d30de7803c94 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -1009,7 +1009,9 @@ static LValue emitFunctionDeclLValue(CIRGenFunction &cgf, const Expr *e, assert(!cir::MissingFeatures::sanitizers()); mlir::Type fnTy = funcOp.getFunctionType(); - mlir::Type ptrTy = cir::PointerType::get(fnTy); + mlir::ptr::MemorySpaceAttrInterface fnAddrSpace = + cgf.cgm.getTypes().getPointerAddressSpace(fd->getType()); + mlir::Type ptrTy = cir::PointerType::get(fnTy, fnAddrSpace); mlir::Value addr; // On the HIP host, a reference to a __global__ kernel must resolve to the @@ -1033,7 +1035,7 @@ static LValue emitFunctionDeclLValue(CIRGenFunction &cgf, const Expr *e, if (funcOp.getFunctionType() != cgf.convertType(fd->getType())) { fnTy = cgf.convertType(fd->getType()); - ptrTy = cir::PointerType::get(fnTy); + ptrTy = cir::PointerType::get(fnTy, fnAddrSpace); addr = cir::CastOp::create(cgf.getBuilder(), addr.getLoc(), ptrTy, cir::CastKind::bitcast, addr); diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index 2082963025238c..633fa8e1933077 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -2653,7 +2653,15 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *ce) { mlir::Value src = Visit(const_cast<Expr *>(subExpr)); mlir::Type dstTy = cgf.convertType(destTy); - assert(!cir::MissingFeatures::addressSpace()); + // Function pointers use the program address space and object pointers the + // default one, so a bitcast between them may change the address space. + auto srcPtrTy = mlir::dyn_cast<cir::PointerType>(src.getType()); + auto dstPtrTy = mlir::dyn_cast<cir::PointerType>(dstTy); + if (srcPtrTy && dstPtrTy && + srcPtrTy.getAddrSpace() != dstPtrTy.getAddrSpace()) + src = cgf.performAddrSpaceCast( + src, cir::PointerType::get(srcPtrTy.getPointee(), + dstPtrTy.getAddrSpace())); if (cgf.sanOpts.has(SanitizerKind::CFIUnrelatedCast)) cgf.getCIRGenModule().errorNYI(subExpr->getSourceRange(), diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 614a731158e455..0b964ed2b7516c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2139,7 +2139,9 @@ void CIRGenModule::replaceUsesOfNonProtoTypeWithRealFunction( // canonicalizer. mlir::Value res = getGlobalOp.getAddr(); const mlir::Type oldResTy = res.getType(); - const auto newPtrTy = cir::PointerType::get(newFn.getFunctionType()); + const auto newPtrTy = cir::PointerType::get( + newFn.getFunctionType(), + mlir::cast<cir::PointerType>(oldResTy).getAddrSpace()); if (oldResTy != newPtrTy) { res.setType(newPtrTy); builder.setInsertionPointAfter(getGlobalOp.getOperation()); diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 4dfbfbb195ca46..c09a60624d505c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -13,6 +13,7 @@ #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "clang/CIR/Dialect/IR/CIRAttrs.h" +#include "clang/CIR/Dialect/IR/CIRDataLayout.h" #include "clang/CIR/Dialect/IR/CIROpsEnums.h" #include "clang/CIR/Dialect/IR/CIRTypes.h" @@ -2759,7 +2760,15 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { << resultType.getPointee() << "' does not match type " << symTy << " of the global @" << getName(); - if (symAddrSpaceAttr != resultType.getAddrSpace()) { + // A function address may also be taken in the program address space. + auto targetAS = + dyn_cast_if_present<TargetAddressSpaceAttr>(resultType.getAddrSpace()); + bool isProgramAS = + isa<FuncOp>(op) && targetAS && + targetAS.getValue() == + CIRDataLayout(getOperation()->getParentOfType<ModuleOp>()) + .getProgramAddressSpace(); + if (symAddrSpaceAttr != resultType.getAddrSpace() && !isProgramAS) { return emitOpError() << "result type address space does not match the address " "space of the global @" diff --git a/clang/test/CIR/CodeGen/address-space.c b/clang/test/CIR/CodeGen/address-space.c index e806e08bca19ff..13d52e5201da4c 100644 --- a/clang/test/CIR/CodeGen/address-space.c +++ b/clang/test/CIR/CodeGen/address-space.c @@ -4,6 +4,8 @@ // RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll // RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG +// RUN: %clang_cc1 -triple avr -fclangir -emit-llvm %s -o %t-avr.ll +// RUN: FileCheck --input-file=%t-avr.ll %s -check-prefix=AVR // Test global variable with address space // CIR: cir.global external @gvar = #cir.ptr<null> : !cir.ptr<!s32i, target_address_space(1)> @@ -58,3 +60,8 @@ struct S { int a[4]; }; void copy_to_as(struct S __attribute__((address_space(1))) *p, struct S *q) { *p = *q; } + +// A function pointer in the program address space cast to an object pointer. +// AVR-LABEL: define{{.*}} ptr @fn_to_void_ptr +// AVR: store ptr addrspacecast (ptr addrspace(1) @copy_to_as to ptr), ptr +void *fn_to_void_ptr(void) { return (void *)copy_to_as; } diff --git a/clang/test/CIR/CodeGen/no-proto-then-def.c b/clang/test/CIR/CodeGen/no-proto-then-def.c index c142c4f0a16f5a..4eb3c8f753569e 100644 --- a/clang/test/CIR/CodeGen/no-proto-then-def.c +++ b/clang/test/CIR/CodeGen/no-proto-then-def.c @@ -4,6 +4,8 @@ // RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -Wno-deprecated-non-prototype %s -o %t.ll // RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fclangir -emit-cir -Wno-deprecated-non-prototype %s -o %t-spirv.cir +// RUN: FileCheck --input-file=%t-spirv.cir %s -check-prefix=SPIRV extern int default_proc(); @@ -22,6 +24,11 @@ int default_proc(int a) { return a; } // CIR: cir.get_global @default_proc : !cir.ptr<!cir.func<(!s32i) -> !s32i>> // CIR: cir.cast bitcast %{{.*}} : !cir.ptr<!cir.func<(!s32i) -> !s32i>> -> !cir.ptr<!cir.func<(...) -> !s32i>> +// The retyped address keeps the program address space. +// SPIRV: cir.func{{.*}}@test_proc_ptr +// SPIRV: cir.get_global @default_proc : !cir.ptr<!cir.func<(!s32i) -> !s32i>, target_address_space(4)> +// SPIRV: cir.cast bitcast %{{.*}} : !cir.ptr<!cir.func<(!s32i) -> !s32i>, target_address_space(4)> -> !cir.ptr<!cir.func<(...) -> !s32i>, target_address_space(4)> + // LLVM: define{{.*}} @test_proc_ptr // LLVM: store ptr @default_proc, ptr // LLVM: define{{.*}} @default_proc diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index 1ddc16332f65ca..a93895fbf1a9e1 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -60,3 +60,12 @@ __device__ int get_shared() { // An explicit opencl_constant qualifier keeps UniformConstant. // LLVM-LABEL: define{{.*}} i32 @_Z6get_ocPU3AS2i(ptr addrspace(2) __device__ int get_oc(__attribute__((opencl_constant)) int *p) { return *p; } + +// Function addresses are in the program address space 4. +// CIR-LABEL: cir.func{{.*}} @_Z7take_fnv +// CIR: cir.get_global @_Z6fn_refv : !cir.ptr<!cir.func<()>, target_address_space(4)> +// LLVM-LABEL: define{{.*}} void @_Z7take_fnv +// LLVM: call{{.*}} @_Z6use_fnPFvvE(ptr addrspace(4) noundef @_Z6fn_refv) +__device__ void fn_ref(); +__device__ void use_fn(void (*)()); +__device__ void take_fn() { use_fn(fn_ref); } `````````` </details> https://github.com/llvm/llvm-project/pull/233755 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
