Author: Srinivasa Ravi Date: 2025-02-14T11:11:44+05:30 New Revision: bd860f986406b6630e49b1836b3c208acd721d3e
URL: https://github.com/llvm/llvm-project/commit/bd860f986406b6630e49b1836b3c208acd721d3e DIFF: https://github.com/llvm/llvm-project/commit/bd860f986406b6630e49b1836b3c208acd721d3e.diff LOG: [NVPTX] Add intrinsics for redux.sync f32 instructions (#126664) Adds NVVM intrinsics, NVPTX codegen and Clang builtins for `redux.sync` f32 instructions introduced in ptx8.6 for sm_100a. Tests added in `CodeGen/NVPTX/redux-sync.ll` and `CodeGenCUDA/redux-builtins.cu` and verified through ptxas 12.8.0. PTX Spec Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-redux-sync Added: clang/test/CodeGenCUDA/redux-f32-builtins.cu llvm/test/CodeGen/NVPTX/redux-sync-f32.ll Modified: clang/include/clang/Basic/BuiltinsNVPTX.td llvm/include/llvm/IR/IntrinsicsNVVM.td llvm/lib/Target/NVPTX/NVPTXIntrinsics.td Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td index 9d24a992563a4..327dc88cffdb4 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.td +++ b/clang/include/clang/Basic/BuiltinsNVPTX.td @@ -669,6 +669,14 @@ def __nvvm_redux_sync_umax : NVPTXBuiltinSMAndPTX<"unsigned int(unsigned int, in def __nvvm_redux_sync_and : NVPTXBuiltinSMAndPTX<"int(int, int)", SM_80, PTX70>; def __nvvm_redux_sync_xor : NVPTXBuiltinSMAndPTX<"int(int, int)", SM_80, PTX70>; def __nvvm_redux_sync_or : NVPTXBuiltinSMAndPTX<"int(int, int)", SM_80, PTX70>; +def __nvvm_redux_sync_fmin : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmin_abs : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmin_NaN : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmin_abs_NaN : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmax : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmax_abs : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmax_NaN : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; +def __nvvm_redux_sync_fmax_abs_NaN : NVPTXBuiltinSMAndPTX<"float(float, int)", SM_100a, PTX86>; // Membar diff --git a/clang/test/CodeGenCUDA/redux-f32-builtins.cu b/clang/test/CodeGenCUDA/redux-f32-builtins.cu new file mode 100644 index 0000000000000..7359fb0006991 --- /dev/null +++ b/clang/test/CodeGenCUDA/redux-f32-builtins.cu @@ -0,0 +1,34 @@ +// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" "-target-feature" "+ptx86" "-target-cpu" "sm_100a" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s +// RUN: %clang_cc1 "-triple" "nvptx64-nvidia-cuda" "-target-feature" "+ptx86" "-target-cpu" "sm_100a" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s + +// CHECK: define{{.*}} void @_Z6kernelPf(ptr noundef %out_f) +__attribute__((global)) void kernel(float* out_f) { + float a = 3.0; + int i = 0; + + out_f[i++] = __nvvm_redux_sync_fmin(a, 0xFF); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmin + + out_f[i++] = __nvvm_redux_sync_fmin_abs(a, 0xFF); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmin.abs + + out_f[i++] = __nvvm_redux_sync_fmin_NaN(a, 0xF0); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmin.NaN + + out_f[i++] = __nvvm_redux_sync_fmin_abs_NaN(a, 0x0F); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmin.abs.NaN + + out_f[i++] = __nvvm_redux_sync_fmax(a, 0xFF); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmax + + out_f[i++] = __nvvm_redux_sync_fmax_abs(a, 0x01); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmax.abs + + out_f[i++] = __nvvm_redux_sync_fmax_NaN(a, 0xF1); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmax.NaN + + out_f[i++] = __nvvm_redux_sync_fmax_abs_NaN(a, 0x10); + // CHECK: call contract float @llvm.nvvm.redux.sync.fmax.abs.NaN + + // CHECK: ret void +} diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 1a6aa17b531c6..7ef270f3256a6 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -4824,6 +4824,18 @@ def int_nvvm_redux_sync_xor : ClangBuiltin<"__nvvm_redux_sync_xor">, def int_nvvm_redux_sync_or : ClangBuiltin<"__nvvm_redux_sync_or">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrInaccessibleMemOnly, IntrNoCallback]>; + +// redux.sync.op.{abs}.{NaN}.f32 dst, src, membermask; +foreach binOp = ["min", "max"] in { + foreach abs = ["", "_abs"] in { + foreach NaN = ["", "_NaN"] in { + def int_nvvm_redux_sync_f # binOp # abs # NaN : + ClangBuiltin<!strconcat("__nvvm_redux_sync_f", binOp, abs, NaN)>, + Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty], + [IntrConvergent, IntrInaccessibleMemOnly, IntrNoCallback]>; + } + } +} // // WGMMA fence instructions diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 6a99a4b3b4f69..f20502521829e 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -328,6 +328,25 @@ defm REDUX_SYNC_AND : REDUX_SYNC<"and", "b32", int_nvvm_redux_sync_and>; defm REDUX_SYNC_XOR : REDUX_SYNC<"xor", "b32", int_nvvm_redux_sync_xor>; defm REDUX_SYNC_OR : REDUX_SYNC<"or", "b32", int_nvvm_redux_sync_or>; +multiclass REDUX_SYNC_F<string BinOp, string abs, string NaN> { + defvar intr_name = "int_nvvm_redux_sync_f" # BinOp # !subst(".", "_", abs) # !subst(".", "_", NaN); + + def : NVPTXInst<(outs Float32Regs:$dst), + (ins Float32Regs:$src, Int32Regs:$mask), + "redux.sync." # BinOp # abs # NaN # ".f32 $dst, $src, $mask;", + [(set f32:$dst, (!cast<Intrinsic>(intr_name) f32:$src, Int32Regs:$mask))]>, + Requires<[hasPTX<86>, hasSM100a]>; +} + +defm REDUX_SYNC_FMIN : REDUX_SYNC_F<"min", "", "">; +defm REDUX_SYNC_FMIN_ABS : REDUX_SYNC_F<"min", ".abs", "">; +defm REDUX_SYNC_FMIN_NAN: REDUX_SYNC_F<"min", "", ".NaN">; +defm REDUX_SYNC_FMIN_ABS_NAN: REDUX_SYNC_F<"min", ".abs", ".NaN">; +defm REDUX_SYNC_FMAX : REDUX_SYNC_F<"max", "", "">; +defm REDUX_SYNC_FMAX_ABS : REDUX_SYNC_F<"max", ".abs", "">; +defm REDUX_SYNC_FMAX_NAN: REDUX_SYNC_F<"max", "", ".NaN">; +defm REDUX_SYNC_FMAX_ABS_NAN: REDUX_SYNC_F<"max", ".abs", ".NaN">; + } // isConvergent = true //----------------------------------- diff --git a/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll new file mode 100644 index 0000000000000..af113e75fd143 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/redux-sync-f32.ll @@ -0,0 +1,139 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck %s +; RUN: %if ptxas-12.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %} + +declare float @llvm.nvvm.redux.sync.fmin(float, i32) +define float @redux_sync_fmin(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmin( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmin_param_1]; +; CHECK-NEXT: redux.sync.min.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmin(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmin.abs(float, i32) +define float @redux_sync_fmin_abs(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmin_abs( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_abs_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmin_abs_param_1]; +; CHECK-NEXT: redux.sync.min.abs.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmin.abs(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmin.NaN(float, i32) +define float @redux_sync_fmin_NaN(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmin_NaN( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_NaN_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmin_NaN_param_1]; +; CHECK-NEXT: redux.sync.min.NaN.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmin.NaN(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmin.abs.NaN(float, i32) +define float @redux_sync_fmin_abs_NaN(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmin_abs_NaN( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmin_abs_NaN_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmin_abs_NaN_param_1]; +; CHECK-NEXT: redux.sync.min.abs.NaN.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmin.abs.NaN(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmax(float, i32) +define float @redux_sync_fmax(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmax( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmax_param_1]; +; CHECK-NEXT: redux.sync.max.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmax(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmax.abs(float, i32) +define float @redux_sync_fmax_abs(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmax_abs( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_abs_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmax_abs_param_1]; +; CHECK-NEXT: redux.sync.max.abs.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmax.abs(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmax.NaN(float, i32) +define float @redux_sync_fmax_NaN(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmax_NaN( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_NaN_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmax_NaN_param_1]; +; CHECK-NEXT: redux.sync.max.NaN.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmax.NaN(float %src, i32 %mask) + ret float %val +} + +declare float @llvm.nvvm.redux.sync.fmax.abs.NaN(float, i32) +define float @redux_sync_fmax_abs_NaN(float %src, i32 %mask) { +; CHECK-LABEL: redux_sync_fmax_abs_NaN( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<2>; +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [redux_sync_fmax_abs_NaN_param_0]; +; CHECK-NEXT: ld.param.u32 %r1, [redux_sync_fmax_abs_NaN_param_1]; +; CHECK-NEXT: redux.sync.max.abs.NaN.f32 %f2, %f1, %r1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %val = call float @llvm.nvvm.redux.sync.fmax.abs.NaN(float %src, i32 %mask) + ret float %val +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
