https://github.com/DharuniRAcharya created https://github.com/llvm/llvm-project/pull/222511
This patch adds support for `pzo` variants and `rz` rounding mode to existing `f32/f16x2/bf16x2` to `FP8` (`e4m3x2`, `e5m2x2`) and `FP6` (`e2m3x2`, `e3m2x2`) conversion intrinsics. Also adds `clang builtins` for the new variants. Tests have been verified through `ptxas-13.4`. PTX ISA Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt >From f66a750b5a3f9074a4200d52fcebeb9c2a4f4a3b Mon Sep 17 00:00:00 2001 From: DharuniRAcharya <[email protected]> Date: Thu, 10 Sep 2026 04:55:12 +0000 Subject: [PATCH] [clang][NVPTX] Add support for pzo in f32/f16x2/bf16x2 to FP8 and FP6 conversions This patch adds support for pzo variants and rz rounding mode to existing f32/f16x2/bf16x2 to FP8 (e4m3x2, e5m2x2) and FP6 (e2m3x2, e3m2x2) conversion intrinsics. Also adds clang builtins for the new variants. Tests have been verified through ptxas-13.4. PTX ISA Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt Signed-off-by: DharuniRAcharya <[email protected]> --- clang/include/clang/Basic/BuiltinsNVPTX.td | 78 ++++ clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 50 +++ clang/test/CodeGen/builtins-nvptx.c | 59 ++- llvm/include/llvm/IR/IntrinsicsNVVM.td | 48 ++- llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 18 +- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 115 ++--- llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll | 407 ++++++++++++++++++ llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll | 407 ++++++++++++++++++ mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 19 +- .../Target/LLVMIR/nvvm/convert_fp6x2.mlir | 36 +- .../Target/LLVMIR/nvvm/convert_fp8x2.mlir | 36 +- 11 files changed, 1140 insertions(+), 133 deletions(-) create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td index 3475b721e95e0..8f1b43f9af65d 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.td +++ b/clang/include/clang/Basic/BuiltinsNVPTX.td @@ -783,6 +783,84 @@ def __nvvm_f32x4_to_e3m2x4_rs_relu_satfinite : NVPTXBuiltinSMAndPTX<"_Vector<4, char>(_Vector<4, float>, uint32_t)", SMa<[100, 103]>, PTX87>; +def __nvvm_ff_to_e4m3x2_rz : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e4m3x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rz : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e4m3x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e4m3x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e4m3x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e4m3x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e5m2x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; + +def __nvvm_f16x2_to_e4m3x2_rz : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e4m3x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rz : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e4m3x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e4m3x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e4m3x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e4m3x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e5m2x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; + +def __nvvm_bf16x2_to_e4m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e4m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e4m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; + +def __nvvm_ff_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; +def __nvvm_ff_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>; + +def __nvvm_f16x2_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; +def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>; + +def __nvvm_bf16x2_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; +def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>; + def __nvvm_ff_to_e2m1x2_rn_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SMa<[100, 101, 120]>, PTX86>; def __nvvm_ff_to_e2m1x2_rn_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SMa<[100, 101, 120]>, PTX86>; diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp index 06c5069d6f984..1c4ea2d5e9143 100644 --- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp @@ -1028,6 +1028,56 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, PZO_CVT(f2bf16_rz_satfinite); PZO_CVT(f2bf16_rz_relu_satfinite); + PZO_CVT(ff_to_e4m3x2_rn); + PZO_CVT(ff_to_e4m3x2_rn_relu); + PZO_CVT(ff_to_e4m3x2_rz); + PZO_CVT(ff_to_e4m3x2_rz_relu); + PZO_CVT(ff_to_e5m2x2_rn); + PZO_CVT(ff_to_e5m2x2_rn_relu); + PZO_CVT(ff_to_e5m2x2_rz); + PZO_CVT(ff_to_e5m2x2_rz_relu); + PZO_CVT(f16x2_to_e4m3x2_rn); + PZO_CVT(f16x2_to_e4m3x2_rn_relu); + PZO_CVT(f16x2_to_e4m3x2_rz); + PZO_CVT(f16x2_to_e4m3x2_rz_relu); + PZO_CVT(f16x2_to_e5m2x2_rn); + PZO_CVT(f16x2_to_e5m2x2_rn_relu); + PZO_CVT(f16x2_to_e5m2x2_rz); + PZO_CVT(f16x2_to_e5m2x2_rz_relu); + PZO_CVT(bf16x2_to_e4m3x2_rn_satfinite); + PZO_CVT(bf16x2_to_e4m3x2_rn_relu_satfinite); + PZO_CVT(bf16x2_to_e4m3x2_rz_satfinite); + PZO_CVT(bf16x2_to_e4m3x2_rz_relu_satfinite); + PZO_CVT(bf16x2_to_e5m2x2_rn_satfinite); + PZO_CVT(bf16x2_to_e5m2x2_rn_relu_satfinite); + PZO_CVT(bf16x2_to_e5m2x2_rz_satfinite); + PZO_CVT(bf16x2_to_e5m2x2_rz_relu_satfinite); + + PZO_CVT(ff_to_e2m3x2_rn_satfinite); + PZO_CVT(ff_to_e2m3x2_rn_relu_satfinite); + PZO_CVT(ff_to_e2m3x2_rz_satfinite); + PZO_CVT(ff_to_e2m3x2_rz_relu_satfinite); + PZO_CVT(ff_to_e3m2x2_rn_satfinite); + PZO_CVT(ff_to_e3m2x2_rn_relu_satfinite); + PZO_CVT(ff_to_e3m2x2_rz_satfinite); + PZO_CVT(ff_to_e3m2x2_rz_relu_satfinite); + PZO_CVT(f16x2_to_e2m3x2_rn_satfinite); + PZO_CVT(f16x2_to_e2m3x2_rn_relu_satfinite); + PZO_CVT(f16x2_to_e2m3x2_rz_satfinite); + PZO_CVT(f16x2_to_e2m3x2_rz_relu_satfinite); + PZO_CVT(f16x2_to_e3m2x2_rn_satfinite); + PZO_CVT(f16x2_to_e3m2x2_rn_relu_satfinite); + PZO_CVT(f16x2_to_e3m2x2_rz_satfinite); + PZO_CVT(f16x2_to_e3m2x2_rz_relu_satfinite); + PZO_CVT(bf16x2_to_e2m3x2_rn_satfinite); + PZO_CVT(bf16x2_to_e2m3x2_rn_relu_satfinite); + PZO_CVT(bf16x2_to_e2m3x2_rz_satfinite); + PZO_CVT(bf16x2_to_e2m3x2_rz_relu_satfinite); + PZO_CVT(bf16x2_to_e3m2x2_rn_satfinite); + PZO_CVT(bf16x2_to_e3m2x2_rn_relu_satfinite); + PZO_CVT(bf16x2_to_e3m2x2_rz_satfinite); + PZO_CVT(bf16x2_to_e3m2x2_rz_relu_satfinite); + #undef PZO_CVT case NVPTX::BI__nvvm_fma_rn_f16: diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index bed1498236b06..b74e9c74d84a9 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -1187,6 +1187,25 @@ __device__ void nvvm_cvt_pzo_sm107f() { __nvvm_f2f16_rz_satfinite_pzo(1); // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rz.relu.satfinite(float 1.000000e+00, i1 true) __nvvm_f2f16_rz_relu_satfinite_pzo(1); + + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 false) + __nvvm_ff_to_e4m3x2_rz_relu(1, 1); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 true) + __nvvm_ff_to_e4m3x2_rz_relu_pzo(1, 1); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> zeroinitializer, i1 true) + __nvvm_f16x2_to_e5m2x2_rn_pzo({0, 0}); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.satfinite(<2 x bfloat> zeroinitializer, i1 true) + __nvvm_bf16x2_to_e4m3x2_rz_satfinite_pzo({0, 0}); + + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + __nvvm_ff_to_e2m3x2_rz_relu_satfinite(1, 1); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true) + __nvvm_ff_to_e2m3x2_rz_relu_satfinite_pzo(1, 1); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> zeroinitializer, i1 true) + __nvvm_f16x2_to_e3m2x2_rn_satfinite_pzo({0, 0}); + // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.satfinite(<2 x bfloat> zeroinitializer, i1 true) + __nvvm_bf16x2_to_e2m3x2_rz_satfinite_pzo({0, 0}); + #endif // CHECK: ret void } @@ -1194,22 +1213,22 @@ __device__ void nvvm_cvt_pzo_sm107f() { // CHECK-LABEL: nvvm_cvt_sm89 __device__ void nvvm_cvt_sm89() { #if (PTX >= 81) && (__CUDA_ARCH__ >= 890) - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e4m3x2_rn(1.0f, 1.0f); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e4m3x2_rn_relu(1.0f, 1.0f); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e5m2x2_rn(1.0f, 1.0f); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e5m2x2_rn_relu(1.0f, 1.0f); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> splat (half 1.000000e+00)) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> splat (half 1.000000e+00), i1 false) __nvvm_f16x2_to_e4m3x2_rn({1.0f16, 1.0f16}); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> splat (half 1.000000e+00)) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> splat (half 1.000000e+00), i1 false) __nvvm_f16x2_to_e4m3x2_rn_relu({1.0f16, 1.0f16}); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> splat (half 1.000000e+00)) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> splat (half 1.000000e+00), i1 false) __nvvm_f16x2_to_e5m2x2_rn({1.0f16, 1.0f16}); - // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> splat (half 1.000000e+00)) + // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> splat (half 1.000000e+00), i1 false) __nvvm_f16x2_to_e5m2x2_rn_relu({1.0f16, 1.0f16}); // CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e4m3x2.to.f16x2.rn(i16 18504) @@ -1259,24 +1278,24 @@ __device__ void nvvm_cvt_sm100a_sm101a_sm120a() { #if (PTX >= 86) && \ (__CUDA_ARCH_FEAT_SM100_ALL || __CUDA_ARCH_FEAT_SM101_ALL || \ __CUDA_ARCH_FEAT_SM120_ALL) - // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e2m3x2_rn_satfinite(1.0f, 1.0f); - // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e2m3x2_rn_relu_satfinite(1.0f, 1.0f); - // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e3m2x2_rn_satfinite(1.0f, 1.0f); - // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) - // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00) + // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) + // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false) __nvvm_ff_to_e3m2x2_rn_relu_satfinite(1.0f, 1.0f); // CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 19532) diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 87ae664ce4c17..bc7b6d58653d6 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -2057,18 +2057,25 @@ let TargetPrefix = "nvvm" in { foreach type = ["e4m3x2", "e5m2x2"] in { foreach relu = ["", "_relu"] in { - def int_nvvm_ff_to_ # type # _rn # relu : NVVMBuiltin, - PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>; + foreach rnd = ["rn", "rz"] in { + def int_nvvm_ff_to_ # type # _ # rnd # relu : NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], + [llvm_float_ty, llvm_float_ty, llvm_i1_ty], + [ImmArg<ArgIndex<2>, DefaultValue<0>>]>; - def int_nvvm_f16x2_to_ # type # _rn # relu : NVVMBuiltin, - PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>; + def int_nvvm_f16x2_to_ # type # _ # rnd # relu : NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty, llvm_i1_ty], + [ImmArg<ArgIndex<1>, DefaultValue<0>>]>; + + def int_nvvm_bf16x2_to_ # type # _ # rnd # relu # _satfinite : + NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i1_ty], + [ImmArg<ArgIndex<1>, DefaultValue<0>>]>; + } def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin, PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>; - - def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite - : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>; - + foreach satfinite = ["", "_satfinite"] in { def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0 : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>; @@ -2115,18 +2122,27 @@ let TargetPrefix = "nvvm" in { // FP6 conversions. foreach type = ["e2m3x2", "e3m2x2"] in { foreach relu = ["", "_relu"] in { - def int_nvvm_ff_to_ # type # _rn # relu # _satfinite : NVVMBuiltin, - PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>; + foreach rnd = ["rn", "rz"] in { + def int_nvvm_ff_to_ # type # _ # rnd # relu # _satfinite : + NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], + [llvm_float_ty, llvm_float_ty, llvm_i1_ty], + [ImmArg<ArgIndex<2>, DefaultValue<0>>]>; + + def int_nvvm_f16x2_to_ # type # _ # rnd # relu # _satfinite : + NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty, llvm_i1_ty], + [ImmArg<ArgIndex<1>, DefaultValue<0>>]>; + + def int_nvvm_bf16x2_to_ # type # _ # rnd # relu # _satfinite : + NVVMBuiltin, + PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i1_ty], + [ImmArg<ArgIndex<1>, DefaultValue<0>>]>; + } def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin, PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>; - - def int_nvvm_f16x2_to_ # type # _rn # relu # _satfinite - : PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>; - def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite - : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>; - foreach satfinite = ["", "_satfinite"] in { def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0 : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>; diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 846bcc0557464..140efeeffe55c 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -741,16 +741,19 @@ let hasSideEffects = false in { def _f32 : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src1, B32:$src2), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f32">, + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + F8Name # "x2.f32">, Requires<[PTX81, SM89]>; def _f16x2 : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f16x2">, + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + F8Name # "x2.f16x2">, Requires<[PTX81, SM89]>; def _bf16x2 : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.bf16x2">, + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + F8Name # "x2.bf16x2">, Requires<[hasFP16X2ToNarrowFPConversionSupport]>; } @@ -849,14 +852,17 @@ let Predicates = [hasS2F6X2ConversionSupport] in { def _f32_sf : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src1, B32:$src2), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.f32">; + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + FP6Name # "x2.f32">; def _f16x2_sf : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.f16x2">, + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + FP6Name # "x2.f16x2">, Requires<[hasFP16X2ToNarrowFPConversionSupport]>; def _bf16x2_sf : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode), - "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.bf16x2">, + "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." # + FP6Name # "x2.bf16x2">, Requires<[hasFP16X2ToNarrowFPConversionSupport]>; } diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 220ef64732830..0ccdb9abacd48 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -2838,39 +2838,42 @@ def : Pat<(int_nvvm_ull2d_rp i64:$a), (CVT_f64_u64 $a, CvtRP)>; def : Pat<(int_nvvm_f2h_rn_ftz f32:$a), (CVT_f16_f32 $a, CvtRN_FTZ)>; def : Pat<(int_nvvm_f2h_rn f32:$a), (CVT_f16_f32 $a, CvtRN)>; -let Predicates = [hasFP8ConversionSupport] in { - def : Pat<(int_nvvm_ff_to_e4m3x2_rn f32:$a, f32:$b), - (CVT_e4m3x2_f32 $a, $b, CvtRN)>; - def : Pat<(int_nvvm_ff_to_e4m3x2_rn_relu f32:$a, f32:$b), - (CVT_e4m3x2_f32 $a, $b, CvtRN_RELU)>; - def : Pat<(int_nvvm_ff_to_e5m2x2_rn f32:$a, f32:$b), - (CVT_e5m2x2_f32 $a, $b, CvtRN)>; - def : Pat<(int_nvvm_ff_to_e5m2x2_rn_relu f32:$a, f32:$b), - (CVT_e5m2x2_f32 $a, $b, CvtRN_RELU)>; -} - -let Predicates = [hasFP16X2ToNarrowFPConversionSupport] in { - foreach dst_type = ["e4m3x2", "e5m2x2"] in { - foreach relu = ["", "_relu"] in { - defvar intrin = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # dst_type # "_rn" # relu # "_satfinite"); - defvar cvt_inst = !cast<NVPTXInst>("CVT_"# dst_type #"_bf16x2"); - defvar cvt_mode = !cast<PatLeaf>("CvtRN" # !toupper(relu)); - def : Pat<(intrin v2bf16:$a), - (cvt_inst $a, cvt_mode)>; +foreach dst = ["e4m3x2", "e5m2x2"] in { + foreach relu = ["", "_relu"] in { + foreach rnd = ["rn", "rz"] in { + defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu)); + defvar Suffix = dst # "_" # rnd # relu; + defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix); + defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix); + defvar BF16X2 = + !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite"); + + foreach pzo = [0, 1] in { + defvar PZO = !if(pzo, -1, 0); + defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value; + defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1)); + defvar FFPreds = + !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>, + [hasFP8ConversionSupport]<Predicate>); + defvar BF16X2Preds = + !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>, + [hasFP16X2ToNarrowFPConversionSupport]<Predicate>); + + def : Pat<(FF f32:$a, f32:$b, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_f32") $a, $b, Mode)>, + Requires<FFPreds>; + def : Pat<(F16X2 v2f16:$a, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_f16x2") $a, Mode)>, + Requires<FFPreds>; + def : Pat<(BF16X2 v2bf16:$a, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>, + Requires<BF16X2Preds>; + } } } } let Predicates = [hasFP8ConversionSupport] in { - def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn v2f16:$a), - (CVT_e4m3x2_f16x2 $a, CvtRN)>; - def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn_relu v2f16:$a), - (CVT_e4m3x2_f16x2 $a, CvtRN_RELU)>; - def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn v2f16:$a), - (CVT_e5m2x2_f16x2 $a, CvtRN)>; - def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn_relu v2f16:$a), - (CVT_e5m2x2_f16x2 $a, CvtRN_RELU)>; - def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn i16:$a), (CVT_f16x2_e4m3x2 $a, CvtRN)>; def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn_relu i16:$a), @@ -2918,16 +2921,41 @@ let Predicates = [hasS2F6X2ConversionSupport] in { (CVT_bf16x2_s2f6x2_sf_scale $a, $scale, CvtRN_RELU)>; } -let Predicates = [hasNarrowFPConversionSupport] in { - def : Pat<(int_nvvm_ff_to_e2m3x2_rn_satfinite f32:$a, f32:$b), - (CVT_e2m3x2_f32_sf $a, $b, CvtRN)>; - def : Pat<(int_nvvm_ff_to_e2m3x2_rn_relu_satfinite f32:$a, f32:$b), - (CVT_e2m3x2_f32_sf $a, $b, CvtRN_RELU)>; - def : Pat<(int_nvvm_ff_to_e3m2x2_rn_satfinite f32:$a, f32:$b), - (CVT_e3m2x2_f32_sf $a, $b, CvtRN)>; - def : Pat<(int_nvvm_ff_to_e3m2x2_rn_relu_satfinite f32:$a, f32:$b), - (CVT_e3m2x2_f32_sf $a, $b, CvtRN_RELU)>; +foreach dst = ["e2m3x2", "e3m2x2"] in { + foreach relu = ["", "_relu"] in { + foreach rnd = ["rn", "rz"] in { + defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu)); + defvar Suffix = dst # "_" # rnd # relu # "_satfinite"; + defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix); + defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix); + defvar BF16X2 = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix); + + foreach pzo = [0, 1] in { + defvar PZO = !if(pzo, -1, 0); + defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value; + defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1)); + defvar FFPreds = + !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>, + [hasNarrowFPConversionSupport]<Predicate>); + defvar FPX2Preds = + !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>, + [hasFP16X2ToNarrowFPConversionSupport]<Predicate>); + + def : Pat<(FF f32:$a, f32:$b, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_f32_sf") $a, $b, Mode)>, + Requires<FFPreds>; + def : Pat<(F16X2 v2f16:$a, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf") $a, Mode)>, + Requires<FPX2Preds>; + def : Pat<(BF16X2 v2bf16:$a, PZO), + (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>, + Requires<FPX2Preds>; + } + } + } +} +let Predicates = [hasNarrowFPConversionSupport] in { def : Pat<(int_nvvm_e2m3x2_to_f16x2_rn i16:$a), (CVT_f16x2_e2m3x2 $a, CvtRN)>; def : Pat<(int_nvvm_e2m3x2_to_f16x2_rn_relu i16:$a), @@ -2938,19 +2966,6 @@ let Predicates = [hasNarrowFPConversionSupport] in { (CVT_f16x2_e3m2x2 $a, CvtRN_RELU)>; } -let Predicates = [hasFP16X2ToNarrowFPConversionSupport] in { - foreach src_type = ["f16x2", "bf16x2"] in { - foreach dst_type = ["e2m3x2", "e3m2x2"] in { - foreach relu = ["", "_relu"] in { - defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_" # dst_type # "_rn" # relu # "_satfinite"); - defvar cvt_inst = !cast<NVPTXInst>("CVT_" # dst_type # "_" # src_type # "_sf"); - defvar cvt_mode = !cast<PatLeaf>("CvtRN" # !toupper(relu)); - def : Pat<(intrin B32:$a), (cvt_inst $a, cvt_mode)>; - } - } - } -} - let Predicates = [hasNarrowFPToBF16x2ConversionSupport] in { foreach src_type = ["e2m3x2", "e3m2x2"] in { foreach relu = ["", "_relu"] in { diff --git a/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll new file mode 100644 index 0000000000000..01c1554e99644 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll @@ -0,0 +1,407 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; 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 %} + +; E2M3X2 conversions from f32 + +define i16 @cvt_rn_pzo_e2m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_pzo_e2m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e2m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e2m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_pzo_e2m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e2m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rz.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e2m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e2m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e2m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e2m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e2m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +; E3M2X2 conversions from f32 + +define i16 @cvt_rn_pzo_e3m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_pzo_e3m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e3m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e3m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_pzo_e3m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e3m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rz.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e3m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e3m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e3m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e3m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rz.relu.satfinite(float %f1, float %f2, i1 true) + ret i16 %val +} + +; E2M3X2 conversions from f16x2 + +define i16 @cvt_rn_pzo_e2m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_pzo_e2m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e2m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_pzo_e2m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rz.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e2m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e2m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e2m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rz.relu.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +; E3M2X2 conversions from f16x2 + +define i16 @cvt_rn_pzo_e3m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_pzo_e3m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e3m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_pzo_e3m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rz.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e3m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e3m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rz.relu.satfinite(<2 x half> %a, i1 true) + ret i16 %val +} + +; E2M3X2 conversions from bf16x2 + +define i16 @cvt_rn_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_pzo_e2m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_pzo_e2m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e2m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +; E3M2X2 conversions from bf16x2 + +define i16 @cvt_rn_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_pzo_e3m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_pzo_e3m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rz.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} diff --git a/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll new file mode 100644 index 0000000000000..01ea380d61403 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll @@ -0,0 +1,407 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; 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 %} + +; E4M3X2 conversions from f32 + +define i16 @cvt_rn_pzo_e4m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_pzo_e4m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e4m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e4m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_pzo_e4m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e4m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rz(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e4m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e4m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e4m3x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e4m3x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e4m3x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float %f1, float %f2, i1 true) + ret i16 %val +} + +; E5M2X2 conversions from f32 + +define i16 @cvt_rn_pzo_e5m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_pzo_e5m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e5m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e5m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_pzo_e5m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e5m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rz(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e5m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e5m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %f1, float %f2, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e5m2x2_f32(float %f1, float %f2) { +; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_f32( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_f32_param_0]; +; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e5m2x2_f32_param_1]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.f32 %rs1, %r1, %r2; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rz.relu(float %f1, float %f2, i1 true) + ret i16 %val +} + +; E4M3X2 conversions from f16x2 + +define i16 @cvt_rn_pzo_e4m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_pzo_e4m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e4m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_pzo_e4m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rz(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e4m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e4m3x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e4m3x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rz.relu(<2 x half> %a, i1 true) + ret i16 %val +} + +; E5M2X2 conversions from f16x2 + +define i16 @cvt_rn_pzo_e5m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_pzo_e5m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e5m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_pzo_e5m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rz(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e5m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e5m2x2_f16x2(<2 x half> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_f16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_f16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.f16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rz.relu(<2 x half> %a, i1 true) + ret i16 %val +} + +; E4M3X2 conversions from bf16x2 + +define i16 @cvt_rn_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_pzo_e4m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_pzo_e4m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e4m3x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +; E5M2X2 conversions from bf16x2 + +define i16 @cvt_rn_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_pzo_e5m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_pzo_e5m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rz.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rn_relu_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} + +define i16 @cvt_rz_relu_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) { +; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_bf16x2( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_bf16x2_param_0]; +; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.bf16x2 %rs1, %r1; +; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; +; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true) + ret i16 %val +} diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 80a0094190383..ffad9fe6d0d04 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -2085,7 +2085,8 @@ def NVVM_ConvertF32x2ToF6x2Op : NVVM_Op<"convert.f32x2.to.f6x2"> { string llvmBuilder = [{ auto intId = NVVM::ConvertF32x2ToF6x2Op::getIntrinsicID($dstTy, $relu); - llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a, $b}); + llvm::Value *packedI16 = + createIntrinsicCall(builder, intId, {$a, $b, builder.getInt1(false)}); if(op.getDst().getType().isInteger(16)) $dst = packedI16; else @@ -2127,7 +2128,8 @@ class NVVM_ConvertFPx2ToF6x2Op<string srcType> string llvmBuilder = [{ auto intId = NVVM::Convert}]#srcType#[{x2ToF6x2Op::getIntrinsicID($dstTy, $relu); - llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$src}); + llvm::Value *packedI16 = + createIntrinsicCall(builder, intId, {$src, builder.getInt1(false)}); if(op.getDst().getType().isInteger(16)) $dst = packedI16; else @@ -2181,7 +2183,10 @@ def NVVM_ConvertF32x2ToF8x2Op : NVVM_Op<"convert.f32x2.to.f8x2"> { string llvmBuilder = [{ auto intId = NVVM::ConvertF32x2ToF8x2Op::getIntrinsicID($dstTy, $rnd, $sat, $relu); - llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a, $b}); + llvm::SmallVector<llvm::Value *> args = {$a, $b}; + if (!mlir::isa<mlir::Float8E8M0FNUType>($dstTy)) + args.push_back(builder.getInt1(false)); + llvm::Value *packedI16 = createIntrinsicCall(builder, intId, args); if(op.getDst().getType().isInteger(16)) $dst = packedI16; else @@ -2227,7 +2232,8 @@ def NVVM_ConvertF16x2ToF8x2Op : NVVM_Op<"convert.f16x2.to.f8x2"> { string llvmBuilder = [{ auto intId = NVVM::ConvertF16x2ToF8x2Op::getIntrinsicID($dstTy, $relu); - llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a}); + llvm::Value *packedI16 = + createIntrinsicCall(builder, intId, {$a, builder.getInt1(false)}); if(op.getDst().getType().isInteger(16)) $dst = packedI16; else @@ -2277,7 +2283,10 @@ def NVVM_ConvertBF16x2ToF8x2Op : NVVM_Op<"convert.bf16x2.to.f8x2"> { string llvmBuilder = [{ auto intId = NVVM::ConvertBF16x2ToF8x2Op::getIntrinsicID($dstTy, $rnd, $sat, $relu); - llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$src}); + llvm::SmallVector<llvm::Value *> args = {$src}; + if (!mlir::isa<mlir::Float8E8M0FNUType>($dstTy)) + args.push_back(builder.getInt1(false)); + llvm::Value *packedI16 = createIntrinsicCall(builder, intId, args); if(op.getDst().getType().isInteger(16)) $dst = packedI16; else diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir index 6cf4e8d2b98fd..21f23521405b1 100644 --- a/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir +++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir @@ -2,28 +2,28 @@ // CHECK-LABEL: @convert_f32x2_to_fp6x2_e2m3 llvm.func @convert_f32x2_to_fp6x2_e2m3(%srcA : f32, %srcB : f32) { - //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false) %res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : i16 (f6E2M3FN) - //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false) %res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB relu = true : i16 (f6E2M3FN) llvm.return } // CHECK-LABEL: @convert_f32x2_to_fp6x2_e3m2 llvm.func @convert_f32x2_to_fp6x2_e3m2(%srcA : f32, %srcB : f32) { - //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false) %res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : i16 (f6E3M2FN) - //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false) %res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB relu = true : i16 (f6E3M2FN) llvm.return } // CHECK-LABEL: @convert_f32x2_to_fp6x2_vector llvm.func @convert_f32x2_to_fp6x2_vector(%srcA : f32, %srcB : f32) { - //CHECK: %[[res0:.*]] = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %[[res0:.*]] = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false) //CHECK-NEXT: %{{.*}} = bitcast i16 %[[res0]] to <2 x i8> %res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : vector<2xi8> (f6E2M3FN) - //CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}) + //CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false) //CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8> %res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : vector<2xi8> (f6E3M2FN) llvm.return @@ -33,28 +33,28 @@ llvm.func @convert_f32x2_to_fp6x2_vector(%srcA : f32, %srcB : f32) { // CHECK-LABEL: @convert_f16x2_to_fp6x2_e2m3 llvm.func @convert_f16x2_to_fp6x2_e2m3(%srcA : vector<2xf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}, i1 false) %res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> i16 (f6E2M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %{{.*}}, i1 false) %res2 = nvvm.convert.f16x2.to.f6x2 %srcA relu = true : vector<2xf16> -> i16 (f6E2M3FN) llvm.return } // CHECK-LABEL: @convert_f16x2_to_fp6x2_e3m2 llvm.func @convert_f16x2_to_fp6x2_e3m2(%srcA : vector<2xf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}, i1 false) %res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> i16 (f6E3M2FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %{{.*}}, i1 false) %res2 = nvvm.convert.f16x2.to.f6x2 %srcA relu = true : vector<2xf16> -> i16 (f6E3M2FN) llvm.return } // CHECK-LABEL: @convert_f16x2_to_fp6x2_vector llvm.func @convert_f16x2_to_fp6x2_vector(%srcA : vector<2xf16>) { - // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}) + // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8> %res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> vector<2xi8> (f6E2M3FN) - // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}) + // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8> %res2 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> vector<2xi8> (f6E3M2FN) llvm.return @@ -64,28 +64,28 @@ llvm.func @convert_f16x2_to_fp6x2_vector(%srcA : vector<2xf16>) { // CHECK-LABEL: @convert_bf16x2_to_fp6x2_e2m3 llvm.func @convert_bf16x2_to_fp6x2_e2m3(%srcA : vector<2xbf16>, %scale_factor : i16) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> i16 (f6E2M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res2 = nvvm.convert.bf16x2.to.f6x2 %srcA relu = true : vector<2xbf16> -> i16 (f6E2M3FN) llvm.return } // CHECK-LABEL: @convert_bf16x2_to_fp6x2_e3m2 llvm.func @convert_bf16x2_to_fp6x2_e3m2(%srcA : vector<2xbf16>, %scale_factor : i16) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> i16 (f6E3M2FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res2 = nvvm.convert.bf16x2.to.f6x2 %srcA relu = true : vector<2xbf16> -> i16 (f6E3M2FN) llvm.return } // CHECK-LABEL: @convert_bf16x2_to_fp6x2_vector llvm.func @convert_bf16x2_to_fp6x2_vector(%srcA : vector<2xbf16>, %scale_factor : i16) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %{{.*}} to <2 x i8> %res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> vector<2xi8> (f6E2M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %{{.*}} to <2 x i8> %res2 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> vector<2xi8> (f6E3M2FN) llvm.return diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir index 6a8745e065902..058e28be90283 100644 --- a/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir +++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir @@ -4,18 +4,18 @@ // CHECK-LABEL: @convert_f32x2_to_f8x2_e4m3 llvm.func @convert_f32x2_to_f8x2_e4m3(%srcA : f32, %srcB : f32) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}, i1 false) %res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : i16 (f8E4M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false) %res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : i16 (f8E4M3FN) llvm.return } // CHECK-LABEL: @convert_f32x2_to_f8x2_e5m2 llvm.func @convert_f32x2_to_f8x2_e5m2(%srcA : f32, %srcB : f32) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %{{.*}}, float %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %{{.*}}, float %{{.*}}, i1 false) %res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : i16 (f8E5M2) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %{{.*}}, float %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false) %res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : i16 (f8E5M2) llvm.return } @@ -35,10 +35,10 @@ llvm.func @convert_f32x2_to_f8x2_ue8m0(%srcA : f32, %srcB : f32) { // CHECK-LABEL: @convert_f32x2_to_f8x2_vector_return llvm.func @convert_f32x2_to_f8x2_vector_return(%srcA : f32, %srcB : f32) { - // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}) + // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8> %res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : vector<2xi8> (f8E4M3FN) - // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}) + // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8> %res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : vector<2xi8> (f8E4M3FN) llvm.return @@ -48,28 +48,28 @@ llvm.func @convert_f32x2_to_f8x2_vector_return(%srcA : f32, %srcB : f32) { // CHECK-LABEL: @convert_f16x2_to_f8x2_e4m3 llvm.func @convert_f16x2_to_f8x2_e4m3(%src : vector<2xf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}, i1 false) %res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> i16 (f8E4M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %{{.*}}, i1 false) %res2 = nvvm.convert.f16x2.to.f8x2 %src relu = true : vector<2xf16> -> i16 (f8E4M3FN) llvm.return } // CHECK-LABEL: @convert_f16x2_to_f8x2_e5m2 llvm.func @convert_f16x2_to_f8x2_e5m2(%src : vector<2xf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}, i1 false) %res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> i16 (f8E5M2) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %{{.*}}, i1 false) %res2 = nvvm.convert.f16x2.to.f8x2 %src relu = true : vector<2xf16> -> i16 (f8E5M2) llvm.return } // CHECK-LABEL: @convert_f16x2_to_f8x2_vector_return llvm.func @convert_f16x2_to_f8x2_vector_return(%src : vector<2xf16>) { - // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}) + // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8> %res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> vector<2xi8> (f8E4M3FN) - // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}) + // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8> %res2 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> vector<2xi8> (f8E5M2) llvm.return @@ -93,18 +93,18 @@ llvm.func @convert_bf16x2_to_f8x2_ue8m0(%src : vector<2xbf16>) { // CHECK-LABEL: @convert_bf16x2_to_f8x2_e4m3 llvm.func @convert_bf16x2_to_f8x2_e4m3(%srcA : vector<2xbf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res1 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> : vector<2xbf16> -> i16 (f8E4M3FN) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res2 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> i16 (f8E4M3FN) llvm.return } // CHECK-LABEL: @convert_bf16x2_to_f8x2_e5m2 llvm.func @convert_bf16x2_to_f8x2_e5m2(%srcA : vector<2xbf16>) { - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res1 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E5M2) - // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false) %res2 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> vector<2xi8> (f8E5M2) llvm.return } @@ -117,10 +117,10 @@ llvm.func @convert_bf16x2_to_f8x2_vector_return(%src : vector<2xbf16>) { // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> %{{.*}}) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8> %res2 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rp> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E8M0FNU) - // CHECK: %[[res3:.*]] = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %[[res3:.*]] = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res3]] to <2 x i8> %res3 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rn> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E4M3FN) - // CHECK: %[[res4:.*]] = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}) + // CHECK: %[[res4:.*]] = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false) // CHECK-NEXT: %{{.*}} = bitcast i16 %[[res4]] to <2 x i8> %res4 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> vector<2xi8> (f8E5M2) llvm.return _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
