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