[clang] 171da44 - [HIPSPV] Fix literals are mapped to Generic address space

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2022-02-05T17:26:52-05:00
New Revision: 171da443d598662371f4101f0ba92c0138011ff6

URL: https://github.com/llvm/llvm-project/commit/171da443d598662371f4101f0ba92c0138011ff6
DIFF: https://github.com/llvm/llvm-project/commit/171da443d598662371f4101f0ba92c0138011ff6.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 5ffb954642a6b..a8f0892f0e8f4 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";
+}