[PATCH] D118876: [HIPSPV] Fix literals are mapped to Generic address space

[Henry Linjamäki via Phabricator via cfe-commits]

linjamaki created this revision.
Herald added a subscriber: yaxunl.
linjamaki edited the summary of this revision.
Herald added a subscriber: Anastasia.
linjamaki published this revision for review.
linjamaki added reviewers: Anastasia, yaxunl.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.

This issue is an oversight in D108621 <https://reviews.llvm.org/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`.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D118876

Files:
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/CodeGenHIP/hipspv-addr-spaces.cpp


Index: clang/test/CodeGenHIP/hipspv-addr-spaces.cpp
===================================================================
--- clang/test/CodeGenHIP/hipspv-addr-spaces.cpp
+++ 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";
+}
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ 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;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D118876.405533.patch
Type: text/x-patch
Size: 1781 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220203/10daf42b/attachment-0001.bin>