[HIPSPV] Fix literals are mapped to Generic address space

[linjamaki]

Summary

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.

[yaxunl]

LGTM. Thanks.

[linjamaki]

Thanks for the review, @yaxunl. Could you push this to the LLVM? And to the LLVM 14 release branch too, if possible?

[yaxunl]

In D118876#3295958, @linjamaki wrote:

Thanks for the review, @yaxunl. Could you push this to the LLVM? And to the LLVM 14 release branch too, if possible?

Sure. @tstellar What is the current procedure for cherry-picking a fix to a release branch? Thanks.

[tstellar]

@yaxunl This looks good to me, you can go ahead and cherry-pick it directly to release/14.x.

[yaxunl]

It has been cherry-picked to 14.x by 02d5b112

clang/lib/CodeGen/CodeGenModule.cpp

 }
 
 LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
   // OpenCL v1.2 s6.5.3: a string literal is in the constant address space.
   if (LangOpts.OpenCL)
     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;
 }
 
 // In address space agnostic languages, string literals are in default address
 // space in AST. However, certain targets (e.g. amdgcn) request them to be
 // emitted in constant address space in LLVM IR. To be consistent with other

clang/test/CodeGenHIP/hipspv-addr-spaces.cpp

 __shared__ int s;
 
 // CHECK: @foo ={{.*}} addrspace(1) externally_initialized global %struct.foo_t
 __device__ struct foo_t {
   int i;
   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;
 }
 
 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_dv()
 __device__ int* baz_d() {
   // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @d to i32 addrspace(4)*
   return &d;
 }
 
 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_cv()
 __device__ int* baz_c() {
   // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @c to i32 addrspace(4)*
   return &c;
 }
 
 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_sv()
 __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";
+}