[llvm] df90dfb - [NVPTX] Add intrinsics for ue5m3x2 to f16x2/bf16x2 conversions (#214411)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 10 03:06:43 PDT 2026
Author: Dharuni R Acharya
Date: 2026-08-10T15:36:38+05:30
New Revision: df90dfb7a2652b85e121b94578a6195b9a74e93f
URL: https://github.com/llvm/llvm-project/commit/df90dfb7a2652b85e121b94578a6195b9a74e93f
DIFF: https://github.com/llvm/llvm-project/commit/df90dfb7a2652b85e121b94578a6195b9a74e93f.diff
LOG: [NVPTX] Add intrinsics for ue5m3x2 to f16x2/bf16x2 conversions (#214411)
This patch adds the following intrinsics for `ue5m3x2` to `f16x2/bf16x2`
conversions introduced in PTX 9.4:
- `cvt.rn.f16x2.ue5m3x2`
- `cvt.rn{.satfinite}{.scaled::n2::ue8m0}.bf16x2.ue5m3x2`
Tests have been verified through `ptxas-13.4`.
---------
Signed-off-by: DharuniRAcharya <dharunira at nvidia.com>
Added:
llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll
Modified:
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/lib/Target/NVPTX/NVPTXSubtarget.h
Removed:
################################################################################
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index a892916454323..d1bca807d897d 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1879,6 +1879,18 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_ue8m0x2_to_bf16x2 : NVVMBuiltin,
PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty]>;
+ // UE5M3x2 conversions.
+ def int_nvvm_ue5m3x2_to_f16x2_rn
+ : PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
+
+ foreach satfinite = ["", "_satfinite"] in {
+ def int_nvvm_ue5m3x2_to_bf16x2_rn # satfinite
+ : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty]>;
+
+ def int_nvvm_ue5m3x2_to_bf16x2_rn # satfinite # _scale_n2_ue8m0
+ : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
+ }
+
//
// Bar.Sync
//
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 83060bb906157..0df5831cce038 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -841,6 +841,20 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
(ins B16:$src),
"cvt.rn.bf16x2.ue8m0x2">;
+ // UE5M3x2 conversions.
+ def CVT_f16x2_ue5m3x2 : BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src), (ins CvtMode:$mode),
+ "cvt${mode:base}.f16x2.ue5m3x2">,
+ Requires<[callSubtarget<"hasUE5M3TypeSupport">]>;
+ def CVT_bf16x2_ue5m3x2 : BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:satfinite}.bf16x2.ue5m3x2">,
+ Requires<[callSubtarget<"hasUE5M3TypeSupport">]>;
+ def CVT_bf16x2_ue5m3x2_scale_n2 : BasicFlagsNVPTXInst<(outs B32:$dst),
+ (ins B16:$src, B16:$src2), (ins CvtMode:$mode),
+ "cvt${mode:base}${mode:satfinite}.scaled::n2::ue8m0.bf16x2.ue5m3x2">,
+ Requires<[callSubtarget<"hasUE5M3TypeSupport">]>;
+
}
def fpround_oneuse : OneUse1<fpround>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index e946f6b6045d0..602bb4faecb9d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2438,6 +2438,22 @@ let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
(CVT_bf16x2_ue8m0x2 $a)>;
}
+// ue5m3x2 to f16x2 / bf16x2 conversions.
+let Predicates = [callSubtarget<"hasUE5M3TypeSupport">] in {
+ def : Pat<(int_nvvm_ue5m3x2_to_f16x2_rn i16:$a),
+ (CVT_f16x2_ue5m3x2 $a, CvtRN)>;
+
+ def : Pat<(int_nvvm_ue5m3x2_to_bf16x2_rn i16:$a),
+ (CVT_bf16x2_ue5m3x2 $a, CvtRN)>;
+ def : Pat<(int_nvvm_ue5m3x2_to_bf16x2_rn_satfinite i16:$a),
+ (CVT_bf16x2_ue5m3x2 $a, CvtRN_SATFINITE)>;
+
+ def : Pat<(int_nvvm_ue5m3x2_to_bf16x2_rn_scale_n2_ue8m0 i16:$a, i16:$b),
+ (CVT_bf16x2_ue5m3x2_scale_n2 $a, $b, CvtRN)>;
+ def : Pat<(int_nvvm_ue5m3x2_to_bf16x2_rn_satfinite_scale_n2_ue8m0 i16:$a, i16:$b),
+ (CVT_bf16x2_ue5m3x2_scale_n2 $a, $b, CvtRN_SATFINITE)>;
+} // let Predicates = [callSubtarget<"hasUE5M3TypeSupport">]
+
def SDT_CVT_F32X4_TO_FPX4_RS_VEC :
SDTypeProfile<1, 6, [SDTCisVec<0>, SDTCisFP<1>, SDTCisFP<2>, SDTCisFP<3>,
SDTCisFP<4>, SDTCisInt<5>, SDTCisInt<6>]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 2057e80a9c5a2..65107d57dae2b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -256,6 +256,11 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
hasAnyFeature({NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f});
}
+ // Checks support for conversions involving ue5m3x2.
+ bool hasUE5M3TypeSupport() const {
+ return PTXVersion >= 94 && hasAnyFeature({NVPTX::SM107f});
+ }
+
bool hasTensormapReplaceSupport() const {
return hasAnyFeature({NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f}) ||
(PTXVersion >= 83 && hasAnyFeature({NVPTX::SM90a}));
diff --git a/llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll b/llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll
new file mode 100644
index 0000000000000..fe544aee19afc
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll
@@ -0,0 +1,80 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | FileCheck %s
+; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %}
+
+define <2 x half> @test_ue5m3x2_to_f16x2_rn(i16 %a) {
+; CHECK-LABEL: test_ue5m3x2_to_f16x2_rn(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_ue5m3x2_to_f16x2_rn_param_0];
+; CHECK-NEXT: cvt.rn.f16x2.ue5m3x2 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ue5m3x2.to.f16x2.rn(i16 %a)
+ ret <2 x half> %val
+}
+
+define <2 x bfloat> @test_ue5m3x2_to_bf16x2_rn(i16 %a) {
+; CHECK-LABEL: test_ue5m3x2_to_bf16x2_rn(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_ue5m3x2_to_bf16x2_rn_param_0];
+; CHECK-NEXT: cvt.rn.bf16x2.ue5m3x2 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ue5m3x2.to.bf16x2.rn(i16 %a)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_ue5m3x2_to_bf16x2_rn_satfinite(i16 %a) {
+; CHECK-LABEL: test_ue5m3x2_to_bf16x2_rn_satfinite(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [test_ue5m3x2_to_bf16x2_rn_satfinite_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.bf16x2.ue5m3x2 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ue5m3x2.to.bf16x2.rn.satfinite(i16 %a)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_ue5m3x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_ue5m3x2_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_ue5m3x2_to_bf16x2_rn_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_ue5m3x2_to_bf16x2_rn_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.ue5m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ue5m3x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @test_ue5m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
+; CHECK-LABEL: test_ue5m3x2_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_ue5m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [test_ue5m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.ue5m3x2 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ue5m3x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
+ ret <2 x bfloat> %val
+}
More information about the llvm-commits
mailing list