[clang] [llvm] [clang][NVPTX] Add support for scaled::n1::ue8m0 in FP8, FP6 and FP4 conversions (PR #227652)
Dharuni R Acharya via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 09:58:12 PDT 2026
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/227652
>From 9d2aab4fc4d02b158d09b1d18bd8a3f37d97eb37 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Wed, 30 Sep 2026 10:39:41 +0000
Subject: [PATCH 1/3] [NVPTX] Add support for scaled::n1::ue8m0 in FP8, FP6 and
FP4 conversions
This patch adds support for scaled::n1::ue8m0 to existing
f32/f16x2/bf16x2 to FP8 (e4m3x2, e5m2x2), FP6 (e2m3x2, e3m2x2)
and FP4 (e2m1x2) conversion intrinsics.
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 <dharunira at nvidia.com>
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 36 ++++
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 42 +++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 36 ++++
llvm/test/CodeGen/NVPTX/convert-fp4-scale.ll | 122 +++++++++++++
llvm/test/CodeGen/NVPTX/convert-fp6-scale.ll | 178 +++++++++++++++++++
llvm/test/CodeGen/NVPTX/convert-fp8-scale.ll | 178 +++++++++++++++++++
6 files changed, 592 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp4-scale.ll
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp6-scale.ll
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp8-scale.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 29bb506b6afba5d..1e03ddab81eb6cd 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2101,15 +2101,27 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ def int_nvvm_ff_to_ # type # _ # rnd # relu # _scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
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_f16x2_to_ # type # _ # rnd # relu # _scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, 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_bf16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
@@ -2138,14 +2150,26 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ def int_nvvm_ff_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<3>, 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_f16x2_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, 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_bf16x2_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
def int_nvvm_e2m1x2_to_f16x2_rn # relu : NVVMBuiltin,
@@ -2173,16 +2197,28 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ def int_nvvm_ff_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<3>, 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_f16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, 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_bf16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
+ : PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 8973b3b9fd66247..31c664c04144ea3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -183,6 +183,7 @@ def hasTensormapReplaceSupport : SubtargetPredicate;
// - TMA S2G im2col_w mode support
// - tcgen05.commit shared mem A variants.
// - conversions involving ue5m3x2
+// - conversions involving scaled::n1::ue8m0
def hasRubinFamilySupport : PredOr<[SM107f]>;
// Checks tcgen05.shift instruction support.
@@ -753,6 +754,16 @@ let hasSideEffects = false in {
defm CVT_f16x2 : CVT_FROM_FLOAT_V2_RS<"f16x2", B32>;
defm CVT_bf16x2 : CVT_FROM_FLOAT_V2_RS<"bf16x2", B32>;
+ class CVT_TO_NARROWFP_SCALE_N1<dag ins, string DstType, string SrcType, string SrcOps>
+ : NVPTXInst<(outs B16:$dst), ins,
+ "{{ \n\t" #
+ ".reg .b8 \t%b8_in; \n\t" #
+ "cvt.u8.u16 \t%b8_in, $scale; \n\t" #
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}.scaled::n1::ue8m0." #
+ DstType # "." # SrcType # " \t$dst, " # SrcOps # ", %b8_in; \n\t" #
+ "}}", []>,
+ Requires<[hasRubinFamilySupport]>;
+
// FP8 conversions.
multiclass CVT_TO_F8X2<string F8Name> {
def _f32 :
@@ -772,6 +783,12 @@ let hasSideEffects = false in {
"cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
F8Name # "x2.bf16x2">,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
+ def _f32_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B32:$src2, B16:$scale, CvtMode:$mode),
+ F8Name # "x2", "f32", "$src1, $src2">;
+ def _f16x2_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ F8Name # "x2", "f16x2", "$src1">;
+ def _bf16x2_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ F8Name # "x2", "bf16x2", "$src1">;
}
defm CVT_e4m3x2 : CVT_TO_F8X2<"e4m3">;
@@ -881,6 +898,12 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
"cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
FP6Name # "x2.bf16x2">,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
+ def _f32_sf_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B32:$src2, B16:$scale, CvtMode:$mode),
+ FP6Name # "x2", "f32", "$src1, $src2">;
+ def _f16x2_sf_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ FP6Name # "x2", "f16x2", "$src1">;
+ def _bf16x2_sf_scale_n1 : CVT_TO_NARROWFP_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ FP6Name # "x2", "bf16x2", "$src1">;
}
defm CVT_e2m3x2 : CVT_TO_FP6X2<"e2m3">;
@@ -937,6 +960,25 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
"}}", []>,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
+ class CVT_TO_E2M1X2_SCALE_N1<dag ins, string SrcType, string SrcOps>
+ : NVPTXInst<(outs B16:$dst), ins,
+ "{{ \n\t" #
+ ".reg .b8 \t%e2m1x2_out; \n\t" #
+ ".reg .b8 \t%b8_in; \n\t" #
+ "cvt.u8.u16 \t%b8_in, $scale; \n\t" #
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}.scaled::n1::ue8m0.e2m1x2." #
+ SrcType # " \t%e2m1x2_out, " # SrcOps # ", %b8_in; \n\t" #
+ "cvt.u16.u8 \t$dst, %e2m1x2_out; \n\t" #
+ "}}", []>,
+ Requires<[hasRubinFamilySupport]>;
+
+ def CVT_e2m1x2_f32_sf_scale_n1 : CVT_TO_E2M1X2_SCALE_N1<(ins B32:$src1, B32:$src2, B16:$scale, CvtMode:$mode),
+ "f32", "$src1, $src2">;
+ def CVT_e2m1x2_f16x2_sf_scale_n1 : CVT_TO_E2M1X2_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ "f16x2", "$src1">;
+ def CVT_e2m1x2_bf16x2_sf_scale_n1 : CVT_TO_E2M1X2_SCALE_N1<(ins B32:$src1, B16:$scale, CvtMode:$mode),
+ "bf16x2", "$src1">;
+
// UE8M0x2 conversions.
class CVT_f32_to_ue8m0x2<string sat = ""> :
BasicFlagsNVPTXInst<(outs B16:$dst),
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index e6ce5fe5be48ac5..ffcbbcea6e65648 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -3105,6 +3105,9 @@ foreach dst = ["e4m3x2", "e5m2x2"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
+ defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_scale_n1_ue8m0");
+ defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_scale_n1_ue8m0");
+ defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3121,6 +3124,16 @@ foreach dst = ["e4m3x2", "e5m2x2"] in {
Requires<F32X2Preds>;
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>,
Requires<BF16X2Preds>;
+
+ def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f32_scale_n1") $a, $b, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(F16Scale v2f16:$a, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_scale_n1") $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_scale_n1") $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
}
}
}
@@ -3182,6 +3195,9 @@ foreach dst = ["e2m3x2", "e3m2x2"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
+ defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3198,6 +3214,16 @@ foreach dst = ["e2m3x2", "e3m2x2"] in {
Requires<FPX2Preds>;
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>,
Requires<FPX2Preds>;
+
+ def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f32_sf_scale_n1") $a, $b, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(F16Scale v2f16:$a, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf_scale_n1") $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf_scale_n1") $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
}
}
}
@@ -3235,6 +3261,9 @@ foreach relu = ["", "relu"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
+ defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3251,6 +3280,13 @@ foreach relu = ["", "relu"] in {
Requires<FPX2Preds>;
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (CVT_e2m1x2_bf16x2_sf $a, Mode)>,
Requires<FPX2Preds>;
+
+ def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO), (CVT_e2m1x2_f32_sf_scale_n1 $a, $b, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(F16Scale v2f16:$a, i16:$s, PZO), (CVT_e2m1x2_f16x2_sf_scale_n1 $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
+ def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO), (CVT_e2m1x2_bf16x2_sf_scale_n1 $a, $s, Mode)>,
+ Requires<[hasRubinFamilySupport]>;
}
}
}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp4-scale.ll b/llvm/test/CodeGen/NVPTX/convert-fp4-scale.ll
new file mode 100644
index 000000000000000..adea6290ac449c2
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp4-scale.ll
@@ -0,0 +1,122 @@
+; 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 scaled::n1::ue8m0 conversions
+
+define i16 @cvt_rn_e2m1x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m1x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m1x2_f32_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_e2m1x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m1x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m1x2.f32 %e2m1x2_out, %r1, %r2, %b8_in;
+; CHECK-NEXT: cvt.u16.u8 %rs2, %e2m1x2_out;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite.scale.n1.ue8m0(float %a, float %b, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m1x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m1x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; 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_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e2m1x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e2m1x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e2m1x2.f32 %e2m1x2_out, %r1, %r2, %b8_in;
+; CHECK-NEXT: cvt.u16.u8 %rs2, %e2m1x2_out;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m1x2.rz.relu.satfinite.scale.n1.ue8m0(float %a, float %b, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e2m1x2_f16x2_scale(<2 x half> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m1x2_f16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m1x2_f16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m1x2_f16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m1x2.f16x2 %e2m1x2_out, %r1, %b8_in;
+; CHECK-NEXT: cvt.u16.u8 %rs2, %e2m1x2_out;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m1x2.rn.satfinite.scale.n1.ue8m0(<2 x half> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e2m1x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m1x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m1x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m1x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m1x2.bf16x2 %e2m1x2_out, %r1, %b8_in;
+; CHECK-NEXT: cvt.u16.u8 %rs2, %e2m1x2_out;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m1x2.rn.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m1x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m1x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e2m1x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e2m1x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_out;
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e2m1x2.bf16x2 %e2m1x2_out, %r1, %b8_in;
+; CHECK-NEXT: cvt.u16.u8 %rs2, %e2m1x2_out;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m1x2.rz.relu.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale, i1 true)
+ ret i16 %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp6-scale.ll b/llvm/test/CodeGen/NVPTX/convert-fp6-scale.ll
new file mode 100644
index 000000000000000..f441600ee159591
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp6-scale.ll
@@ -0,0 +1,178 @@
+; 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 scaled::n1::ue8m0 conversions
+
+define i16 @cvt_rn_e2m3x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m3x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m3x2_f32_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_e2m3x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m3x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m3x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite.scale.n1.ue8m0(float %a, float %b, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m3x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; 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_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e2m3x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e2m3x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e2m3x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite.scale.n1.ue8m0(float %a, float %b, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e2m3x2_f16x2_scale(<2 x half> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m3x2_f16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m3x2_f16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m3x2_f16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m3x2.f16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite.scale.n1.ue8m0(<2 x half> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e2m3x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e2m3x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e2m3x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e2m3x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e2m3x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m3x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; 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_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e2m3x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e2m3x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.relu.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+; E3M2X2 scaled::n1::ue8m0 conversions
+
+define i16 @cvt_rn_e3m2x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e3m2x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e3m2x2_f32_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_e3m2x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e3m2x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e3m2x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite.scale.n1.ue8m0(float %a, float %b, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_pzo_e3m2x2_f16x2_scale(<2 x half> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_pzo_e3m2x2_f16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_f16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_pzo_e3m2x2_f16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.pzo.scaled::n1::ue8m0.e3m2x2.f16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite.scale.n1.ue8m0(<2 x half> %a, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_e3m2x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rz_e3m2x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_e3m2x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_e3m2x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.scaled::n1::ue8m0.e3m2x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rz.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale)
+ ret i16 %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp8-scale.ll b/llvm/test/CodeGen/NVPTX/convert-fp8-scale.ll
new file mode 100644
index 000000000000000..2ca1fa0bb7c0695
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp8-scale.ll
@@ -0,0 +1,178 @@
+; 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 scaled::n1::ue8m0 conversions
+
+define i16 @cvt_rn_e4m3x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e4m3x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e4m3x2_f32_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_e4m3x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e4m3x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e4m3x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.scale.n1.ue8m0(float %a, float %b, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e4m3x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; 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_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e4m3x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e4m3x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e4m3x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu.scale.n1.ue8m0(float %a, float %b, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e4m3x2_f16x2_scale(<2 x half> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e4m3x2_f16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e4m3x2_f16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e4m3x2_f16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e4m3x2.f16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.scale.n1.ue8m0(<2 x half> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_e4m3x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e4m3x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e4m3x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e4m3x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e4m3x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e4m3x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; 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_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_relu_pzo_e4m3x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.scaled::n1::ue8m0.e4m3x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.relu.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+; E5M2X2 scaled::n1::ue8m0 conversions
+
+define i16 @cvt_rn_e5m2x2_f32_scale(float %a, float %b, i16 %scale) {
+; CHECK-LABEL: cvt_rn_e5m2x2_f32_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_e5m2x2_f32_scale_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_e5m2x2_f32_scale_param_1];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_e5m2x2_f32_scale_param_2];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n1::ue8m0.e5m2x2.f32 %rs2, %r1, %r2, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.scale.n1.ue8m0(float %a, float %b, i16 %scale)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_pzo_e5m2x2_f16x2_scale(<2 x half> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rn_pzo_e5m2x2_f16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_f16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rn_pzo_e5m2x2_f16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.pzo.scaled::n1::ue8m0.e5m2x2.f16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.scale.n1.ue8m0(<2 x half> %a, i16 %scale, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_e5m2x2_bf16x2_scale(<2 x bfloat> %a, i16 %scale) {
+; CHECK-LABEL: cvt_rz_e5m2x2_bf16x2_scale(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_e5m2x2_bf16x2_scale_param_0];
+; CHECK-NEXT: ld.param::func.b16 %rs1, [cvt_rz_e5m2x2_bf16x2_scale_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %b8_in;
+; CHECK-NEXT: cvt.u8.u16 %b8_in, %rs1;
+; CHECK-NEXT: cvt.rz.satfinite.scaled::n1::ue8m0.e5m2x2.bf16x2 %rs2, %r1, %b8_in;
+; CHECK-NEXT: }
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rz.satfinite.scale.n1.ue8m0(<2 x bfloat> %a, i16 %scale)
+ ret i16 %val
+}
>From f4720d4bd44c74b2d6cc0edf0d6b4299bf6c1471 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Wed, 30 Sep 2026 14:06:06 +0000
Subject: [PATCH 2/3] Address comments
---
clang/include/clang/Basic/BuiltinsNVPTX.td | 129 +++++++++++++++++++++
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 63 ++++++++++
clang/test/CodeGen/builtins-nvptx.c | 23 ++++
llvm/include/llvm/IR/IntrinsicsNVVM.td | 37 +++---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 46 ++++----
5 files changed, 260 insertions(+), 38 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 05d47cb893f1b8e..8915d89f01be385 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -907,6 +907,135 @@ def __nvvm_bf16x2_to_ue8m0x2_rp_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<
def __nvvm_ue8m0x2_to_bf16x2 : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(short)", SMa<[100, 101, 120]>, PTX86>;
+def __nvvm_ff_to_e4m3x2_rn_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rn_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rn_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rn_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e4m3x2_rn_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rn_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_relu_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rn_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rn_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_relu_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e4m3x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+
+def __nvvm_ff_to_e2m3x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+
+def __nvvm_ff_to_e2m1x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(float, float, short)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>, short)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0 : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>, short)", SM_107f, PTX94>;
+
// FNS
let Attributes = [NoThrow] in {
def __nvvm_fns : NVPTXBuiltinPTX<"unsigned int(unsigned int, unsigned int, int)", PTX60>;
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index b14996e79c38d0a..e304eed8a1aab26 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -1061,6 +1061,31 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
PZO_CVT(bf16x2_to_e5m2x2_rz_satfinite);
PZO_CVT(bf16x2_to_e5m2x2_rz_relu_satfinite);
+ PZO_CVT(ff_to_e4m3x2_rn_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e4m3x2_rn_relu_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e4m3x2_rz_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e4m3x2_rz_relu_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e5m2x2_rn_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e5m2x2_rn_relu_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e5m2x2_rz_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e5m2x2_rz_relu_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e4m3x2_rn_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e4m3x2_rn_relu_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e4m3x2_rz_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e4m3x2_rz_relu_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e5m2x2_rn_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e5m2x2_rn_relu_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e5m2x2_rz_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e5m2x2_rz_relu_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e4m3x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e4m3x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e4m3x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e4m3x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e5m2x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e5m2x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e5m2x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e5m2x2_rz_relu_satfinite_scale_n1_ue8m0);
+
PZO_CVT(ff_to_e2m3x2_rn_satfinite);
PZO_CVT(ff_to_e2m3x2_rn_relu_satfinite);
PZO_CVT(ff_to_e2m3x2_rz_satfinite);
@@ -1086,6 +1111,31 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
PZO_CVT(bf16x2_to_e3m2x2_rz_satfinite);
PZO_CVT(bf16x2_to_e3m2x2_rz_relu_satfinite);
+ PZO_CVT(ff_to_e2m3x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m3x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e3m2x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e3m2x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m3x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m3x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m3x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e3m2x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e3m2x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e3m2x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0);
+
PZO_CVT(ff_to_e2m1x2_rn_satfinite);
PZO_CVT(ff_to_e2m1x2_rn_relu_satfinite);
PZO_CVT(ff_to_e2m1x2_rz_satfinite);
@@ -1099,6 +1149,19 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
PZO_CVT(bf16x2_to_e2m1x2_rz_satfinite);
PZO_CVT(bf16x2_to_e2m1x2_rz_relu_satfinite);
+ PZO_CVT(ff_to_e2m1x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m1x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(ff_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(f16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m1x2_rn_relu_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0);
+ PZO_CVT(bf16x2_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0);
+
#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 7ecbf3b5895eb58..acd7d40277bfa28 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -1239,6 +1239,29 @@ __device__ void nvvm_cvt_pzo_sm107f() {
// 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});
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.scale.n1.ue8m0(float 1.000000e+00, float 1.000000e+00, i16 1, i1 false)
+ __nvvm_ff_to_e4m3x2_rn_scale_n1_ue8m0(1, 1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.scale.n1.ue8m0(float 1.000000e+00, float 1.000000e+00, i16 1, i1 true)
+ __nvvm_ff_to_e4m3x2_rn_scale_n1_ue8m0_pzo(1, 1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rz.relu.scale.n1.ue8m0(<2 x half> zeroinitializer, i16 1, i1 false)
+ __nvvm_f16x2_to_e5m2x2_rz_relu_scale_n1_ue8m0({0, 0}, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.satfinite.scale.n1.ue8m0(<2 x bfloat> zeroinitializer, i16 1, i1 true)
+ __nvvm_bf16x2_to_e4m3x2_rz_satfinite_scale_n1_ue8m0_pzo({0, 0}, 1);
+
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite.scale.n1.ue8m0(float 1.000000e+00, float 1.000000e+00, i16 1, i1 false)
+ __nvvm_ff_to_e2m3x2_rn_satfinite_scale_n1_ue8m0(1, 1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e3m2x2.rz.relu.satfinite.scale.n1.ue8m0(<2 x half> zeroinitializer, i16 1, i1 true)
+ __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite_scale_n1_ue8m0_pzo({0, 0}, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite.scale.n1.ue8m0(<2 x bfloat> zeroinitializer, i16 1, i1 false)
+ __nvvm_bf16x2_to_e2m3x2_rn_satfinite_scale_n1_ue8m0({0, 0}, 1);
+
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m1x2.rz.relu.satfinite.scale.n1.ue8m0(float 1.000000e+00, float 1.000000e+00, i16 1, i1 true)
+ __nvvm_ff_to_e2m1x2_rz_relu_satfinite_scale_n1_ue8m0_pzo(1, 1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e2m1x2.rn.satfinite.scale.n1.ue8m0(<2 x half> zeroinitializer, i16 1, i1 false)
+ __nvvm_f16x2_to_e2m1x2_rn_satfinite_scale_n1_ue8m0({0, 0}, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e2m1x2.rz.satfinite.scale.n1.ue8m0(<2 x bfloat> zeroinitializer, i16 1, i1 true)
+ __nvvm_bf16x2_to_e2m1x2_rz_satfinite_scale_n1_ue8m0_pzo({0, 0}, 1);
+
#endif
// CHECK: ret void
}
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 1e03ddab81eb6cd..4e2aaf9ab441318 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2101,25 +2101,26 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
- def int_nvvm_ff_to_ # type # _ # rnd # relu # _scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
- [llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
- [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
+ def int_nvvm_ff_to_ # type # _ # rnd # relu # _scale_n1_ue8m0 : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
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_f16x2_to_ # type # _ # rnd # relu # _scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
- [llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
- [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ def int_nvvm_f16x2_to_ # type # _ # rnd # relu # _scale_n1_ue8m0 : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
+ [llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, 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_bf16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
@@ -2151,7 +2152,8 @@ let TargetPrefix = "nvvm" in {
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
def int_nvvm_ff_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
@@ -2159,7 +2161,8 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
def int_nvvm_f16x2_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
@@ -2167,7 +2170,8 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
def int_nvvm_bf16x2_to_e2m1x2_ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
@@ -2198,7 +2202,8 @@ let TargetPrefix = "nvvm" in {
[llvm_float_ty, llvm_float_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
def int_nvvm_ff_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_float_ty, llvm_float_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
@@ -2207,7 +2212,8 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
def int_nvvm_f16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_v2f16_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
@@ -2216,7 +2222,8 @@ let TargetPrefix = "nvvm" in {
PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
def int_nvvm_bf16x2_to_ # type # _ # rnd # relu # _satfinite_scale_n1_ue8m0
- : PureIntrinsic<[llvm_i16_ty],
+ : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
[llvm_v2bf16_ty, llvm_i16_ty, llvm_i1_ty],
[ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index ffcbbcea6e65648..cfee082f9ca8d12 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2974,11 +2974,11 @@ def : Pat<(int_nvvm_ui2f_rp i32:$a), (CVT_f32_u32 $a, CvtRP)>;
// Utility to look up base and satfinite variants of a conversion intrinsic.
// Missing variants are left unset so this can be used when only one exists.
-class ToCvtIntrinsics<string suffix> {
- Intrinsic base = !if(!exists<Intrinsic>("int_nvvm_" # suffix),
- !cast<Intrinsic>("int_nvvm_" # suffix), ?);
- Intrinsic sf = !if(!exists<Intrinsic>("int_nvvm_" # suffix # "_satfinite"),
- !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite"), ?);
+class ToCvtIntrinsics<string suffix, string tail = ""> {
+ Intrinsic base = !if(!exists<Intrinsic>("int_nvvm_" # suffix # tail),
+ !cast<Intrinsic>("int_nvvm_" # suffix # tail), ?);
+ Intrinsic sf = !if(!exists<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # tail),
+ !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # tail), ?);
}
foreach rnd = ["rn", "rz"] in {
@@ -3105,9 +3105,9 @@ foreach dst = ["e4m3x2", "e5m2x2"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
- defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_scale_n1_ue8m0");
- defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_scale_n1_ue8m0");
- defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar F32Scale = ToCvtIntrinsics<"ff_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar F16Scale = ToCvtIntrinsics<"f16x2_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar BF16Scale = ToCvtIntrinsics<"bf16x2_to_" # Suffix, "_scale_n1_ue8m0">;
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3125,13 +3125,13 @@ foreach dst = ["e4m3x2", "e5m2x2"] in {
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>,
Requires<BF16X2Preds>;
- def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO),
+ def : Pat<(F32Scale.base f32:$a, f32:$b, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_f32_scale_n1") $a, $b, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(F16Scale v2f16:$a, i16:$s, PZO),
+ def : Pat<(F16Scale.base v2f16:$a, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_f16x2_scale_n1") $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO),
+ def : Pat<(BF16Scale.sf v2bf16:$a, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_scale_n1") $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
}
@@ -3195,9 +3195,9 @@ foreach dst = ["e2m3x2", "e3m2x2"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
- defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
- defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
- defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar F32Scale = ToCvtIntrinsics<"ff_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar F16Scale = ToCvtIntrinsics<"f16x2_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar BF16Scale = ToCvtIntrinsics<"bf16x2_to_" # Suffix, "_scale_n1_ue8m0">;
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3215,13 +3215,13 @@ foreach dst = ["e2m3x2", "e3m2x2"] in {
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>,
Requires<FPX2Preds>;
- def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO),
+ def : Pat<(F32Scale.sf f32:$a, f32:$b, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_f32_sf_scale_n1") $a, $b, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(F16Scale v2f16:$a, i16:$s, PZO),
+ def : Pat<(F16Scale.sf v2f16:$a, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf_scale_n1") $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO),
+ def : Pat<(BF16Scale.sf v2bf16:$a, i16:$s, PZO),
(!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf_scale_n1") $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
}
@@ -3261,9 +3261,9 @@ foreach relu = ["", "relu"] in {
defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
- defvar F32Scale = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
- defvar F16Scale = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
- defvar BF16Scale = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite_scale_n1_ue8m0");
+ defvar F32Scale = ToCvtIntrinsics<"ff_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar F16Scale = ToCvtIntrinsics<"f16x2_to_" # Suffix, "_scale_n1_ue8m0">;
+ defvar BF16Scale = ToCvtIntrinsics<"bf16x2_to_" # Suffix, "_scale_n1_ue8m0">;
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
@@ -3281,11 +3281,11 @@ foreach relu = ["", "relu"] in {
def : Pat<(BF16X2.sf v2bf16:$a, PZO), (CVT_e2m1x2_bf16x2_sf $a, Mode)>,
Requires<FPX2Preds>;
- def : Pat<(F32Scale f32:$a, f32:$b, i16:$s, PZO), (CVT_e2m1x2_f32_sf_scale_n1 $a, $b, $s, Mode)>,
+ def : Pat<(F32Scale.sf f32:$a, f32:$b, i16:$s, PZO), (CVT_e2m1x2_f32_sf_scale_n1 $a, $b, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(F16Scale v2f16:$a, i16:$s, PZO), (CVT_e2m1x2_f16x2_sf_scale_n1 $a, $s, Mode)>,
+ def : Pat<(F16Scale.sf v2f16:$a, i16:$s, PZO), (CVT_e2m1x2_f16x2_sf_scale_n1 $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
- def : Pat<(BF16Scale v2bf16:$a, i16:$s, PZO), (CVT_e2m1x2_bf16x2_sf_scale_n1 $a, $s, Mode)>,
+ def : Pat<(BF16Scale.sf v2bf16:$a, i16:$s, PZO), (CVT_e2m1x2_bf16x2_sf_scale_n1 $a, $s, Mode)>,
Requires<[hasRubinFamilySupport]>;
}
}
>From d4a562ee4a69e15c6e9858ba3bfa79a099345d44 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 1 Oct 2026 16:57:45 +0000
Subject: [PATCH 3/3] Address Comments
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 +++++-----
1 file changed, 5 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cfee082f9ca8d12..f0d8abd2417fd0d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2974,11 +2974,11 @@ def : Pat<(int_nvvm_ui2f_rp i32:$a), (CVT_f32_u32 $a, CvtRP)>;
// Utility to look up base and satfinite variants of a conversion intrinsic.
// Missing variants are left unset so this can be used when only one exists.
-class ToCvtIntrinsics<string suffix, string tail = ""> {
- Intrinsic base = !if(!exists<Intrinsic>("int_nvvm_" # suffix # tail),
- !cast<Intrinsic>("int_nvvm_" # suffix # tail), ?);
- Intrinsic sf = !if(!exists<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # tail),
- !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # tail), ?);
+class ToCvtIntrinsics<string suffix, string scale_suffix = ""> {
+ Intrinsic base = !if(!exists<Intrinsic>("int_nvvm_" # suffix # scale_suffix),
+ !cast<Intrinsic>("int_nvvm_" # suffix # scale_suffix), ?);
+ Intrinsic sf = !if(!exists<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # scale_suffix),
+ !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite" # scale_suffix), ?);
}
foreach rnd = ["rn", "rz"] in {
More information about the llvm-commits
mailing list