diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4381,6 +4381,14 @@ 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 --- a/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp +++ b/clang/test/CodeGenHIP/hipspv-addr-spaces.cpp @@ -22,6 +22,9 @@ 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 @@ // 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"; +}