Author: Alex Duran Date: 2026-05-08T12:14:46+02:00 New Revision: ca7fe087557033802371ea74210a8ada737d3719
URL: https://github.com/llvm/llvm-project/commit/ca7fe087557033802371ea74210a8ada737d3719 DIFF: https://github.com/llvm/llvm-project/commit/ca7fe087557033802371ea74210a8ada737d3719.diff LOG: [llvm][OpenMP][SPIRV] Fix assertion for GPU reductions (#194879) 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 Added: clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c Modified: llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp Removed: ################################################################################ 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..bddd5548b9b8b --- /dev/null +++ b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c @@ -0,0 +1,33 @@ +// 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. +// CHECK: define internal void @{{.*}}reduction{{.*}}func + +int main() { + int x = 0; + + #pragma omp target teams num_teams(2) reduction(+ : x) + { + x += 2; + } + + return x; +} diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index f17602e8e786c..e3d5bf0663490 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -4735,7 +4735,15 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU( &LHSPtr, &RHSPtr, CurFunc)); // Fix the CallBack code genereated to use the correct Values for the LHS - // and RHS + // and RHS. Cast to match types before replacing (necessary to handle + // diff erent address spaces). + if (LHSPtr->getType() != RedValue->getType()) + RedValue = Builder.CreatePointerBitCastOrAddrSpaceCast( + RedValue, LHSPtr->getType()); + if (RHSPtr->getType() != RHS->getType()) + RHS = + Builder.CreatePointerBitCastOrAddrSpaceCast(RHS, RHSPtr->getType()); + LHSPtr->replaceUsesWithIf(RedValue, [ReductionFunc](const Use &U) { return cast<Instruction>(U.getUser())->getParent()->getParent() == ReductionFunc; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
