maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Yaxun (Sam) Liu 2022-02-04 18:18:14 -05:00
parent 73f55fba76
commit 171da443d5
2 changed files with 16 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -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;

8
clang/test/CodeGenHIP/hipspv-addr-spaces.cpp
View File

@ -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";
}