llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Discussed in https://github.com/llvm/llvm-project/pull/230813 --- Full diff: https://github.com/llvm/llvm-project/pull/231098.diff 4 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+1-8) - (modified) clang/lib/CIR/CodeGen/TargetInfo.h (+7) - (modified) clang/lib/CIR/CodeGen/Targets/SPIRV.cpp (+12) - (modified) clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip (+11) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index db505dc9824d0..0c1a2dfccd0dc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2330,14 +2330,7 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s, mlir::ptr::MemorySpaceAttrInterface CIRGenModule::getGlobalCIRAddressSpace(LangAS as) { - // The cir::LangAddressSpace fold loses the SPIRDefIsGenMap split between - // cuda_constant and opencl_constant, so use the target map like classic. - if (langOpts.CUDAIsDevice && getTriple().isSPIRV() && - (as == LangAS::cuda_device || as == LangAS::cuda_constant || - as == LangAS::cuda_shared)) - return cir::TargetAddressSpaceAttr::get( - &getMLIRContext(), astContext.getTargetAddressSpace(as)); - return cir::toCIRAddressSpaceAttr(getMLIRContext(), as); + return getTargetCIRGenInfo().getGlobalCIRAddressSpace(as); } // TODO(cir): this could be a common AST helper for both CIR and LLVM codegen. diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 91d585fe00a19..09f41fcd3cf9b 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -70,6 +70,13 @@ class TargetCIRGenInfo { cir::LangAddressSpace::Default); } + /// Get the CIR address space attribute for a global in AST address space + /// \p as. + virtual mlir::ptr::MemorySpaceAttrInterface + getGlobalCIRAddressSpace(clang::LangAS as) const { + return cir::toCIRAddressSpaceAttr(info->cgt.getMLIRContext(), as); + } + /// Get the CIR value of a null pointer of type \p ptrTy, where \p qt is the /// source pointer type. virtual mlir::Value getNullPointer(CIRGenModule &cgm, cir::PointerType ptrTy, diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 8e36b25301dc9..2231d96550cdc 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -40,6 +40,18 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { cir::LangAddressSpace::OffloadPrivate); } + mlir::ptr::MemorySpaceAttrInterface + getGlobalCIRAddressSpace(LangAS as) const override { + // The cir::LangAddressSpace fold loses the SPIRDefIsGenMap split between + // cuda_constant and opencl_constant, so use the target map like classic. + if (as == LangAS::cuda_device || as == LangAS::cuda_constant || + as == LangAS::cuda_shared) + return cir::TargetAddressSpaceAttr::get( + &getABIInfo().cgt.getMLIRContext(), + getABIInfo().cgt.getASTContext().getTargetAddressSpace(as)); + return TargetCIRGenInfo::getGlobalCIRAddressSpace(as); + } + void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, CIRGenModule &cgm) const override { auto func = mlir::dyn_cast<cir::FuncOp>(global); diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index 92d10dcbf1e86..8c3cbe27254af 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -7,12 +7,18 @@ #define __device__ __attribute__((device)) #define __constant__ __attribute__((constant)) +#define __shared__ __attribute__((shared)) // CIR-DAG: cir.global {{.*}} target_address_space(1) @tab // CIR-DAG: cir.global {{.*}} target_address_space(1) @_ZZ10static_fmtiE3fmt // LLVM-DAG: @tab = addrspace(1) externally_initialized constant [2 x i32] [i32 1, i32 2] +// CIR-DAG: cir.global {{.*}} target_address_space(1) @dev +// CIR-DAG: cir.global {{.*}} target_address_space(3) @_ZZ10get_sharedvE2sh // LLVM-DAG: @_ZZ10static_fmtiE3fmt = internal addrspace(1) constant [3 x i8] c"%d\00" +// LLVM-DAG: @dev = addrspace(1) externally_initialized global i32 3 +// LLVM-DAG: @_ZZ10get_sharedvE2sh = internal addrspace(3) global [4 x i32] __constant__ int tab[2] = {1, 2}; +__device__ int dev = 3; // CIR-LABEL: cir.func{{.*}} @_Z10static_fmti // CIR: %[[G:.*]] = cir.get_global @_ZZ10static_fmtiE3fmt : !cir.ptr<!cir.array<!s8i x 3>, target_address_space(1)> @@ -33,6 +39,11 @@ __device__ const char *static_fmt(int i) { // LLVM: addrspacecast (ptr addrspace(1) @tab to ptr addrspace(4)) __device__ const int *get_tab() { return tab; } +__device__ int get_shared() { + __shared__ int sh[4]; + return dev + sh[0]; +} + // 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; } `````````` </details> https://github.com/llvm/llvm-project/pull/231098 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
