[clang] [llvm] [clang][OpenMP][SPIR-V] Fix addrspace of globals and global constants (PR #134399)
Nick Sarnie via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 7 11:30:20 PDT 2025
https://github.com/sarnex updated https://github.com/llvm/llvm-project/pull/134399
>From 3812b132c83e4a2e7ae9bd0b5ecefe7232f86af1 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Thu, 3 Apr 2025 09:08:44 -0700
Subject: [PATCH 1/2] [clang][OpenMP][SPIR-V] Fix addrspace of globals and
global constants
Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
clang/lib/CodeGen/CodeGenModule.cpp | 9 +++++++++
clang/test/OpenMP/spirv_target_addrspace.c | 20 ++++++++++++++++++++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 6 ++++++
3 files changed, 35 insertions(+)
create mode 100644 clang/test/OpenMP/spirv_target_addrspace.c
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 8f9cf965af2b9..cc6d726445cbb 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5384,6 +5384,11 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
LangAS AS;
if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
return AS;
+ if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+ // SPIR-V globals should map to CrossWorkGroup instead of default
+ // AS, as generic/no address space is invalid. This is similar
+ // to what is done for HIPSPV.
+ return LangAS::opencl_global;
}
return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
}
@@ -5402,6 +5407,10 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
// 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 (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
+ // OpenMP SPIR-V global constants should map to UniformConstant, different
+ // from the HIPSPV case above.
+ return LangAS::opencl_constant;
if (auto AS = getTarget().getConstantAddressSpace())
return *AS;
return LangAS::Default;
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c b/clang/test/OpenMP/spirv_target_addrspace.c
new file mode 100644
index 0000000000000..8430a30efe0c8
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_addrspace.c
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=spirv64 -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -O0 -fopenmp -fopenmp-targets=spirv64 -fopenmp-is-target-device -triple spirv64 -fopenmp-host-ir-file-path %t-host.bc -emit-llvm %s -o - | FileCheck %s
+
+extern int printf(char[]);
+
+#pragma omp declare target
+// CHECK: @global = addrspace(1) global i32 0, align 4
+// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1
+int global = 0;
+#pragma omp end declare target
+int main() {
+ // CHECK: = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @__omp_offloading_{{.*}}_kernel_environment to ptr), ptr %{{.*}})
+#pragma omp target
+ {
+ for(int i = 0; i < 5; i++)
+ global++;
+ printf("foo");
+ }
+ return global;
+}
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 68b1fa42934ad..998702c1af3cd 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -6295,6 +6295,12 @@ OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetInit(
: ConstantExpr::getAddrSpaceCast(KernelEnvironmentGV,
KernelEnvironmentPtr);
Value *KernelLaunchEnvironment = DebugKernelWrapper->getArg(0);
+ Type *KernelLaunchEnvParamTy = Fn->getFunctionType()->getParamType(1);
+ KernelLaunchEnvironment =
+ KernelLaunchEnvironment->getType() == KernelLaunchEnvParamTy
+ ? KernelLaunchEnvironment
+ : Builder.CreateAddrSpaceCast(KernelLaunchEnvironment,
+ KernelLaunchEnvParamTy);
CallInst *ThreadKind =
Builder.CreateCall(Fn, {KernelEnvironment, KernelLaunchEnvironment});
>From 93ff204e9417f4a09fa124add1a517452112cbba Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sarnie at intel.com>
Date: Mon, 7 Apr 2025 11:29:54 -0700
Subject: [PATCH 2/2] do it in the target
Signed-off-by: Sarnie, Nick <nick.sarnie at intel.com>
---
clang/lib/Basic/Targets/SPIR.h | 9 +++++++--
clang/lib/CodeGen/CodeGenModule.cpp | 9 ---------
clang/test/CodeGenHLSL/GlobalDestructors.hlsl | 6 ++++--
.../builtins/StructuredBuffers-constructors.hlsl | 10 ++++++----
clang/test/OpenMP/spirv_target_addrspace.c | 2 +-
5 files changed, 18 insertions(+), 18 deletions(-)
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 78505d66d6f2f..36187ff5b9b4e 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -37,8 +37,8 @@ static const unsigned SPIRDefIsPrivMap[] = {
0, // cuda_device
0, // cuda_constant
0, // cuda_shared
- // SYCL address space values for this map are dummy
- 0, // sycl_global
+ // Most SYCL address space values for this map are dummy
+ 1, // sycl_global
0, // sycl_global_device
0, // sycl_global_host
0, // sycl_local
@@ -374,6 +374,11 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public BaseSPIRVTargetInfo {
const llvm::omp::GV &getGridValue() const override {
return llvm::omp::SPIRVGridValues;
}
+
+ std::optional<LangAS> getConstantAddressSpace() const override {
+ // opencl_constant will map to UniformConstant in SPIR-V
+ return LangAS::opencl_constant;
+ }
};
class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index cc6d726445cbb..8f9cf965af2b9 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5384,11 +5384,6 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
LangAS AS;
if (OpenMPRuntime->hasAllocateAttributeForGlobalVar(D, AS))
return AS;
- if (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
- // SPIR-V globals should map to CrossWorkGroup instead of default
- // AS, as generic/no address space is invalid. This is similar
- // to what is done for HIPSPV.
- return LangAS::opencl_global;
}
return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
}
@@ -5407,10 +5402,6 @@ LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
// 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 (LangOpts.OpenMPIsTargetDevice && getTriple().isSPIRV())
- // OpenMP SPIR-V global constants should map to UniformConstant, different
- // from the HIPSPV case above.
- return LangAS::opencl_constant;
if (auto AS = getTarget().getConstantAddressSpace())
return *AS;
return LangAS::Default;
diff --git a/clang/test/CodeGenHLSL/GlobalDestructors.hlsl b/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
index 9f90971bafd05..fcb0c9b20e052 100644
--- a/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
+++ b/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
@@ -87,8 +87,10 @@ void main(unsigned GI : SV_GroupIndex) {
// NOINLINE-SPIRV: define internal spir_func void @_GLOBAL__D_a() [[IntAttr:\#[0-9]+]]
// NOINLINE-SPIRV-NEXT: entry:
// NOINLINE-SPIRV-NEXT: %0 = call token @llvm.experimental.convergence.entry()
-// NOINLINE-SPIRV-NEXT: call spir_func void @_ZN4TailD1Ev(ptr @_ZZ3WagvE1T) [ "convergencectrl"(token %0) ]
-// NOINLINE-SPIRV-NEXT: call spir_func void @_ZN6PupperD1Ev(ptr @GlobalPup) [ "convergencectrl"(token %0) ]
+// NOINLINE-SPIRV-NEXT: call spir_func void @_ZN4TailD1Ev(ptr addrspacecast (ptr addrspace(1) @_ZZ3WagvE1T to ptr)) [ "convergencectrl"(token %0) ]
+
+// NOINLINE-SPIRV-NEXT: call spir_func void @_ZN6PupperD1Ev(ptr addrspacecast (ptr addrspace(1) @GlobalPup to ptr)) [ "convergencectrl"(token %0) ]
+
// NOINLINE-SPIRV-NEXT: ret void
// NOINLINE: attributes [[IntAttr]] = {{.*}} alwaysinline
diff --git a/clang/test/CodeGenHLSL/builtins/StructuredBuffers-constructors.hlsl b/clang/test/CodeGenHLSL/builtins/StructuredBuffers-constructors.hlsl
index 8a1429fd1a6fc..62993a332d205 100644
--- a/clang/test/CodeGenHLSL/builtins/StructuredBuffers-constructors.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/StructuredBuffers-constructors.hlsl
@@ -22,8 +22,10 @@ RasterizerOrderedStructuredBuffer<float> Buf5 : register(u1, space2);
// CHECK-SPIRV: %"class.hlsl::RWStructuredBuffer" = type { target("spirv.VulkanBuffer", [0 x float], 12, 1) }
-// CHECK: @_ZL3Buf = internal global %"class.hlsl::StructuredBuffer" poison
-// CHECK: @_ZL4Buf2 = internal global %"class.hlsl::RWStructuredBuffer" poison
+// CHECK-SPIRV: @_ZL3Buf = internal addrspace(1) global %"class.hlsl::StructuredBuffer" poison
+// CHECK-SPIRV: @_ZL4Buf2 = internal addrspace(1) global %"class.hlsl::RWStructuredBuffer" poison
+// CHECK-DXIL: @_ZL3Buf = internal{{.*}}global %"class.hlsl::StructuredBuffer" poison
+// CHECK-DXIL: @_ZL4Buf2 = internal{{.*}}global %"class.hlsl::RWStructuredBuffer" poison
// CHECK-DXIL: @_ZL4Buf3 = internal global %"class.hlsl::AppendStructuredBuffer" poison, align 4
// CHECK-DXIL: @_ZL4Buf4 = internal global %"class.hlsl::ConsumeStructuredBuffer" poison, align 4
// CHECK-DXIL: @_ZL4Buf5 = internal global %"class.hlsl::RasterizerOrderedStructuredBuffer" poison, align 4
@@ -32,13 +34,13 @@ RasterizerOrderedStructuredBuffer<float> Buf5 : register(u1, space2);
// CHECK-DXIL: [[H:%.*]] = call target("dx.RawBuffer", float, 0, 0) @llvm.dx.resource.handlefrombinding.tdx.RawBuffer_f32_0_0t(i32 0, i32 10, i32 1, i32 0, i1 false)
// CHECK-DXIL: store target("dx.RawBuffer", float, 0, 0) [[H]], ptr @_ZL3Buf, align 4
// CHECK-SPIRV: [[H:%.*]] = call target("spirv.VulkanBuffer", [0 x float], 12, 0) @llvm.spv.resource.handlefrombinding.tspirv.VulkanBuffer_a0f32_12_0t(i32 0, i32 10, i32 1, i32 0, i1 false)
-// CHECK-SPIRV: store target("spirv.VulkanBuffer", [0 x float], 12, 0) [[H]], ptr @_ZL3Buf, align 8
+// CHECK-SPIRV: store target("spirv.VulkanBuffer", [0 x float], 12, 0) [[H]], ptr addrspace(1) @_ZL3Buf, align 8
// CHECK: define internal void @_init_resource__ZL4Buf2()
// CHECK-DXIL: [[H:%.*]] = call target("dx.RawBuffer", float, 1, 0) @llvm.dx.resource.handlefrombinding.tdx.RawBuffer_f32_1_0t(i32 1, i32 5, i32 1, i32 0, i1 false)
// CHECK-DXIL: store target("dx.RawBuffer", float, 1, 0) [[H]], ptr @_ZL4Buf2, align 4
// CHECK-SPIRV: [[H:%.*]] = call target("spirv.VulkanBuffer", [0 x float], 12, 1) @llvm.spv.resource.handlefrombinding.tspirv.VulkanBuffer_a0f32_12_1t(i32 1, i32 5, i32 1, i32 0, i1 false)
-// CHECK-SPIRV: store target("spirv.VulkanBuffer", [0 x float], 12, 1) [[H]], ptr @_ZL4Buf2, align 8
+// CHECK-SPIRV: store target("spirv.VulkanBuffer", [0 x float], 12, 1) [[H]], ptr addrspace(1) @_ZL4Buf2, align 8
// CHECK-DXIL: define internal void @_init_resource__ZL4Buf3()
// CHECK-DXIL: [[H:%.*]] = call target("dx.RawBuffer", float, 1, 0) @llvm.dx.resource.handlefrombinding.tdx.RawBuffer_f32_1_0t(i32 0, i32 3, i32 1, i32 0, i1 false)
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c b/clang/test/OpenMP/spirv_target_addrspace.c
index 8430a30efe0c8..92ac029e8a3b7 100644
--- a/clang/test/OpenMP/spirv_target_addrspace.c
+++ b/clang/test/OpenMP/spirv_target_addrspace.c
@@ -5,7 +5,7 @@ extern int printf(char[]);
#pragma omp declare target
// CHECK: @global = addrspace(1) global i32 0, align 4
-// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1
+// CHECK: @.str = private unnamed_addr addrspace(2) constant [4 x i8] c"foo\00", align 1f
int global = 0;
#pragma omp end declare target
int main() {
More information about the cfe-commits
mailing list