Author: Robert Imschweiler Date: 2026-08-10T20:58:54+02:00 New Revision: e03cfd917b2febacc907a64fe515f2f8eb69ead9
URL: https://github.com/llvm/llvm-project/commit/e03cfd917b2febacc907a64fe515f2f8eb69ead9 DIFF: https://github.com/llvm/llvm-project/commit/e03cfd917b2febacc907a64fe515f2f8eb69ead9.diff LOG: [OpenMP][offload] Fix target reduction for narrow non-integer types (#215294) Cast the result of the shuffle function back to the width that corresponds to the integer that is used to hold the bits of the actual type. Prevents corresponding reductions from returning NaN or otherwise incorrect results at higher optimization levels. Claude assisted with this patch. Added: clang/test/OpenMP/target_parallel_reduction_half_codegen.c offload/test/offloading/xteam_fp16_reduction.c Modified: llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp Removed: ################################################################################ diff --git a/clang/test/OpenMP/target_parallel_reduction_half_codegen.c b/clang/test/OpenMP/target_parallel_reduction_half_codegen.c new file mode 100644 index 0000000000000..bb14038be42fb --- /dev/null +++ b/clang/test/OpenMP/target_parallel_reduction_half_codegen.c @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu \ +// RUN: -fopenmp-targets=amdgpu-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc +// RUN: %clang_cc1 -verify -fopenmp -x c -triple amdgpu-amd-amdhsa \ +// RUN: -fopenmp-targets=amdgpu-amd-amdhsa -emit-llvm %s \ +// RUN: -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - \ +// RUN: | FileCheck %s + +// expected-no-diagnostics + +// Check that the value returned by the warp shuffle is narrowed back to the +// width of the reduction element before it is stored into the reduction slot. +// The shuffle runtime functions always return a 32- or 64-bit value, so storing +// it unnarrowed writes past the end of a slot for a narrower element such as +// 'half'. + +_Float16 half_reduction(_Float16 *a, int n) { + _Float16 s = 0; +#pragma omp target parallel for map(tofrom : s) reduction(+ : s) + for (int i = 0; i < n; ++i) + s += a[i]; + return s; +} + +// CHECK-LABEL: define internal void @_omp_reduction_shuffle_and_reduce_func( +// CHECK: %[[ELEM:.+]] = alloca half, align 2 +// CHECK: %[[SHUFFLE:.+]] = call i32 @__kmpc_shuffle_int32( +// CHECK-NEXT: %[[NARROWED:.+]] = trunc i32 %[[SHUFFLE]] to i16 +// AMDGPU allocas live in addrspace(5) and are accessed through a cast to the +// generic address space, so the store goes through '<name>.ascast' there. +// CHECK-NEXT: store i16 %[[NARROWED]], ptr %[[ELEM]]{{(\.ascast)?}}, align 2 diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 38a56e975a807..3560cfef096fe 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -1838,7 +1838,8 @@ class OpenMPIRBuilder { Value *castValueToType(InsertPointTy AllocaIP, Value *From, Type *ToType); /// This function creates calls to one of two shuffle functions to copy - /// variables between lanes in a warp. + /// variables between lanes in a warp. The returned value has \p ElementType, + /// even though the shuffle runtime functions operate on 32- or 64-bit values. Value *createRuntimeShuffleFunction(InsertPointTy AllocaIP, Value *Element, Type *ElementType, Value *Offset); diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 63eab978b1db0..16575a5c3eb12 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -3398,7 +3398,10 @@ Value *OpenMPIRBuilder::createRuntimeShuffleFunction(InsertPointTy AllocaIP, Builder.CreateIntCast(WarpSize, Builder.getInt16Ty(), /*isSigned=*/true); Value *ShuffleCall = createRuntimeFunctionCall(ShuffleFunc, {ElemCast, Offset, WarpSizeCast}); - return castValueToType(AllocaIP, ShuffleCall, CastTy); + // The shuffle runtime functions return a 32- or 64-bit value. Cast it back + // down to the requested element type, otherwise storing the result would + // write past the end of an element narrower than the shuffle width. + return castValueToType(AllocaIP, ShuffleCall, ElementType); } void OpenMPIRBuilder::shuffleAndStore(InsertPointTy AllocaIP, Value *SrcAddr, @@ -3472,11 +3475,10 @@ void OpenMPIRBuilder::shuffleAndStore(InsertPointTy AllocaIP, Value *SrcAddr, emitBranch(PreCondBB); emitBlock(ExitBB, CurFunc); } else { + // The shuffled value comes back as the chunk's integer type, so the + // store covers exactly this chunk regardless of what ElemType is. Value *Res = createRuntimeShuffleFunction( AllocaIP, Builder.CreateLoad(IntType, Ptr), IntType, Offset); - if (ElemType->isIntegerTy() && ElemType->getScalarSizeInBits() < - Res->getType()->getScalarSizeInBits()) - Res = Builder.CreateTrunc(Res, ElemType); Builder.CreateStore(Res, ElemPtr); Ptr = Builder.CreateGEP(IntType, Ptr, {ConstantInt::get(IndexTy, 1)}); ElemPtr = diff --git a/offload/test/offloading/xteam_fp16_reduction.c b/offload/test/offloading/xteam_fp16_reduction.c new file mode 100644 index 0000000000000..c60fa0cbad1ec --- /dev/null +++ b/offload/test/offloading/xteam_fp16_reduction.c @@ -0,0 +1,37 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// RUN: %libomptarget-compileopt-run-and-check-generic +// UNSUPPORTED: intelgpu + +#include <stdio.h> + +int main(void) { + _Float16 sum = 0; +#pragma omp target teams distribute parallel for map(tofrom : sum) \ + reduction(+ : sum) + for (int i = 0; i < 1024; ++i) + sum += (_Float16)1; + + _Float16 maximum = 0; +#pragma omp target teams distribute parallel for map(tofrom : maximum) \ + reduction(max : maximum) + for (int i = 0; i < 1024; ++i) { + _Float16 v = (_Float16)(i % 10); + maximum = v > maximum ? v : maximum; + } + + // Control: 'short' is also 2 bytes but takes the integer path, which was + // never broken. It must stay correct. + short int_sum = 0; +#pragma omp target teams distribute parallel for map(tofrom : int_sum) \ + reduction(+ : int_sum) + for (int i = 0; i < 1024; ++i) + int_sum += 1; + + // CHECK: sum = 1024 + // CHECK: maximum = 9 + // CHECK: int_sum = 1024 + printf("sum = %g\n", (double)sum); + printf("maximum = %g\n", (double)maximum); + printf("int_sum = %d\n", int_sum); + return 0; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
