This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG171da443d598: [HIPSPV] Fix literals are mapped to Generic
address space (authored by yaxunl).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D118876/new/
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
