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;
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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
