Author: Yaxun (Sam) Liu Date: 2022-02-07T10:00:54-05:00 New Revision: 02d5b112138e7e9f30dec685afb380c1b9593a84
URL: https://github.com/llvm/llvm-project/commit/02d5b112138e7e9f30dec685afb380c1b9593a84 DIFF: https://github.com/llvm/llvm-project/commit/02d5b112138e7e9f30dec685afb380c1b9593a84.diff LOG: [HIPSPV] Fix literals are mapped to Generic address space This issue is an oversight in D108621. Literals in HIP are emitted as global constant variables with default address space which maps to Generic address space for HIPSPV. In SPIR-V such variables translate to OpVariable instructions with Generic storage class which are not legal. Fix by mapping literals to CrossWorkGroup address space. The literals are not mapped to UniformConstant because the “flat” pointers in HIP may reference them and “flat” pointers are modeled as Generic pointers in SPIR-V. In SPIR-V/OpenCL UniformConstant pointers may not be casted to Generic. Patch by: Henry Linjamäki Reviewed by: Yaxun Liu Differential Revision: https://reviews.llvm.org/D118876 Added: Modified: clang/lib/CodeGen/CodeGenModule.cpp clang/test/CodeGenHIP/hipspv-addr-spaces.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 2346176a15628..29806b65e984e 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4381,6 +4381,14 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const { return LangAS::opencl_constant; if (LangOpts.SYCLIsDevice) return LangAS::sycl_global; + if (LangOpts.HIP && LangOpts.CUDAIsDevice && getTriple().isSPIRV()) + // For HIPSPV map literals to cuda_device (maps to CrossWorkGroup in SPIR-V) + // instead of default AS (maps to Generic in SPIR-V). Otherwise, we end up + // with OpVariable instructions with Generic storage class which is not + // allowed (SPIR-V V1.6 s3.42.8). Also, mapping literals to SPIR-V + // UniformConstant storage class is not viable as pointers to it may not be + // casted to Generic pointers which are used to model HIP's "flat" pointers. + return LangAS::cuda_device; if (auto AS = getTarget().getConstantAddressSpace()) return AS.getValue(); return LangAS::Default; diff --git a/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp b/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp index 8f56f2104ecbd..bde360eec8cd9 100644 --- a/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp +++ b/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp @@ -22,6 +22,9 @@ __device__ struct foo_t { int* pi; } foo; +// Check literals are placed in address space 1 (CrossWorkGroup/__global). +// CHECK: @.str ={{.*}} unnamed_addr addrspace(1) constant + // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z3barPi(i32 addrspace(4)* __device__ int* bar(int *x) { return x; @@ -44,3 +47,8 @@ __device__ int* baz_s() { // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @s to i32 addrspace(4)* return &s; } + +// CHECK: define{{.*}} spir_func noundef i8 addrspace(4)* @_Z3quzv() +__device__ const char* quz() { + return "abc"; +} _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits