llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-nvptx @llvm/pr-subscribers-mlir Author: Srinivasa Ravi (Wolfram70) <details> <summary>Changes</summary> This change removes the `pzo` modifier from the `fp16x2` conversion intrinsics with `rs` rounding mode. This was added in https://github.com/llvm/llvm-project/pull/214667 due to a bug in the 13.4 PTX ISA developer-preview, and this change aligns it with the latest PTX ISA docs. PTX ISA reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt --- Patch is 22.54 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/225012.diff 9 Files Affected: - (modified) clang/test/CodeGen/builtins-nvptx.c (+16-16) - (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+2-2) - (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+2-2) - (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+11-18) - (modified) llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll (+2-2) - (removed) llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll (-135) - (modified) llvm/test/CodeGen/NVPTX/convert-sm103a.ll (+2) - (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+5-2) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir (+8-8) ``````````diff diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 607095bce481c..cd3eaff9d4c53 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -1394,51 +1394,51 @@ __device__ void nvvm_cvt_sm100a_sm103a() { typedef __bf16 bf16x2 __attribute__((ext_vector_type(2))); typedef char uint8x4 __attribute__((ext_vector_type(4))); -// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x half> %[[R1]], ptr %r1 -// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x half> %[[R1]], ptr %r1 f16x2 r1 = __nvvm_ff2f16x2_rs(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x half> %[[R2]], ptr %r2 -// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x half> %[[R2]], ptr %r2 f16x2 r2 = __nvvm_ff2f16x2_rs_relu(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x half> %[[R3]], ptr %r3 -// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x half> %[[R3]], ptr %r3 f16x2 r3 = __nvvm_ff2f16x2_rs_satfinite(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x half> %[[R4]], ptr %r4 -// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x half> %[[R4]], ptr %r4 f16x2 r4 = __nvvm_ff2f16x2_rs_relu_satfinite(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x bfloat> %[[R5]], ptr %r5 -// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x bfloat> %[[R5]], ptr %r5 bf16x2 r5 = __nvvm_ff2bf16x2_rs(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x bfloat> %[[R6]], ptr %r6 -// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x bfloat> %[[R6]], ptr %r6 bf16x2 r6 = __nvvm_ff2bf16x2_rs_relu(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x bfloat> %[[R7]], ptr %r7 -// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x bfloat> %[[R7]], ptr %r7 bf16x2 r7 = __nvvm_ff2bf16x2_rs_satfinite(1.0f, 1.0f, 0); -// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM100a: store <2 x bfloat> %[[R8]], ptr %r8 -// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false) +// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0) // CHECK_PTX87_SM103a: store <2 x bfloat> %[[R8]], ptr %r8 bf16x2 r8 = __nvvm_ff2bf16x2_rs_relu_satfinite(1.0f, 1.0f, 0); diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index ff12175180d14..162dbe482d0ef 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -2064,10 +2064,10 @@ let TargetPrefix = "nvvm" in { foreach relu = ["", "_relu"] in { foreach satfinite = ["", "_satfinite"] in { def int_nvvm_ff2f16x2_rs # relu # satfinite : NVVMBuiltin, - PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>; + PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>; def int_nvvm_ff2bf16x2_rs # relu # satfinite : NVVMBuiltin, - PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>; + PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>; } } diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index c130a2f351ff3..11b7314010806 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -739,13 +739,13 @@ let hasSideEffects = false in { BasicFlagsNVPTXInst<(outs RC:$dst), (ins B32:$src1, B32:$src2, B32:$src3), (ins CvtMode:$mode), - "cvt${mode:base}${mode:relu}${mode:pzo}." # FromName # ".f32">; + "cvt${mode:base}${mode:relu}." # FromName # ".f32">; def _f32_rs_sf : BasicFlagsNVPTXInst<(outs RC:$dst), (ins B32:$src1, B32:$src2, B32:$src3), (ins CvtMode:$mode), - "cvt${mode:base}${mode:relu}.satfinite${mode:pzo}." # FromName # ".f32">; + "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">; } defm CVT_f16x2 : CVT_FROM_FLOAT_V2_RS<"f16x2", B32>; diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index f54fd1b348af8..38ce9fa8e69b5 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -3027,28 +3027,21 @@ foreach rnd = ["rn", "rz"] in { } } -foreach relu = ["", "_relu"] in { - defvar BaseMode = !cast<PatLeaf>("CvtRS" # !toupper(relu)); - - foreach pzo = [0, 1] in { - defvar PZO = !if(pzo, -1, 0); - defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value; - defvar Preds = - !listconcat( - [hasConvertWithStochasticRounding]<Predicate>, - !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>)); +let Predicates = [hasConvertWithStochasticRounding] in { + foreach relu = ["", "_relu"] in { + defvar Mode = !cast<PatLeaf>("CvtRS" # !toupper(relu)); defvar BF16X2 = ToCvtIntrinsics<"ff2bf16x2_rs" # relu>; defvar F16X2 = ToCvtIntrinsics<"ff2f16x2_rs" # relu>; - def : Pat<(BF16X2.base f32:$a, f32:$b, i32:$c, PZO), - (CVT_bf16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>; - def : Pat<(BF16X2.sf f32:$a, f32:$b, i32:$c, PZO), - (CVT_bf16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>; - def : Pat<(F16X2.base f32:$a, f32:$b, i32:$c, PZO), - (CVT_f16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>; - def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c, PZO), - (CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>; + def : Pat<(BF16X2.base f32:$a, f32:$b, i32:$c), + (CVT_bf16x2_f32_rs $a, $b, $c, Mode)>; + def : Pat<(BF16X2.sf f32:$a, f32:$b, i32:$c), + (CVT_bf16x2_f32_rs_sf $a, $b, $c, Mode)>; + def : Pat<(F16X2.base f32:$a, f32:$b, i32:$c), + (CVT_f16x2_f32_rs $a, $b, $c, Mode)>; + def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c), + (CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>; } } diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll index 6861daf25aacf..339984387d9ca 100644 --- a/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll +++ b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll @@ -2,8 +2,8 @@ ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | FileCheck %s ; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %} -; Conversions to fp16 types with the preserve-zero-order modifier. The RS -; (stochastic rounding) variants need sm_107a and live in convert-fp16-rs-pzo.ll. +; Conversions to fp16 types with the preserve-zero-order modifier. The modifier +; is not supported with the RS (stochastic rounding) rounding mode. ; F16X2 conversions diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll deleted file mode 100644 index 58eacc36bcf55..0000000000000 --- a/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll +++ /dev/null @@ -1,135 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 -; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | FileCheck %s -; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %} - -; F16X2 conversions - -define <2 x half> @cvt_rs_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_pzo_f16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_f16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_f16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_f16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.pzo.f16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x half> %val -} - -define <2 x half> @cvt_rs_relu_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_relu_pzo_f16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_f16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_f16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_f16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.relu.pzo.f16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x half> %val -} - -define <2 x half> @cvt_rs_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_sf_pzo_f16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_f16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_f16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_f16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x half> %val -} - -define <2 x half> @cvt_rs_relu_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_relu_sf_pzo_f16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_f16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_f16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_f16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x half> %val -} - -; BF16X2 conversions - -define <2 x bfloat> @cvt_rs_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_pzo_bf16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_bf16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_bf16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_bf16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.pzo.bf16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x bfloat> %val -} - -define <2 x bfloat> @cvt_rs_relu_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_relu_pzo_bf16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_bf16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_bf16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_bf16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.relu.pzo.bf16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x bfloat> %val -} - -define <2 x bfloat> @cvt_rs_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_sf_pzo_bf16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_bf16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_bf16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_bf16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x bfloat> %val -} - -define <2 x bfloat> @cvt_rs_relu_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) { -; CHECK-LABEL: cvt_rs_relu_sf_pzo_bf16x2_f32( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_0]; -; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_1]; -; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_2]; -; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3; -; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; - %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true) - ret <2 x bfloat> %val -} diff --git a/llvm/test/CodeGen/NVPTX/convert-sm103a.ll b/llvm/test/CodeGen/NVPTX/convert-sm103a.ll index 75a394e4764ed..6a452d1d66baa 100644 --- a/llvm/test/CodeGen/NVPTX/convert-sm103a.ll +++ b/llvm/test/CodeGen/NVPTX/convert-sm103a.ll @@ -1,8 +1,10 @@ ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx87 | FileCheck %s ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck %s +; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | FileCheck %s ; RUN: %if ptxas-sm_100a && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx87 | %ptxas-verify -arch=sm_100a %} ; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %} +; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %} ; F16X2 conversions diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp index 46e9ab29bc4a0..f9f8a3c4387d7 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp @@ -5674,8 +5674,10 @@ ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op, if (op.getRandomBits()) args.push_back(mt.lookupValue(op.getRandomBits())); + // The PZO modifier is not supported with the RS rounding mode. // TODO: Add support for PZO modifier - args.push_back(builder.getInt1(false)); + if (op.getRnd() != FPRoundingMode::RS) + args.push_back(builder.getInt1(false)); switch (op.getRnd()) { case FPRoundingMode::RN: @@ -5726,7 +5728,8 @@ ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op, args.push_back(mt.lookupValue(op.getRandomBits())); // TODO: Add support for PZO modifier - args.push_back(builder.getInt1(false)); + if (op.getRnd() != FPRoundingMode::RS) + args.push_back(builder.getInt1(false)); switch (op.getRnd()) { case FPRoundingMode::RN: diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir index ddb6defe270bb..6338f6a15030e 100644 --- a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir +++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir @@ -30,13 +30,13 @@ llvm.func @convert_f32x2_to_f16x2_rz(%srcA : f32, %srcB : f32) { // CHECK-LABEL: @convert_f32x2_to_f16x2_rs_stochastic llvm.func @convert_f32x2_to_f16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits : i32) { - // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false) + // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}) %res1 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> : vector<2xf16> - // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false) + // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/225012 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
