[clang] [llvm] [mlir] [clang][NVPTX] Add support for pzo in f32/f16x2/bf16x2 to FP4 conversions (PR #225318)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 22 00:51:40 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Dharuni R Acharya (DharuniRAcharya)
<details>
<summary>Changes</summary>
This patch adds support for `pzo` variants and `rz` rounding mode to
existing `f32/f16x2/bf16x2` to `FP4` (`e2m1x2`) 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
---
Patch is 26.54 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/225318.diff
9 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsNVPTX.td (+21)
- (modified) clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp (+12)
- (modified) clang/test/CodeGen/builtins-nvptx.c (+15-6)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+14-8)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+3-3)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+28-16)
- (added) llvm/test/CodeGen/NVPTX/convert-fp4-pzo.ll (+253)
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+3)
- (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp4x2.mlir (+6-6)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 8f1b43f9af65d..05d47cb893f1b 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -861,6 +861,27 @@ def __nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(
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_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e2m1x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e2m1x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_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 76f6757326eca..0b927d958ab2c 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -1077,6 +1077,18 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
PZO_CVT(bf16x2_to_e3m2x2_rn_relu_satfinite);
PZO_CVT(bf16x2_to_e3m2x2_rz_satfinite);
PZO_CVT(bf16x2_to_e3m2x2_rz_relu_satfinite);
+ PZO_CVT(ff_to_e2m1x2_rn_satfinite);
+ PZO_CVT(ff_to_e2m1x2_rn_relu_satfinite);
+ PZO_CVT(ff_to_e2m1x2_rz_satfinite);
+ PZO_CVT(ff_to_e2m1x2_rz_relu_satfinite);
+ PZO_CVT(f16x2_to_e2m1x2_rn_satfinite);
+ PZO_CVT(f16x2_to_e2m1x2_rn_relu_satfinite);
+ PZO_CVT(f16x2_to_e2m1x2_rz_satfinite);
+ PZO_CVT(f16x2_to_e2m1x2_rz_relu_satfinite);
+ PZO_CVT(bf16x2_to_e2m1x2_rn_satfinite);
+ PZO_CVT(bf16x2_to_e2m1x2_rn_relu_satfinite);
+ PZO_CVT(bf16x2_to_e2m1x2_rz_satfinite);
+ PZO_CVT(bf16x2_to_e2m1x2_rz_relu_satfinite);
#undef PZO_CVT
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index 607095bce481c..111770a0010f8 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -1206,6 +1206,15 @@ __device__ void nvvm_cvt_pzo_sm107f() {
// 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});
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m1x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ __nvvm_ff_to_e2m1x2_rz_relu_satfinite(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m1x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff_to_e2m1x2_rz_relu_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e2m1x2.rn.satfinite(<2 x half> zeroinitializer, i1 true)
+ __nvvm_f16x2_to_e2m1x2_rn_satfinite_pzo({0, 0});
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e2m1x2.rz.satfinite(<2 x bfloat> zeroinitializer, i1 true)
+ __nvvm_bf16x2_to_e2m1x2_rz_satfinite_pzo({0, 0});
+
#endif
// CHECK: ret void
}
@@ -1318,14 +1327,14 @@ __device__ void nvvm_cvt_sm100a_sm101a_sm120a() {
// CHECK_PTX86_SM120a: call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn.relu(i16 19532)
__nvvm_e3m2x2_to_f16x2_rn_relu(0x4C4C);
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e2m1x2_rn_satfinite(1.0f, 1.0f);
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e2m1x2_rn_relu_satfinite(1.0f, 1.0f);
// CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn(i16 76)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index d9b73418394ed..d1777b85eabef 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2136,18 +2136,24 @@ let TargetPrefix = "nvvm" in {
// FP4 conversions.
foreach relu = ["", "_relu"] in {
- def int_nvvm_ff_to_e2m1x2_rn # relu # _satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;
+ foreach rnd = ["rn", "rz"] in {
+ def int_nvvm_ff_to_e2m1x2_ # 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_e2m1x2_ # rnd # relu # _satfinite : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
+
+ def int_nvvm_bf16x2_to_e2m1x2_ # rnd # relu # _satfinite : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
+ }
def int_nvvm_e2m1x2_to_f16x2_rn # relu : NVVMBuiltin,
PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
-
- def int_nvvm_f16x2_to_e2m1x2_rn # relu # _satfinite
- : PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;
- def int_nvvm_bf16x2_to_e2m1x2_rn # relu # _satfinite
- : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
-
foreach satfinite = ["", "_satfinite"] in {
def int_nvvm_e2m1x2_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 f794002971905..755c7884d6c1a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -891,7 +891,7 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
(ins B32:$src1, B32:$src2, CvtMode:$mode),
!strconcat("{{ \n\t",
".reg .b8 \t%e2m1x2_out; \n\t",
- "cvt${mode:base}.satfinite${mode:relu}.e2m1x2.f32 \t%e2m1x2_out, $src1, $src2; \n\t",
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}.e2m1x2.f32 \t%e2m1x2_out, $src1, $src2; \n\t",
"cvt.u16.u8 \t$dst, %e2m1x2_out; \n\t",
"}}"), []>;
@@ -923,7 +923,7 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
(ins B32:$src, CvtMode:$mode),
"{{ \n\t" #
".reg .b8 \t%e2m1x2_out; \n\t" #
- "cvt${mode:base}.satfinite${mode:relu}.e2m1x2.f16x2 \t%e2m1x2_out, $src; \n\t" #
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}.e2m1x2.f16x2 \t%e2m1x2_out, $src; \n\t" #
"cvt.u16.u8 \t$dst, %e2m1x2_out; \n\t" #
"}}", []>,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
@@ -932,7 +932,7 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
(ins B32:$src, CvtMode:$mode),
"{{ \n\t" #
".reg .b8 \t%e2m1x2_out; \n\t" #
- "cvt${mode:base}.satfinite${mode:relu}.e2m1x2.bf16x2 \t%e2m1x2_out, $src; \n\t" #
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}.e2m1x2.bf16x2 \t%e2m1x2_out, $src; \n\t" #
"cvt.u16.u8 \t$dst, %e2m1x2_out; \n\t" #
"}}", []>,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 89a9adf4426d8..492783cf2d981 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -3244,29 +3244,41 @@ let Predicates = [hasNarrowFPToBF16x2ConversionSupport] in {
}
} // let Predicates = [hasNarrowFPToBF16x2ConversionSupport]
-let Predicates = [hasNarrowFPConversionSupport] in {
- def : Pat<(int_nvvm_ff_to_e2m1x2_rn_satfinite f32:$a, f32:$b),
- (CVT_e2m1x2_f32_sf $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff_to_e2m1x2_rn_relu_satfinite f32:$a, f32:$b),
- (CVT_e2m1x2_f32_sf $a, $b, CvtRN_RELU)>;
+foreach relu = ["", "relu"] in {
+ foreach rnd = ["rn", "rz"] in {
+ defvar Relu = !if(!empty(relu), "", "_" # relu);
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # Relu));
+ defvar Suffix = "e2m1x2_" # rnd # Relu;
+ defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
+ defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
+ defvar BF16X2 = ToCvtIntrinsics<"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 F32X2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasNarrowFPConversionSupport]<Predicate>);
+ defvar FPX2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
+
+ def : Pat<(F32X2.sf f32:$a, f32:$b, PZO), (CVT_e2m1x2_f32_sf $a, $b, Mode)>,
+ Requires<F32X2Preds>;
+ def : Pat<(F16X2.sf v2f16:$a, PZO), (CVT_e2m1x2_f16x2_sf $a, Mode)>,
+ Requires<FPX2Preds>;
+ def : Pat<(BF16X2.sf v2bf16:$a, PZO), (CVT_e2m1x2_bf16x2_sf $a, Mode)>,
+ Requires<FPX2Preds>;
+ }
+ }
+}
+let Predicates = [hasNarrowFPConversionSupport] in {
def : Pat<(int_nvvm_e2m1x2_to_f16x2_rn i16:$a),
(CVT_f16x2_e2m1x2 $a, CvtRN)>;
def : Pat<(int_nvvm_e2m1x2_to_f16x2_rn_relu i16:$a),
(CVT_f16x2_e2m1x2 $a, CvtRN_RELU)>;
}
-let Predicates = [hasFP16X2ToNarrowFPConversionSupport] in {
- foreach src_type = ["f16x2", "bf16x2"] in {
- foreach relu = ["", "_relu"] in {
- defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_e2m1x2_rn" # relu # "_satfinite");
- defvar cvt_inst = !cast<NVPTXInst>("CVT_e2m1x2_" # 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 relu = ["", "_relu"] in {
foreach satfinite = ["", "_satfinite"] in {
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp4-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp4-pzo.ll
new file mode 100644
index 0000000000000..f276f664b4c16
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp4-pzo.ll
@@ -0,0 +1,253 @@
+; 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 %}
+
+; E2M1X2 conversions from f32
+
+define i16 @cvt_rn_pzo_e2m1x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_e2m1x2_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_e2m1x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e2m1x2_f32_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m1x2.f32 %e2m1x2_out, %r1, %r2;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rn.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e2m1x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_e2m1x2_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_e2m1x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e2m1x2_f32_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m1x2.f32 %e2m1x2_out, %r1, %r2;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rz.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e2m1x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e2m1x2_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_e2m1x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e2m1x2_f32_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m1x2.f32 %e2m1x2_out, %r1, %r2;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rn.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m1x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m1x2_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_e2m1x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e2m1x2_f32_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m1x2.f32 %e2m1x2_out, %r1, %r2;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rz.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+; E2M1X2 conversions from f16x2
+
+define i16 @cvt_rn_pzo_e2m1x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e2m1x2_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_e2m1x2_f16x2_param_0];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m1x2.f16x2 %e2m1x2_out, %r1;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rn.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e2m1x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e2m1x2_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_e2m1x2_f16x2_param_0];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m1x2.f16x2 %e2m1x2_out, %r1;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rz.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e2m1x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e2m1x2_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_e2m1x2_f16x2_param_0];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m1x2.f16x2 %e2m1x2_out, %r1;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rn.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m1x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m1x2_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_e2m1x2_f16x2_param_0];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m1x2.f16x2 %e2m1x2_out, %r1;
+; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;
+; CHECK-NEXT: }
+; 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.e2m1x2.rz.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+; E2M1X2 conversions from bf16x2
+
+define i16 @cvt_rn_pzo_e2m1x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e2m1x2_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_e2m1x2_bf16x2_param_...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/225318
More information about the llvm-commits
mailing list