https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/227652
>From 9d2aab4fc4d02b158d09b1d18bd8a3f37d97eb37 Mon Sep 17 00:00:00 2001 From: DharuniRAcharya <[email protected]> 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 <[email protected]> --- 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 <[email protected]> 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 <[email protected]> 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 { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
