[NVPTX] Add intrinsics for narrow-fp to bf16 conversions (#191376)
Adds the following intrinsics for narrow-fp to bf16 conversions
introduced in PTX 9.2:
-
llvm.nvvm.{e4m3x2/e5m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
-
llvm.nvvm.{e2m3x2/e3m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
- llvm.nvvm.e2m1x2.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0
Tests have been verified through `ptxas-13.2`.
PTX ISA Reference:
https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt
This commit is contained in:
Srinivasa Ravi
GitHub
parent
cd14598958
commit
83b4c5cd71
11 changed files with 487 additions and 19 deletions
|
|
@ -1134,6 +1134,14 @@ The following table describes the rounding modes used across these intrinsics:
|
|||
| | the input. |
|
||||
+-----------------------+---------------------------------------------------+
|
||||
|
||||
.. _scale-factor:
|
||||
|
||||
Some conversions involve a scale factor which is provided as a packed 16-bit
|
||||
integer containing two scaling factors of type ``ue8m0``, one for each input.
|
||||
For down conversion, inputs are divided by ``scale_factor`` and then the
|
||||
conversion is performed. For up-conversion, inputs are converted to destination
|
||||
type and then multiplied by ``scale_factor``.
|
||||
|
||||
``fp8`` Conversion Intrinsics
|
||||
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
|
||||
|
||||
|
|
@ -1148,6 +1156,7 @@ Syntax:
|
|||
declare i16 @llvm.bf16x2.to{.e4m3x2, .e5m2x2}.rn{.relu}.satfinite(<2 x bfloat> %a)
|
||||
declare i16 @llvm.bf16x2.to.ue8m0x2{.rz, .rp}{.satfinite}(<2 x bfloat> %a)
|
||||
declare <2 x half> @llvm.nvvm{.e4m3x2, .e5m2x2}.to.f16x2.rn{.relu}(i16 %a)
|
||||
declare <2 x bfloat> @llvm.nvvm{.e4m3x2, .e5m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
|
||||
declare <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 %a)
|
||||
declare <4 x i8> @llvm.nvvm.f32x4.to{.e4m3x4, .e5m2x4}.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
|
||||
|
||||
|
|
@ -1172,6 +1181,8 @@ specified destination format. The ``satfinite`` modifier is assumed to be
|
|||
present for conversions involving ``e4m3`` and ``e5m2`` types as the
|
||||
destination.
|
||||
|
||||
For scale factor, see :ref:`scale-factor <scale-factor>`.
|
||||
|
||||
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
|
||||
|
||||
``s2f6`` Conversion Intrinsics
|
||||
|
|
@ -1204,11 +1215,7 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
|
|||
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
|
||||
destination format.
|
||||
|
||||
The operand ``%scale_factor`` stores two packed scaling factors of type
|
||||
``ue8m0``, one for each input. For down conversion, inputs are divided by
|
||||
``scale_factor`` and then the conversion is performed. For up-conversion,
|
||||
inputs are converted to destination type and then multiplied by
|
||||
``scale_factor``.
|
||||
For scale factor, see :ref:`scale-factor <scale-factor>`.
|
||||
|
||||
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
|
||||
|
||||
|
|
@ -1224,6 +1231,7 @@ Syntax:
|
|||
declare i16 @llvm.nvvm.f16x2.to{.e2m3x2, .e3m2x2}.rn{.relu}.satfinite(<2 x half> %a)
|
||||
declare i16 @llvm.nvvm.bf16x2.to{.e2m3x2, .e3m2x2}.rn{.relu}.satfinite(<2 x bfloat> %a)
|
||||
declare <2 x half> @llvm.nvvm{.e2m3x2, .e3m2x2}.to.f16x2.rn{.relu}(i16 %a)
|
||||
declare <2 x bfloat> @llvm.nvvm{.e2m3x2, .e3m2x2}.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
|
||||
declare <4 x i8> @llvm.nvvm.f32x4.to{.e2m3x4, .e3m2x4}.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
|
||||
|
||||
Overview:
|
||||
|
|
@ -1245,6 +1253,8 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
|
|||
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
|
||||
destination format.
|
||||
|
||||
For scale factor, see :ref:`scale-factor <scale-factor>`.
|
||||
|
||||
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
|
||||
|
||||
``fp4`` Conversion Intrinsics
|
||||
|
|
@ -1259,6 +1269,7 @@ Syntax:
|
|||
declare i16 @llvm.nvvm.f16x2.to.e2m1x2.rn{.relu}.satfinite(<2 x half> %a)
|
||||
declare i16 @llvm.nvvm.bf16x2.to.e2m1x2.rn{.relu}.satfinite(<2 x bfloat> %a)
|
||||
declare <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn{.relu}(i16 %a)
|
||||
declare <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn{.relu}{.satfinite}.scale.n2.ue8m0(i16 %a, i16 %scale_factor)
|
||||
declare i16 @llvm.nvvm.f32x4.to.e2m1x4.rs{.relu}.satfinite(<4 x f32> %a, i32 %rnd_bits)
|
||||
|
||||
Overview:
|
||||
|
|
@ -1281,6 +1292,8 @@ result is sign-preserved ``MAX_NORM`` of the destination format. Also, if the
|
|||
input is ``NaN``, then the result is the positive ``MAX_NORM`` of the
|
||||
destination format.
|
||||
|
||||
For scale factor, see :ref:`scale-factor <scale-factor>`.
|
||||
|
||||
For more information, see `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt>`__.
|
||||
|
||||
Arithmetic Intrinsics
|
||||
|
|
|
|||
|
|
@ -1795,6 +1795,11 @@ let TargetPrefix = "nvvm" in {
|
|||
|
||||
def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite
|
||||
: PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
|
||||
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
|
||||
: PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -1820,6 +1825,11 @@ let TargetPrefix = "nvvm" in {
|
|||
|
||||
def int_nvvm_bf16x2_to_e2m1x2_rn # relu # _satfinite
|
||||
: PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
|
||||
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
def int_nvvm_e2m1x2_to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
|
||||
: PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
|
||||
}
|
||||
}
|
||||
|
||||
// RS rounding mode (Stochastic Rounding) conversions for f4x4 type
|
||||
|
|
@ -1843,6 +1853,11 @@ let TargetPrefix = "nvvm" in {
|
|||
|
||||
def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite
|
||||
: PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
|
||||
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
|
||||
: PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -118,6 +118,11 @@ void NVPTXInstPrinter::printCvtMode(const MCInst *MI, int OpNum,
|
|||
if (Imm & NVPTX::PTXCvtMode::SAT_FLAG)
|
||||
O << ".sat";
|
||||
return;
|
||||
} else if (Modifier == "satfinite") {
|
||||
// SATFINITE flag
|
||||
if (Imm & NVPTX::PTXCvtMode::SATFINITE_FLAG)
|
||||
O << ".satfinite";
|
||||
return;
|
||||
} else if (Modifier == "relu") {
|
||||
// RELU flag
|
||||
if (Imm & NVPTX::PTXCvtMode::RELU_FLAG)
|
||||
|
|
|
|||
|
|
@ -236,7 +236,8 @@ enum CvtMode {
|
|||
BASE_MASK = 0x0F,
|
||||
FTZ_FLAG = 0x10,
|
||||
SAT_FLAG = 0x20,
|
||||
RELU_FLAG = 0x40
|
||||
RELU_FLAG = 0x40,
|
||||
SATFINITE_FLAG = 0x80
|
||||
};
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -106,7 +106,7 @@ foreach sm = [20, 21, 30, 32, 35, 37, 50, 52, 53, 60,
|
|||
|
||||
foreach version = [32, 40, 41, 42, 43, 50, 60, 61, 62, 63, 64, 65, 70, 71, 72,
|
||||
73, 74, 75, 76, 77, 78, 80, 81, 82, 83, 84, 85, 86, 87, 88,
|
||||
90, 91] in
|
||||
90, 91, 92] in
|
||||
def PTX#version : FeaturePTX<version>;
|
||||
|
||||
def Is64Bit : Predicate<"Subtarget->getTargetTriple().getArch() == Triple::nvptx64">;
|
||||
|
|
|
|||
|
|
@ -55,6 +55,10 @@ def CvtRN_RELU : PatLeaf<(i32 0x45)>;
|
|||
def CvtRZ_RELU : PatLeaf<(i32 0x46)>;
|
||||
def CvtRS_RELU : PatLeaf<(i32 0x4A)>;
|
||||
|
||||
def CvtNONE_SATFINITE : PatLeaf<(i32 0x80)>;
|
||||
def CvtRN_SATFINITE : PatLeaf<(i32 0x85)>;
|
||||
def CvtRN_RELU_SATFINITE : PatLeaf<(i32 0xC5)>;
|
||||
|
||||
def CvtMode : Operand<i32> {
|
||||
let PrintMethod = "printCvtMode";
|
||||
}
|
||||
|
|
@ -666,15 +670,22 @@ let hasSideEffects = false in {
|
|||
|
||||
defm CVT_e4m3x2 : CVT_TO_F8X2<"e4m3">;
|
||||
defm CVT_e5m2x2 : CVT_TO_F8X2<"e5m2">;
|
||||
|
||||
multiclass CVT_F8x2_TO_FP16x2<string F8Name> {
|
||||
def f16x2_ # F8Name # x2 :
|
||||
BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}.f16x2." # F8Name # "x2">,
|
||||
Requires<[hasPTX<81>, hasSM<89>]>;
|
||||
def bf16x2_ # F8Name # x2_scale :
|
||||
BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F8Name # "x2">,
|
||||
Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
|
||||
}
|
||||
|
||||
class CVT_f16x2_fp8<string F8Name> :
|
||||
BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}.f16x2." # F8Name # "x2">,
|
||||
Requires<[hasPTX<81>, hasSM<89>]>;
|
||||
|
||||
def CVT_f16x2_e4m3x2 : CVT_f16x2_fp8<"e4m3">;
|
||||
def CVT_f16x2_e5m2x2 : CVT_f16x2_fp8<"e5m2">;
|
||||
defm CVT_ : CVT_F8x2_TO_FP16x2<"e4m3">;
|
||||
defm CVT_ : CVT_F8x2_TO_FP16x2<"e5m2">;
|
||||
|
||||
class CVT_TO_FP8X4<string F8Name> :
|
||||
NVPTXInst<(outs B32:$dst),
|
||||
|
|
@ -723,11 +734,21 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
|
|||
}
|
||||
|
||||
// FP6 conversions.
|
||||
foreach type = ["e2m3x2", "e3m2x2"] in {
|
||||
def CVT_f16x2_ # type : BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}.f16x2." # type>;
|
||||
multiclass CVT_F6x2_TO_FP16x2<string F6Name> {
|
||||
def f16x2_ # F6Name # x2 :
|
||||
BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}.f16x2." # F6Name # "x2">,
|
||||
Requires<[callSubtarget<"hasNarrowFPConversionSupport">]>;
|
||||
def bf16x2_ # F6Name # x2_scale :
|
||||
BasicFlagsNVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src, B16:$src2), (ins CvtMode:$mode),
|
||||
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2." # F6Name # "x2">,
|
||||
Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
|
||||
}
|
||||
|
||||
defm CVT_ : CVT_F6x2_TO_FP16x2<"e2m3">;
|
||||
defm CVT_ : CVT_F6x2_TO_FP16x2<"e3m2">;
|
||||
|
||||
class CVT_TO_FP6X4<string F6Name> :
|
||||
NVPTXInst<(outs B32:$dst),
|
||||
|
|
@ -772,6 +793,16 @@ let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
|
|||
"cvt.u8.u16 \t%e2m1x2_in, $src; \n\t",
|
||||
"cvt${mode:base}${mode:relu}.f16x2.e2m1x2 \t$dst, %e2m1x2_in; \n\t",
|
||||
"}}"), []>;
|
||||
|
||||
def CVT_bf16x2_e2m1x2_scale :
|
||||
NVPTXInst<(outs B32:$dst),
|
||||
(ins B16:$src1, B16:$src2, CvtMode:$mode),
|
||||
"{{ \n\t" #
|
||||
".reg .b8 \t%e2m1x2_in; \n\t" #
|
||||
"cvt.u8.u16 \t%e2m1x2_in, $src1; \n\t" #
|
||||
"cvt${mode:base}${mode:relu}${mode:satfinite}.scaled::n2::ue8m0.bf16x2.e2m1x2 \t$dst, %e2m1x2_in, $src2; \n\t" #
|
||||
"}}", []>,
|
||||
Requires<[callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]>;
|
||||
|
||||
def CVT_e2m1x4_f32x4_rs_sf :
|
||||
NVPTXInst<(outs B16:$dst),
|
||||
|
|
|
|||
|
|
@ -2262,6 +2262,21 @@ let Predicates = [callSubtarget<"hasFP8ConversionSupport">] in {
|
|||
(CVT_f16x2_e5m2x2 $a, CvtRN_RELU)>;
|
||||
}
|
||||
|
||||
let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
|
||||
foreach src_type = ["e4m3x2", "e5m2x2"] in {
|
||||
foreach relu = ["", "_relu"] in {
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
|
||||
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
|
||||
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
|
||||
|
||||
def : Pat<(intrin i16:$a, i16:$b),
|
||||
(inst $a, $b, mode)>;
|
||||
}
|
||||
}
|
||||
}
|
||||
} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
|
||||
|
||||
let Predicates = [callSubtarget<"hasS2F6X2ConversionSupport">] in {
|
||||
def : Pat<(int_nvvm_ff_to_s2f6x2_rn_satfinite_scale_n2_ue8m0 f32:$a, f32:$b, i16:$scale),
|
||||
(CVT_s2f6x2_f32_sf_scale $a, $b, $scale, CvtRN)>;
|
||||
|
|
@ -2317,6 +2332,20 @@ let Predicates = [callSubtarget<"hasFP16X2ToNarrowFPConversionSupport">] in {
|
|||
}
|
||||
}
|
||||
|
||||
let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
|
||||
foreach src_type = ["e2m3x2", "e3m2x2"] in {
|
||||
foreach relu = ["", "_relu"] in {
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
|
||||
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_" # src_type # "_scale");
|
||||
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
|
||||
def : Pat<(intrin i16:$a, i16:$b),
|
||||
(inst $a, $b, mode)>;
|
||||
}
|
||||
}
|
||||
}
|
||||
} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
|
||||
|
||||
let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
|
||||
def : Pat<(int_nvvm_ff_to_e2m1x2_rn_satfinite f32:$a, f32:$b),
|
||||
(CVT_e2m1x2_f32_sf $a, $b, CvtRN)>;
|
||||
|
|
@ -2340,6 +2369,18 @@ let Predicates = [callSubtarget<"hasFP16X2ToNarrowFPConversionSupport">] in {
|
|||
}
|
||||
}
|
||||
|
||||
let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">] in {
|
||||
foreach relu = ["", "_relu"] in {
|
||||
foreach satfinite = ["", "_satfinite"] in {
|
||||
defvar intrin = !cast<Intrinsic>("int_nvvm_e2m1x2_to_bf16x2_rn" # relu # satfinite # "_scale_n2_ue8m0");
|
||||
defvar inst = !cast<NVPTXInst>("CVT_bf16x2_e2m1x2_scale");
|
||||
defvar mode = !cast<PatLeaf>("CvtRN" # !toupper(relu) # !toupper(satfinite));
|
||||
def : Pat<(intrin i16:$a, i16:$b),
|
||||
(inst $a, $b, mode)>;
|
||||
}
|
||||
}
|
||||
} // let Predicates = [callSubtarget<"hasNarrowFPToBF16x2ConversionSupport">]
|
||||
|
||||
let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
|
||||
def : Pat<(int_nvvm_ff_to_ue8m0x2_rz f32:$a, f32:$b),
|
||||
(CVT_ue8m0x2_f32 $a, $b, CvtRZ)>;
|
||||
|
|
|
|||
|
|
@ -241,6 +241,11 @@ public:
|
|||
return hasPTXWithAccelSMs(91, {100, 103, 110, 120, 121});
|
||||
}
|
||||
|
||||
// Checks support for conversions from narrow FP types to bf16x2.
|
||||
bool hasNarrowFPToBF16x2ConversionSupport() const {
|
||||
return hasPTXWithFamilySMs(92, {100, 110, 120});
|
||||
}
|
||||
|
||||
bool hasTensormapReplaceSupport() const {
|
||||
return hasPTXWithFamilySMs(90, {90, 100, 110, 120}) ||
|
||||
hasPTXWithFamilySMs(88, {90, 100, 101, 120}) ||
|
||||
|
|
|
|||
87
llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll
Normal file
87
llvm/test/CodeGen/NVPTX/convert_fp4x2_to_bf16x2.ll
Normal file
|
|
@ -0,0 +1,87 @@
|
|||
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
|
||||
; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
|
||||
; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
|
||||
|
||||
define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: {
|
||||
; CHECK-NEXT: .reg .b8 %e2m1x2_in;
|
||||
; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
|
||||
; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
|
||||
; CHECK-NEXT: }
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: {
|
||||
; CHECK-NEXT: .reg .b8 %e2m1x2_in;
|
||||
; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
|
||||
; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
|
||||
; CHECK-NEXT: }
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: {
|
||||
; CHECK-NEXT: .reg .b8 %e2m1x2_in;
|
||||
; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
|
||||
; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
|
||||
; CHECK-NEXT: }
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m1x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: {
|
||||
; CHECK-NEXT: .reg .b8 %e2m1x2_in;
|
||||
; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;
|
||||
; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e2m1x2 %r1, %e2m1x2_in, %rs2;
|
||||
; CHECK-NEXT: }
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m1x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
135
llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll
Normal file
135
llvm/test/CodeGen/NVPTX/convert_fp6x2_to_bf16x2.ll
Normal file
|
|
@ -0,0 +1,135 @@
|
|||
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
|
||||
; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
|
||||
; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
|
||||
|
||||
define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e2m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e2m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e2m3x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e3m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e3m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e3m2x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
135
llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll
Normal file
135
llvm/test/CodeGen/NVPTX/convert_fp8x2_to_bf16x2.ll
Normal file
|
|
@ -0,0 +1,135 @@
|
|||
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | FileCheck %s
|
||||
; RUN: %if ptxas-sm_100f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100f -mattr=+ptx92 | %ptxas-verify -arch=sm_100f %}
|
||||
; RUN: %if ptxas-sm_110f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_110f -mattr=+ptx92 | %ptxas-verify -arch=sm_110f %}
|
||||
; RUN: %if ptxas-sm_120f && ptxas-isa-9.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120f -mattr=+ptx92 | %ptxas-verify -arch=sm_120f %}
|
||||
|
||||
define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e4m3x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e4m3x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e4m3x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_relu_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.relu.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.satfinite.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
|
||||
define <2 x bfloat> @test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(i16 %a, i16 %b) {
|
||||
; CHECK-LABEL: test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0(
|
||||
; CHECK: {
|
||||
; CHECK-NEXT: .reg .b16 %rs<3>;
|
||||
; CHECK-NEXT: .reg .b32 %r<2>;
|
||||
; CHECK-EMPTY:
|
||||
; CHECK-NEXT: // %bb.0:
|
||||
; CHECK-NEXT: ld.param.b16 %rs1, [test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_0];
|
||||
; CHECK-NEXT: ld.param.b16 %rs2, [test_e5m2x2_to_bf16x2_rn_relu_satfinite_scale_ue8m0_param_1];
|
||||
; CHECK-NEXT: cvt.rn.relu.satfinite.scaled::n2::ue8m0.bf16x2.e5m2x2 %r1, %rs1, %rs2;
|
||||
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
||||
; CHECK-NEXT: ret;
|
||||
%val = call <2 x bfloat> @llvm.nvvm.e5m2x2.to.bf16x2.rn.relu.satfinite.scale.n2.ue8m0(i16 %a, i16 %b)
|
||||
ret <2 x bfloat> %val
|
||||
}
|
||||
Loading…
Reference in a new issue