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" <[email protected]> Date: Wed, 29 Apr 2026 07:59:29 -0700 Subject: [PATCH 1/3] [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" <[email protected]> Date: Wed, 29 Apr 2026 08:12:38 -0700 Subject: [PATCH 2/3] 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" <[email protected]> Date: Tue, 5 May 2026 08:11:04 -0700 Subject: [PATCH 3/3] 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
