[clang] [llvm] [llvm][OpenMP][SPIRV] Fix assertion for GPU reductions (PR #194879)
Alex Duran via cfe-commits
cfe-commits at lists.llvm.org
Tue May 5 09:22:37 PDT 2026
https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/194879
>From 4560134032728638746f3bdd8fc33332b8e5019b Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Wed, 29 Apr 2026 07:59:29 -0700
Subject: [PATCH 1/6] [llvm][OpenMP][SPIRV] Fix assertion for GPU reductions
Currenty compiling a target reduction results in the following assert for spirv64-intel target:
Assertion `New->getType() == getType() && "replaceUses of value with new value of different type!"' failed.
This patch fixes it by adding an addrespace cast where necessary to make the types of the expressions match.
Assisted-by: claude-sonnet-4-5
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 16 +++++++++++++---
1 file changed, 13 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 5a4f12d91d540..6a1832a34cd9c 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4731,12 +4731,22 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
&LHSPtr, &RHSPtr, CurFunc));
// Fix the CallBack code genereated to use the correct Values for the LHS
- // and RHS
- LHSPtr->replaceUsesWithIf(RedValue, [ReductionFunc](const Use &U) {
+ // and RHS. Cast to match types before replacing (necessary to handle SPIRV address
+ // spaces).
+ Value *CastRedValue = RedValue;
+ if (LHSPtr->getType() != RedValue->getType())
+ CastRedValue = Builder.CreatePointerBitCastOrAddrSpaceCast(
+ RedValue, LHSPtr->getType());
+ Value *CastRHS = RHS;
+ if (RHSPtr->getType() != RHS->getType())
+ CastRHS =
+ Builder.CreatePointerBitCastOrAddrSpaceCast(RHS, RHSPtr->getType());
+
+ LHSPtr->replaceUsesWithIf(CastRedValue, [ReductionFunc](const Use &U) {
return cast<Instruction>(U.getUser())->getParent()->getParent() ==
ReductionFunc;
});
- RHSPtr->replaceUsesWithIf(RHS, [ReductionFunc](const Use &U) {
+ RHSPtr->replaceUsesWithIf(CastRHS, [ReductionFunc](const Use &U) {
return cast<Instruction>(U.getUser())->getParent()->getParent() ==
ReductionFunc;
});
>From 38826b1a77db66d96c0d4777b83ece1a9ebb85e1 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Wed, 29 Apr 2026 08:12:38 -0700
Subject: [PATCH 2/6] format
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 6a1832a34cd9c..30bac4097027c 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4731,8 +4731,8 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
&LHSPtr, &RHSPtr, CurFunc));
// Fix the CallBack code genereated to use the correct Values for the LHS
- // and RHS. Cast to match types before replacing (necessary to handle SPIRV address
- // spaces).
+ // and RHS. Cast to match types before replacing (necessary to handle
+ // SPIRV address spaces).
Value *CastRedValue = RedValue;
if (LHSPtr->getType() != RedValue->getType())
CastRedValue = Builder.CreatePointerBitCastOrAddrSpaceCast(
>From 843c327b639609ef5d4f0ea110c09a1a1cfe6025 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Tue, 5 May 2026 08:11:04 -0700
Subject: [PATCH 3/6] add test
---
.../spirv_target_teams_reduction_addrspace.c | 34 +++++++++++++++++++
1 file changed, 34 insertions(+)
create mode 100644 clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
diff --git a/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
new file mode 100644
index 0000000000000..8d85ed45401d1
--- /dev/null
+++ b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
@@ -0,0 +1,34 @@
+// Test that target teams reduction codegen handles address space casts correctly.
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+// Verify the kernel is generated
+// CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}}_main_{{.*}}
+
+// Verify __kmpc_alloc_shared is called for reduction variable
+// The return type should be ptr addrspace(4) (generic pointer)
+// CHECK: call spir_func align 8 addrspace(9) ptr addrspace(4) @__kmpc_alloc_shared(i64 4)
+
+// Verify the reduction runtime function is called
+// CHECK: call spir_func addrspace(9) i32 @__kmpc_nvptx_teams_reduce_nowait_v2(
+
+// Verify __kmpc_free_shared is called
+// CHECK: call spir_func addrspace(9) void @__kmpc_free_shared(ptr addrspace(4)
+
+// Verify the reduction function is generated
+// This is where the address space cast fix is critical
+// CHECK: define internal void @{{.*}}reduction{{.*}}func
+
+int main() {
+ int x = 0;
+
+ #pragma omp target teams num_teams(2) reduction(+ : x)
+ {
+ x += 2;
+ }
+
+ return x;
+}
\ No newline at end of file
>From e15d79756805c277060329b009bfc2e3f95764d8 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Tue, 5 May 2026 08:12:33 -0700
Subject: [PATCH 4/6] fix comment
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 30bac4097027c..30fe6a28e37b2 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4732,7 +4732,7 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
// Fix the CallBack code genereated to use the correct Values for the LHS
// and RHS. Cast to match types before replacing (necessary to handle
- // SPIRV address spaces).
+ // different address spaces).
Value *CastRedValue = RedValue;
if (LHSPtr->getType() != RedValue->getType())
CastRedValue = Builder.CreatePointerBitCastOrAddrSpaceCast(
>From 198667d65859b21f32ab16ffa6dbd6bf1aedf49a Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Tue, 5 May 2026 09:19:23 -0700
Subject: [PATCH 5/6] fix test comments
---
.../spirv_target_teams_reduction_addrspace.c | 15 +++++++--------
1 file changed, 7 insertions(+), 8 deletions(-)
diff --git a/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
index 8d85ed45401d1..bddd5548b9b8b 100644
--- a/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
+++ b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c
@@ -5,21 +5,20 @@
// expected-no-diagnostics
-// Verify the kernel is generated
+// Verify the kernel is generated.
// CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}}_main_{{.*}}
-// Verify __kmpc_alloc_shared is called for reduction variable
-// The return type should be ptr addrspace(4) (generic pointer)
+// Verify __kmpc_alloc_shared is called for reduction variable.
+// The return type should be ptr addrspace(4) (generic pointer).
// CHECK: call spir_func align 8 addrspace(9) ptr addrspace(4) @__kmpc_alloc_shared(i64 4)
-// Verify the reduction runtime function is called
+// Verify the reduction runtime function is called.
// CHECK: call spir_func addrspace(9) i32 @__kmpc_nvptx_teams_reduce_nowait_v2(
-// Verify __kmpc_free_shared is called
+// Verify __kmpc_free_shared is called.
// CHECK: call spir_func addrspace(9) void @__kmpc_free_shared(ptr addrspace(4)
-// Verify the reduction function is generated
-// This is where the address space cast fix is critical
+// Verify the reduction function is generated.
// CHECK: define internal void @{{.*}}reduction{{.*}}func
int main() {
@@ -31,4 +30,4 @@ int main() {
}
return x;
-}
\ No newline at end of file
+}
>From d1a5ce138e4d18c8a4763c1f8805058677e3f611 Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Tue, 5 May 2026 09:22:22 -0700
Subject: [PATCH 6/6] small refactor
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 30fe6a28e37b2..ce3bfaee9898d 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -4733,16 +4733,15 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU(
// Fix the CallBack code genereated to use the correct Values for the LHS
// and RHS. Cast to match types before replacing (necessary to handle
// different address spaces).
- Value *CastRedValue = RedValue;
if (LHSPtr->getType() != RedValue->getType())
- CastRedValue = Builder.CreatePointerBitCastOrAddrSpaceCast(
+ RedValue = Builder.CreatePointerBitCastOrAddrSpaceCast(
RedValue, LHSPtr->getType());
Value *CastRHS = RHS;
if (RHSPtr->getType() != RHS->getType())
CastRHS =
Builder.CreatePointerBitCastOrAddrSpaceCast(RHS, RHSPtr->getType());
- LHSPtr->replaceUsesWithIf(CastRedValue, [ReductionFunc](const Use &U) {
+ LHSPtr->replaceUsesWithIf(RedValue, [ReductionFunc](const Use &U) {
return cast<Instruction>(U.getUser())->getParent()->getParent() ==
ReductionFunc;
});
More information about the cfe-commits
mailing list