https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/195911
>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/8] [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/8] 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/8] 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" <[email protected]> Date: Tue, 5 May 2026 08:12:33 -0700 Subject: [PATCH 4/8] 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" <[email protected]> Date: Tue, 5 May 2026 09:19:23 -0700 Subject: [PATCH 5/8] 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" <[email protected]> Date: Tue, 5 May 2026 09:22:22 -0700 Subject: [PATCH 6/8] 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; }); >From 3a0be79c01575125b94220e0d88a290952ffb372 Mon Sep 17 00:00:00 2001 From: "Duran, Alex" <[email protected]> Date: Tue, 5 May 2026 11:39:04 -0700 Subject: [PATCH 7/8] minor 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 ce3bfaee9898d..adb290ecfeeef 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -4736,16 +4736,15 @@ OpenMPIRBuilder::InsertPointOrErrorTy OpenMPIRBuilder::createReductionsGPU( if (LHSPtr->getType() != RedValue->getType()) RedValue = Builder.CreatePointerBitCastOrAddrSpaceCast( RedValue, LHSPtr->getType()); - Value *CastRHS = RHS; if (RHSPtr->getType() != RHS->getType()) - CastRHS = + RHS = Builder.CreatePointerBitCastOrAddrSpaceCast(RHS, RHSPtr->getType()); LHSPtr->replaceUsesWithIf(RedValue, [ReductionFunc](const Use &U) { return cast<Instruction>(U.getUser())->getParent()->getParent() == ReductionFunc; }); - RHSPtr->replaceUsesWithIf(CastRHS, [ReductionFunc](const Use &U) { + RHSPtr->replaceUsesWithIf(RHS, [ReductionFunc](const Use &U) { return cast<Instruction>(U.getUser())->getParent()->getParent() == ReductionFunc; }); >From 4f6e13a6097baaf5f730c0de26530c03422156e6 Mon Sep 17 00:00:00 2001 From: "Duran, Alex" <[email protected]> Date: Tue, 5 May 2026 12:18:00 -0700 Subject: [PATCH 8/8] [clang][OpenMP][SPIRV] Use the right calling convention for reduction helpers --- clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c | 6 ++++-- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 3 +++ 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c index bddd5548b9b8b..7217ef9400a6b 100644 --- a/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c +++ b/clang/test/OpenMP/spirv_target_teams_reduction_addrspace.c @@ -18,8 +18,10 @@ // 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 +// Verify the reduction helper functions are generated. +// CHECK: define internal spir_func void @{{.*}}reduction{{.*}}func +// CHECK: define internal spir_func void @{{.*}}shuffle_and_reduce_func +// CHECK: define internal spir_func void @{{.*}}inter_warp_copy_func int main() { int x = 0; diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index adb290ecfeeef..f179316a822d5 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -3426,6 +3426,7 @@ Expected<Function *> OpenMPIRBuilder::emitInterWarpCopyFunction( Function *WcFunc = Function::Create(FuncTy, GlobalVariable::InternalLinkage, "_omp_reduction_inter_warp_copy_func", &M); + WcFunc->setCallingConv(Config.getRuntimeCC()); WcFunc->setAttributes(FuncAttrs); WcFunc->addParamAttr(0, Attribute::NoUndef); WcFunc->addParamAttr(1, Attribute::NoUndef); @@ -3686,6 +3687,7 @@ Expected<Function *> OpenMPIRBuilder::emitShuffleAndReduceFunction( Function *SarFunc = Function::Create(FuncTy, GlobalVariable::InternalLinkage, "_omp_reduction_shuffle_and_reduce_func", &M); + SarFunc->setCallingConv(Config.getRuntimeCC()); SarFunc->setAttributes(FuncAttrs); SarFunc->addParamAttr(0, Attribute::NoUndef); SarFunc->addParamAttr(1, Attribute::NoUndef); @@ -4382,6 +4384,7 @@ Expected<Function *> OpenMPIRBuilder::createReductionFunction( std::string Name = getReductionFuncName(ReducerName); Function *ReductionFunc = Function::Create(FuncTy, GlobalVariable::InternalLinkage, Name, &M); + ReductionFunc->setCallingConv(Config.getRuntimeCC()); ReductionFunc->setAttributes(FuncAttrs); ReductionFunc->addParamAttr(0, Attribute::NoUndef); ReductionFunc->addParamAttr(1, Attribute::NoUndef); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
