[clang] ca7fe08 - [llvm][OpenMP][SPIRV] Fix assertion for GPU reductions (#194879)

[via cfe-commits]

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;