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.sar...@intel.com>
Date: Thu, 3 Apr 2025 09:08:44 -0700
Subject: [PATCH 1/4] [clang][OpenMP][SPIR-V] Fix addrspace of globals and
global constants
Signed-off-by: Sarnie, Nick <nick.sar...@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.sar...@intel.com>
Date: Mon, 7 Apr 2025 11:29:54 -0700
Subject: [PATCH 2/4] do it in the target
Signed-off-by: Sarnie, Nick <nick.sar...@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() {
>From 06bfbba877c26630b6c5b0ffef7f6623aa2e9ee8 Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sar...@intel.com>
Date: Mon, 7 Apr 2025 11:31:16 -0700
Subject: [PATCH 3/4] whitespace
Signed-off-by: Sarnie, Nick <nick.sar...@intel.com>
---
clang/test/CodeGenHLSL/GlobalDestructors.hlsl | 4 +---
1 file changed, 1 insertion(+), 3 deletions(-)
diff --git a/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
b/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
index fcb0c9b20e052..f83f26b31fd79 100644
--- a/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
+++ b/clang/test/CodeGenHLSL/GlobalDestructors.hlsl
@@ -88,9 +88,7 @@ void main(unsigned GI : SV_GroupIndex) {
// NOINLINE-SPIRV-NEXT: entry:
// NOINLINE-SPIRV-NEXT: %0 = call token
@llvm.experimental.convergence.entry()
// 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: 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
>From c4941ea7b9d17001f583faa65493311f2ad4f88a Mon Sep 17 00:00:00 2001
From: "Sarnie, Nick" <nick.sar...@intel.com>
Date: Mon, 7 Apr 2025 12:39:37 -0700
Subject: [PATCH 4/4] opencl_constant only for ocl
Signed-off-by: Sarnie, Nick <nick.sar...@intel.com>
---
clang/lib/Basic/Targets/SPIR.h | 11 ++++++++++-
clang/test/OpenMP/spirv_target_addrspace.c | 2 +-
2 files changed, 11 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 36187ff5b9b4e..b6694c9c52f37 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -376,9 +376,18 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public
BaseSPIRVTargetInfo {
}
std::optional<LangAS> getConstantAddressSpace() const override {
+ return ConstantAS;
+ }
+ void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
+ BaseSPIRVTargetInfo::adjust(Diags, Opts);
// opencl_constant will map to UniformConstant in SPIR-V
- return LangAS::opencl_constant;
+ if (Opts.OpenCL)
+ ConstantAS = LangAS::opencl_constant;
}
+
+private:
+ // opencl_global will map to CrossWorkgroup in SPIR-V
+ LangAS ConstantAS = LangAS::opencl_global;
};
class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
diff --git a/clang/test/OpenMP/spirv_target_addrspace.c
b/clang/test/OpenMP/spirv_target_addrspace.c
index 92ac029e8a3b7..a2a20d0a37e04 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 1f
+// CHECK: @.str = private unnamed_addr addrspace(1) constant [4 x i8]
c"foo\00", align 1
int global = 0;
#pragma omp end declare target
int main() {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
