[llvm] [NVPTX] Add intrinsics for narrow-fp to bf16 conversions (PR #191376)
Srinivasa Ravi via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 13 07:04:23 PDT 2026
https://github.com/Wolfram70 updated https://github.com/llvm/llvm-project/pull/191376
>From 8203a7be1cdc0dcdf85883985f8ab883c8734b2d Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Thu, 9 Apr 2026 06:50:13 +0000
Subject: [PATCH 1/6] [NVPTX] Add intrinsics for narrow-fp to bf16 conversions
Adds the following intrinsics for narrow-fp to bf16 conversions introduced
in PTX 9.2:
- llvm.nvvm.{e4m3x2/e5m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
- llvm.nvvm.{e2m3x2/e3m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
- llvm.nvvm.e2m1x2.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
PTX ISA Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt
---
llvm/docs/NVPTXUsage.rst | 23 ++-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 15 ++
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 5 +
llvm/lib/Target/NVPTX/NVPTX.h | 3 +-
llvm/lib/Target/NVPTX/NVPTX.td | 2 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 33 +++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 41 ++++++
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 5 +
.../CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll | 87 +++++++++++
.../CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll | 135 ++++++++++++++++++
.../CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll | 135 ++++++++++++++++++
11 files changed, 477 insertions(+), 7 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll
create mode 100644 llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll
create mode 100644 llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 70341e87c4770..68144a4669c5e 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -1134,6 +1134,14 @@ The following table describes the rounding modes used across these intrinsics:
| | the input. |
+-----------------------+---------------------------------------------------+
+.. _scale-factor:
+
+Some conversions involve a scale factor which is provided as a packed 16-bit
+integer containing two scaling factors of type ``ue8m0``, one for each input.
+For down conversion, inputs are divided by ``scale_factor`` and then the
+conversion is performed. For up-conversion, inputs are converted to destination
+type and then multiplied by ``scale_factor``.
+
``fp8`` Conversion Intrinsics
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1148,6 +1156,7 @@ Syntax:
declare i16 @llvm.bf16x2.to{.e4m3x2, .e5m2x2}.rn{.relu}.satfinite(<2 x bfloat> %a)
declare i16 @llvm.bf16x2.to.ue8m0x2{.rz, .rp}{.satfinite}(<2 x bfloat> %a)
declare <2 x half> @llvm.nvvm{.e4m3x2, .e5m2x2}.to.f16x2.rn{.relu}(i16 %a)
+ declare <2 x bfloat> @llvm.nvvm{.e4m3x2, .e5m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
declare <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 %a)
declare <4 x i8> @llvm.nvvm.f32x4.to{.e4m3x4, .e5m2x4}.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
@@ -1172,6 +1181,8 @@ specified destination format. The ``satfinite`` modifier is assumed to be
present for conversions involving ``e4m3`` and ``e5m2`` types as the
destination.
+For scale factor, see :ref:`scale-factor`.
+
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
``s2f6`` Conversion Intrinsics
@@ -1204,11 +1215,7 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
-The operand ``%scale_factor`` stores two packed scaling factors of type
-``ue8m0``, one for each input. For down conversion, inputs are divided by
-``scale_factor`` and then the conversion is performed. For up-conversion,
-inputs are converted to destination type and then multiplied by
-``scale_factor``.
+For scale factor, see :ref:`scale-factor`.
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
@@ -1224,6 +1231,7 @@ Syntax:
declare i16 @llvm.nvvm.f16x2.to{.e2m3x2, .e3m2x2}.rn{.relu}.satfinite(<2 x half> %a)
declare i16 @llvm.nvvm.bf16x2.to{.e2m3x2, .e3m2x2}.rn{.relu}.satfinite(<2 x bfloat> %a)
declare <2 x half> @llvm.nvvm{.e2m3x2, .e3m2x2}.to.f16x2.rn{.relu}(i16 %a)
+ declare <2 x bfloat> @llvm.nvvm{.e2m3x2, .e3m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
declare <4 x i8> @llvm.nvvm.f32x4.to{.e2m3x4, .e3m2x4}.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
Overview:
@@ -1245,6 +1253,8 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
+For scale factor, see :ref:`scale-factor`.
+
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
``fp4`` Conversion Intrinsics
@@ -1259,6 +1269,7 @@ Syntax:
declare i16 @llvm.nvvm.f16x2.to.e2m1x2.rn{.relu}.satfinite(<2 x half> %a)
declare i16 @llvm.nvvm.bf16x2.to.e2m1x2.rn{.relu}.satfinite(<2 x bfloat> %a)
declare <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn{.relu}(i16 %a)
+ declare <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
declare i16 @llvm.nvvm.f32x4.to.e2m1x4.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
Overview:
@@ -1281,6 +1292,8 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
+For scale factor, see :ref:`scale-factor`.
+
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
Arithmetic Intrinsics
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 3fa219ee89325..a146791612da0 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1795,6 +1795,11 @@ let TargetPrefix = "nvvm" in {
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]>;
+ }
}
}
@@ -1820,6 +1825,11 @@ let TargetPrefix = "nvvm" in {
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]>;
+ }
}
// RS rounding mode (Stochastic Rounding) conversions for f4x4 type
@@ -1843,6 +1853,11 @@ let TargetPrefix = "nvvm" in {
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/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 64419d5cceefe..def709e2cf23e 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -118,6 +118,11 @@ void NVPTXInstPrinter::printCvtMode(const MCInst *MI, int OpNum,
if (Imm & NVPTX::PTXCvtMode::SAT_FLAG)
O << ".sat";
return;
+ } else if (Modifier == "satfinite") {
+ // SATFINITE flag
+ if (Imm & NVPTX::PTXCvtMode::SATFINITE_FLAG)
+ O << ".satfinite";
+ return;
} else if (Modifier == "relu") {
// RELU flag
if (Imm & NVPTX::PTXCvtMode::RELU_FLAG)
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 09a94034894eb..d751f95fe3db0 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -236,7 +236,8 @@ enum CvtMode {
BASE_MASK = 0x0F,
FTZ_FLAG = 0x10,
SAT_FLAG = 0x20,
- RELU_FLAG = 0x40
+ RELU_FLAG = 0x40,
+ SATFINITE_FLAG = 0x80
};
}
diff --git a/llvm/lib/Target/NVPTX/NVPTX.td b/llvm/lib/Target/NVPTX/NVPTX.td
index 17090dca0856d..1c169b05841c3 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.td
+++ b/llvm/lib/Target/NVPTX/NVPTX.td
@@ -106,7 +106,7 @@ foreach sm = [20, 21, 30, 32, 35, 37, 50, 52, 53, 60,
foreach version = [32, 40, 41, 42, 43, 50, 60, 61, 62, 63, 64, 65, 70, 71, 72,
73, 74, 75, 76, 77, 78, 80, 81, 82, 83, 84, 85, 86, 87, 88,
- 90, 91] in
+ 90, 91, 92] in
def PTX#version : FeaturePTX<version>;
def Is64Bit : Predicate<"Subtarget->getTargetTriple().getArch() == Triple::nvptx64">;
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 73ad1b979d8c8..100eaf2dd622c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -55,6 +55,10 @@ def CvtRN_RELU : PatLeaf<(i32 0x45)>;
def CvtRZ_RELU : PatLeaf<(i32 0x46)>;
def CvtRS_RELU : PatLeaf<(i32 0x4A)>;
+def CvtNONE_SATFINITE : PatLeaf<(i32 0x80)>;
+def CvtRN_SATFINITE : PatLeaf<(i32 0x85)>;
+def CvtRN_RELU_SATFINITE : PatLeaf<(i32 0xC5)>;
+
def CvtMode : Operand<i32> {
let PrintMethod = "printCvtMode";
}
@@ -666,6 +670,16 @@ let hasSideEffects = false in {
defm CVT_e4m3x2 : CVT_TO_F8X2<"e4m3">;
defm CVT_e5m2x2 : CVT_TO_F8X2<"e5m2">;
+
+ multiclass CVT_F8x2_TO_BF16x2<string F8Name> {
+ def _ # F8Name # x2_scale :
+ BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src, B16:$src2), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F8Name # "x2">;
+ }
+
+ defm CVT_bf16x2 : CVT_F8x2_TO_BF16x2<"e4m3">;
+ defm CVT_bf16x2 : CVT_F8x2_TO_BF16x2<"e5m2">;
class CVT_f16x2_fp8<string F8Name> :
BasicFlagsNVPTXInst<(outs B32:$dst),
@@ -729,6 +743,16 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
"cvt${mode:base}${mode:relu}.f16x2." # type>;
}
+ multiclass CVT_F6x2_TO_BF16x2<string F6Name> {
+ def _ # F6Name # x2_scale :
+ BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src, B16:$src2), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F6Name # "x2">;
+ }
+
+ defm CVT_bf16x2 : CVT_F6x2_TO_BF16x2<"e2m3">;
+ defm CVT_bf16x2 : CVT_F6x2_TO_BF16x2<"e3m2">;
+
class CVT_TO_FP6X4<string F6Name> :
NVPTXInst<(outs B32:$dst),
(ins B32:$src1, B32:$src2, B32:$src3, B32:$src4, B32:$src5, CvtMode:$mode),
@@ -772,6 +796,15 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
"cvt.u8.u16 \t%e2m1x2_in, $src; \n\t",
"cvt${mode:base}${mode:relu}.f16x2.e2m1x2 \t$dst, %e2m1x2_in; \n\t",
"}}"), []>;
+
+ def CVT_bf16x2_e2m1x2_scale :
+ NVPTXInst<(outs B32:$dst),
+ (ins B16:$src1, B16:$src2, CvtMode:$mode),
+ "{{ \n\t" #
+ ".reg .b8 \t%e2m1x2_in; \n\t" #
+ "cvt.u8.u16 \t%e2m1x2_in, $src1; \n\t" #
+ "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2.e2m1x2 \t$dst, %e2m1x2_in, $src2; \n\t" #
+ "}}", []>;
def CVT_e2m1x4_f32x4_rs_sf :
NVPTXInst<(outs B16:$dst),
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cde6b2c8760f4..160d8c3ca2bb3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2262,6 +2262,21 @@ let Predicates = [callSubtarget<"hasFP8ConversionSupport">] in {
(CVT_f16x2_e5m2x2 $a, CvtRN_RELU)>;
}
+let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
+ foreach src_type = ["e4m3x2", "e5m2x2"] in {
+ foreach relu = ["", "_relu"] in {
+ foreach satfinite = ["", "_satfinite"] in {
+ defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
+ defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
+ defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
+
+ def : Pat<(intrin B16:$a, B16:$b),
+ (inst $a, $b, mode)>;
+ }
+ }
+ }
+} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
+
let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
def : Pat<(int_nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0 f32:$a, f32:$b, i16:$scale),
(CVT_s2f6x2_f32_sf_scale $a, $b, $scale, CvtRN)>;
@@ -2317,6 +2332,20 @@ let Predicates = [callSubtarget<"hasFP16X2ToNarrowFPConversionSupport">] in {
}
}
+let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
+ foreach src_type = ["e2m3x2", "e3m2x2"] in {
+ foreach relu = ["", "_relu"] in {
+ foreach satfinite = ["", "_satfinite"] in {
+ defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
+ defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
+ defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
+ def : Pat<(intrin B16:$a, B16:$b),
+ (inst $a, $b, mode)>;
+ }
+ }
+ }
+} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
+
let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
def : Pat<(int_nvvm_ff_to_e2m1x2_rn_satfinite f32:$a, f32:$b),
(CVT_e2m1x2_f32_sf $a, $b, CvtRN)>;
@@ -2340,6 +2369,18 @@ let Predicates = [callSubtarget<"hasFP16X2ToNarrowFPConversionSupport">] in {
}
}
+let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
+ foreach relu = ["", "_relu"] in {
+ foreach satfinite = ["", "_satfinite"] in {
+ defvar intrin = !cast<Intrinsic>("int_nvvm_e2m1x2_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
+ defvar inst = !cast<NVPTXInst>("CVT_bf16x2_e2m1x2_scale");
+ defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
+ def : Pat<(intrin B16:$a, B16:$b),
+ (inst $a, $b, mode)>;
+ }
+ }
+} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
+
let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
def : Pat<(int_nvvm_ff_to_ue8m0x2_rz f32:$a, f32:$b),
(CVT_ue8m0x2_f32 $a, $b, CvtRZ)>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 8ff34284e6b16..4398a4b3fd628 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -240,6 +240,11 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
bool hasS2F6X2ConversionSupport() const {
return hasPTXWithAccelSMs(91, {100, 103, 110, 120, 121});
}
+
+ // Checks support for conversions from narrow FP types to bf16x2.
+ bool hasNarrowFPToBF16x2ConversionSupport() const {
+ return hasPTXWithFamilySMs(92, {100, 110, 120});
+ }
bool hasTensormapReplaceSupport() const {
return hasPTXWithFamilySMs(90, {90, 100, 110, 120}) ||
diff --git a/llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll b/llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll
new file mode 100644
index 0000000000000..3560ef5b57936
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll
@@ -0,0 +1,87 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
+; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
+
+define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_in;
+; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
+; CHECK-NEXT: }
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_in;
+; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
+; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
+; CHECK-NEXT: }
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_in;
+; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
+; CHECK-NEXT: }
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b8 %e2m1x2_in;
+; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
+; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
+; CHECK-NEXT: }
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll b/llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll
new file mode 100644
index 0000000000000..56476252641df
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll
@@ -0,0 +1,135 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
+; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
+
+define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll b/llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll
new file mode 100644
index 0000000000000..1199df2fb5697
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll
@@ -0,0 +1,135 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
+; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
+; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
+; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
+
+define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
>From 9b80bf9d3f64c6cc11f441c9a68d4e7e29874eec Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Fri, 10 Apr 2026 10:09:03 +0000
Subject: [PATCH 2/6] fix formatting
---
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 4398a4b3fd628..1df5d326f63a6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -240,7 +240,7 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
bool hasS2F6X2ConversionSupport() const {
return hasPTXWithAccelSMs(91, {100, 103, 110, 120, 121});
}
-
+
// Checks support for conversions from narrow FP types to bf16x2.
bool hasNarrowFPToBF16x2ConversionSupport() const {
return hasPTXWithFamilySMs(92, {100, 110, 120});
>From 08210d5e69579c9f5ab5c8813cef7395d422dce9 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 13 Apr 2026 08:53:31 +0000
Subject: [PATCH 3/6] fix refs in docs
---
llvm/docs/NVPTXUsage.rst | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 68144a4669c5e..72eaafce38482 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -1181,7 +1181,7 @@ specified destination format. The ``satfinite`` modifier is assumed to be
present for conversions involving ``e4m3`` and ``e5m2`` types as the
destination.
-For scale factor, see :ref:`scale-factor`.
+For scale factor, see :ref:`scale-factor <scale-factor>`.
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
@@ -1215,7 +1215,7 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
-For scale factor, see :ref:`scale-factor`.
+For scale factor, see :ref:`scale-factor <scale-factor>`.
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
@@ -1253,7 +1253,7 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
-For scale factor, see :ref:`scale-factor`.
+For scale factor, see :ref:`scale-factor <scale-factor>`.
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
@@ -1292,7 +1292,7 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
destination format.
-For scale factor, see :ref:`scale-factor`.
+For scale factor, see :ref:`scale-factor <scale-factor>`.
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
>From 0c3852b29c10f2c2b06ede1c84b58962e6b52126 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 13 Apr 2026 09:22:45 +0000
Subject: [PATCH 4/6] address comments and add predicates
---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 50 ++++++++++++-------------
1 file changed, 24 insertions(+), 26 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 100eaf2dd622c..c7e921b4a3bc2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -671,24 +671,21 @@ let hasSideEffects = false in {
defm CVT_e4m3x2 : CVT_TO_F8X2<"e4m3">;
defm CVT_e5m2x2 : CVT_TO_F8X2<"e5m2">;
- multiclass CVT_F8x2_TO_BF16x2<string F8Name> {
- def _ # F8Name # x2_scale :
+ multiclass CVT_F8x2_TO_FP16x2<string F8Name> {
+ def f16x2_ # F8Name # x2 :
+ BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:relu}.f16x2." # F8Name # "x2">,
+ Requires<[hasPTX<81>, hasSM<89>]>;
+ def bf16x2_ # F8Name # x2_scale :
BasicFlagsNVPTXInst<(outs B32:$dst),
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F8Name # "x2">;
+ "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F8Name # "x2">,
+ Requires<[callSubTarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
}
- defm CVT_bf16x2 : CVT_F8x2_TO_BF16x2<"e4m3">;
- defm CVT_bf16x2 : CVT_F8x2_TO_BF16x2<"e5m2">;
-
- class CVT_f16x2_fp8<string F8Name> :
- BasicFlagsNVPTXInst<(outs B32:$dst),
- (ins B16:$src), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.f16x2." # F8Name # "x2">,
- Requires<[hasPTX<81>, hasSM<89>]>;
-
- def CVT_f16x2_e4m3x2 : CVT_f16x2_fp8<"e4m3">;
- def CVT_f16x2_e5m2x2 : CVT_f16x2_fp8<"e5m2">;
+ defm CVT_ : CVT_F8x2_TO_FP16x2<"e4m3">;
+ defm CVT_ : CVT_F8x2_TO_FP16x2<"e5m2">;
class CVT_TO_FP8X4<string F8Name> :
NVPTXInst<(outs B32:$dst),
@@ -737,21 +734,21 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
}
// FP6 conversions.
- foreach type = ["e2m3x2", "e3m2x2"] in {
- def CVT_f16x2_ # type : BasicFlagsNVPTXInst<(outs B32:$dst),
- (ins B16:$src), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.f16x2." # type>;
- }
-
- multiclass CVT_F6x2_TO_BF16x2<string F6Name> {
- def _ # F6Name # x2_scale :
+ multiclass CVT_F6x2_TO_FP16x2<string F6Name> {
+ def f16x2_ # F6Name # x2 :
+ BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:relu}.f16x2." # F6Name # "x2">,
+ Requires<[callSubTarget<"hasNarrowFPConversionSupport">]>;
+ def bf16x2_ # F6Name # x2_scale :
BasicFlagsNVPTXInst<(outs B32:$dst),
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F6Name # "x2">;
+ "cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F6Name # "x2">,
+ Requires<[callSubTarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
}
- defm CVT_bf16x2 : CVT_F6x2_TO_BF16x2<"e2m3">;
- defm CVT_bf16x2 : CVT_F6x2_TO_BF16x2<"e3m2">;
+ defm CVT_ : CVT_F6x2_TO_FP16x2<"e2m3">;
+ defm CVT_ : CVT_F6x2_TO_FP16x2<"e3m2">;
class CVT_TO_FP6X4<string F6Name> :
NVPTXInst<(outs B32:$dst),
@@ -804,7 +801,8 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
".reg .b8 \t%e2m1x2_in; \n\t" #
"cvt.u8.u16 \t%e2m1x2_in, $src1; \n\t" #
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2.e2m1x2 \t$dst, %e2m1x2_in, $src2; \n\t" #
- "}}", []>;
+ "}}", []>,
+ Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
def CVT_e2m1x4_f32x4_rs_sf :
NVPTXInst<(outs B16:$dst),
>From 416286822991b14e52d1b0db721b758a03f9ba5c Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 13 Apr 2026 09:48:08 +0000
Subject: [PATCH 5/6] fix errors
---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index c7e921b4a3bc2..4e978c6ab030e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -681,7 +681,7 @@ let hasSideEffects = false in {
BasicFlagsNVPTXInst<(outs B32:$dst),
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F8Name # "x2">,
- Requires<[callSubTarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
+ Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
}
defm CVT_ : CVT_F8x2_TO_FP16x2<"e4m3">;
@@ -739,12 +739,12 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
BasicFlagsNVPTXInst<(outs B32:$dst),
(ins B16:$src), (ins CvtMode:$mode),
"cvt${mode:base}${mode:relu}.f16x2." # F6Name # "x2">,
- Requires<[callSubTarget<"hasNarrowFPConversionSupport">]>;
+ Requires<[callSubtarget<"hasNarrowFPConversionSupport">]>;
def bf16x2_ # F6Name # x2_scale :
BasicFlagsNVPTXInst<(outs B32:$dst),
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F6Name # "x2">,
- Requires<[callSubTarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
+ Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
}
defm CVT_ : CVT_F6x2_TO_FP16x2<"e2m3">;
>From d48bd5a83d072902da4ab5e9f27d66b343f8b589 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Mon, 13 Apr 2026 14:03:46 +0000
Subject: [PATCH 6/6] replace B16 with i16 in patterns
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 160d8c3ca2bb3..7acc9b9cda34a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2270,7 +2270,7 @@ let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
- def : Pat<(intrin B16:$a, B16:$b),
+ def : Pat<(intrin i16:$a, i16:$b),
(inst $a, $b, mode)>;
}
}
@@ -2339,7 +2339,7 @@ let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
- def : Pat<(intrin B16:$a, B16:$b),
+ def : Pat<(intrin i16:$a, i16:$b),
(inst $a, $b, mode)>;
}
}
@@ -2375,7 +2375,7 @@ let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
defvar intrin = !cast<Intrinsic>("int_nvvm_e2m1x2_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_e2m1x2_scale");
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
- def : Pat<(intrin B16:$a, B16:$b),
+ def : Pat<(intrin i16:$a, i16:$b),
(inst $a, $b, mode)>;
}
}
More information about the llvm-commits
mailing list