[llvm] [NVPTX] Add intrinsics for ue5m3x2 to f16x2/bf16x2 conversions (PR #214411)
Dharuni R Acharya via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 6 00:02:05 PDT 2026
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/214411
>From 46c8480cb75735396a8caafc05314858b8a1fdc6 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 6 Aug 2026 06:37:54 +0000
Subject: [PATCH] [NVPTX] Add intrinsics for ue5m3x2 to f16x2/bf16x2
conversions
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>
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 12 ++++
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 14 ++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 16 +++++
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 3 +
llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll | 80 ++++++++++++++++++++++
5 files changed, 125 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 40f3c40a76bc6..8dce8203ce632 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 6d9ef347d616d..52c8e2294e183 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -869,6 +869,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 5230e01261fd6..bf3ca9a660a0b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2404,6 +2404,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 22affe4f40759..884015fd5048e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -246,6 +246,9 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
return hasPTXWithFamilySMs(92, {100, 110, 120});
}
+ // Checks support for conversions involving ue5m3x2.
+ bool hasUE5M3TypeSupport() const { return hasPTXWithFamilySMs(94, {107}); }
+
bool hasTensormapReplaceSupport() const {
return hasPTXWithFamilySMs(90, {90, 100, 110, 120}) ||
hasPTXWithFamilySMs(88, {90, 100, 101, 120}) ||
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