[llvm] [NVPTX] Add intrinsics for ue5m3x2 to f16x2/bf16x2 conversions (PR #214411)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 5 23:51:44 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Dharuni R Acharya (DharuniRAcharya)
<details>
<summary>Changes</summary>
This patch adds 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`.
PTX ISA Reference: https://docs.nvidia.com/cuda/developer-preview/13.4/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt
TODO: The remaining intrinsics in this conversion family, along with the corresponding documentation, will be added in future patches.
---
Full diff: https://github.com/llvm/llvm-project/pull/214411.diff
5 Files Affected:
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+12)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+14)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+16)
- (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.h (+5)
- (added) llvm/test/CodeGen/NVPTX/convert-ue5m3x2.ll (+80)
``````````diff
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..b4828fe4614ca 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -246,6 +246,11 @@ 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
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/214411
More information about the llvm-commits
mailing list