[clang] [llvm] [mlir] [clang][NVPTX] Add support for pzo in f32/f16x2/bf16x2 to FP8 and FP6 conversions (PR #222511)
Dharuni R Acharya via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 10 03:53:37 PDT 2026
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/222511
>From f66a750b5a3f9074a4200d52fcebeb9c2a4f4a3b Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 10 Sep 2026 04:55:12 +0000
Subject: [PATCH 1/2] [clang][NVPTX] Add support for pzo in f32/f16x2/bf16x2 to
FP8 and FP6 conversions
This patch adds support for pzo variants and rz rounding mode to existing
f32/f16x2/bf16x2 to FP8 (e4m3x2, e5m2x2) and FP6 (e2m3x2, e3m2x2) conversion intrinsics.
Also adds clang builtins for the new variants.
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 <dharunira at nvidia.com>
---
clang/include/clang/Basic/BuiltinsNVPTX.td | 78 ++++
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 50 +++
clang/test/CodeGen/builtins-nvptx.c | 59 ++-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 48 ++-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 18 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 115 ++---
llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll | 407 ++++++++++++++++++
llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll | 407 ++++++++++++++++++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 19 +-
.../Target/LLVMIR/nvvm/convert_fp6x2.mlir | 36 +-
.../Target/LLVMIR/nvvm/convert_fp8x2.mlir | 36 +-
11 files changed, 1140 insertions(+), 133 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll
create mode 100644 llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 3475b721e95e0..8f1b43f9af65d 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -783,6 +783,84 @@ def __nvvm_f32x4_to_e3m2x4_rs_relu_satfinite :
NVPTXBuiltinSMAndPTX<"_Vector<4, char>(_Vector<4, float>, uint32_t)",
SMa<[100, 103]>, PTX87>;
+def __nvvm_ff_to_e4m3x2_rz : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e4m3x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e5m2x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e4m3x2_rz : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_relu : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e4m3x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e5m2x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e4m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e4m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e5m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+
+def __nvvm_ff_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+def __nvvm_ff_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(float, float)", SM_107f, PTX94>;
+
+def __nvvm_f16x2_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+def __nvvm_f16x2_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __fp16>)", SM_107f, PTX94>;
+
+def __nvvm_bf16x2_to_e2m3x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e2m3x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+def __nvvm_bf16x2_to_e3m2x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"short(_Vector<2, __bf16>)", SM_107f, PTX94>;
+
def __nvvm_ff_to_e2m1x2_rn_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SMa<[100, 101, 120]>, PTX86>;
def __nvvm_ff_to_e2m1x2_rn_relu_satfinite : NVPTXBuiltinSMAndPTX<"short(float, float)", SMa<[100, 101, 120]>, PTX86>;
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index 06c5069d6f984..1c4ea2d5e9143 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -1028,6 +1028,56 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
PZO_CVT(f2bf16_rz_satfinite);
PZO_CVT(f2bf16_rz_relu_satfinite);
+ PZO_CVT(ff_to_e4m3x2_rn);
+ PZO_CVT(ff_to_e4m3x2_rn_relu);
+ PZO_CVT(ff_to_e4m3x2_rz);
+ PZO_CVT(ff_to_e4m3x2_rz_relu);
+ PZO_CVT(ff_to_e5m2x2_rn);
+ PZO_CVT(ff_to_e5m2x2_rn_relu);
+ PZO_CVT(ff_to_e5m2x2_rz);
+ PZO_CVT(ff_to_e5m2x2_rz_relu);
+ PZO_CVT(f16x2_to_e4m3x2_rn);
+ PZO_CVT(f16x2_to_e4m3x2_rn_relu);
+ PZO_CVT(f16x2_to_e4m3x2_rz);
+ PZO_CVT(f16x2_to_e4m3x2_rz_relu);
+ PZO_CVT(f16x2_to_e5m2x2_rn);
+ PZO_CVT(f16x2_to_e5m2x2_rn_relu);
+ PZO_CVT(f16x2_to_e5m2x2_rz);
+ PZO_CVT(f16x2_to_e5m2x2_rz_relu);
+ PZO_CVT(bf16x2_to_e4m3x2_rn_satfinite);
+ PZO_CVT(bf16x2_to_e4m3x2_rn_relu_satfinite);
+ PZO_CVT(bf16x2_to_e4m3x2_rz_satfinite);
+ PZO_CVT(bf16x2_to_e4m3x2_rz_relu_satfinite);
+ PZO_CVT(bf16x2_to_e5m2x2_rn_satfinite);
+ PZO_CVT(bf16x2_to_e5m2x2_rn_relu_satfinite);
+ PZO_CVT(bf16x2_to_e5m2x2_rz_satfinite);
+ PZO_CVT(bf16x2_to_e5m2x2_rz_relu_satfinite);
+
+ PZO_CVT(ff_to_e2m3x2_rn_satfinite);
+ PZO_CVT(ff_to_e2m3x2_rn_relu_satfinite);
+ PZO_CVT(ff_to_e2m3x2_rz_satfinite);
+ PZO_CVT(ff_to_e2m3x2_rz_relu_satfinite);
+ PZO_CVT(ff_to_e3m2x2_rn_satfinite);
+ PZO_CVT(ff_to_e3m2x2_rn_relu_satfinite);
+ PZO_CVT(ff_to_e3m2x2_rz_satfinite);
+ PZO_CVT(ff_to_e3m2x2_rz_relu_satfinite);
+ PZO_CVT(f16x2_to_e2m3x2_rn_satfinite);
+ PZO_CVT(f16x2_to_e2m3x2_rn_relu_satfinite);
+ PZO_CVT(f16x2_to_e2m3x2_rz_satfinite);
+ PZO_CVT(f16x2_to_e2m3x2_rz_relu_satfinite);
+ PZO_CVT(f16x2_to_e3m2x2_rn_satfinite);
+ PZO_CVT(f16x2_to_e3m2x2_rn_relu_satfinite);
+ PZO_CVT(f16x2_to_e3m2x2_rz_satfinite);
+ PZO_CVT(f16x2_to_e3m2x2_rz_relu_satfinite);
+ PZO_CVT(bf16x2_to_e2m3x2_rn_satfinite);
+ PZO_CVT(bf16x2_to_e2m3x2_rn_relu_satfinite);
+ PZO_CVT(bf16x2_to_e2m3x2_rz_satfinite);
+ PZO_CVT(bf16x2_to_e2m3x2_rz_relu_satfinite);
+ PZO_CVT(bf16x2_to_e3m2x2_rn_satfinite);
+ PZO_CVT(bf16x2_to_e3m2x2_rn_relu_satfinite);
+ PZO_CVT(bf16x2_to_e3m2x2_rz_satfinite);
+ PZO_CVT(bf16x2_to_e3m2x2_rz_relu_satfinite);
+
#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 bed1498236b06..b74e9c74d84a9 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -1187,6 +1187,25 @@ __device__ void nvvm_cvt_pzo_sm107f() {
__nvvm_f2f16_rz_satfinite_pzo(1);
// CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rz.relu.satfinite(float 1.000000e+00, i1 true)
__nvvm_f2f16_rz_relu_satfinite_pzo(1);
+
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
+ __nvvm_ff_to_e4m3x2_rz_relu(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff_to_e4m3x2_rz_relu_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> zeroinitializer, i1 true)
+ __nvvm_f16x2_to_e5m2x2_rn_pzo({0, 0});
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.satfinite(<2 x bfloat> zeroinitializer, i1 true)
+ __nvvm_bf16x2_to_e4m3x2_rz_satfinite_pzo({0, 0});
+
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ __nvvm_ff_to_e2m3x2_rz_relu_satfinite(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff_to_e2m3x2_rz_relu_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> zeroinitializer, i1 true)
+ __nvvm_f16x2_to_e3m2x2_rn_satfinite_pzo({0, 0});
+ // CHECK_PTX94_SM107f: call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.satfinite(<2 x bfloat> zeroinitializer, i1 true)
+ __nvvm_bf16x2_to_e2m3x2_rz_satfinite_pzo({0, 0});
+
#endif
// CHECK: ret void
}
@@ -1194,22 +1213,22 @@ __device__ void nvvm_cvt_pzo_sm107f() {
// CHECK-LABEL: nvvm_cvt_sm89
__device__ void nvvm_cvt_sm89() {
#if (PTX >= 81) && (__CUDA_ARCH__ >= 890)
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e4m3x2_rn(1.0f, 1.0f);
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e4m3x2_rn_relu(1.0f, 1.0f);
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e5m2x2_rn(1.0f, 1.0f);
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e5m2x2_rn_relu(1.0f, 1.0f);
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> splat (half 1.000000e+00))
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> splat (half 1.000000e+00), i1 false)
__nvvm_f16x2_to_e4m3x2_rn({1.0f16, 1.0f16});
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> splat (half 1.000000e+00))
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> splat (half 1.000000e+00), i1 false)
__nvvm_f16x2_to_e4m3x2_rn_relu({1.0f16, 1.0f16});
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> splat (half 1.000000e+00))
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> splat (half 1.000000e+00), i1 false)
__nvvm_f16x2_to_e5m2x2_rn({1.0f16, 1.0f16});
- // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> splat (half 1.000000e+00))
+ // CHECK_PTX81_SM89: call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> splat (half 1.000000e+00), i1 false)
__nvvm_f16x2_to_e5m2x2_rn_relu({1.0f16, 1.0f16});
// CHECK_PTX81_SM89: call <2 x half> @llvm.nvvm.e4m3x2.to.f16x2.rn(i16 18504)
@@ -1259,24 +1278,24 @@ __device__ void nvvm_cvt_sm100a_sm101a_sm120a() {
#if (PTX >= 86) && \
(__CUDA_ARCH_FEAT_SM100_ALL || __CUDA_ARCH_FEAT_SM101_ALL || \
__CUDA_ARCH_FEAT_SM120_ALL)
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e2m3x2_rn_satfinite(1.0f, 1.0f);
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e2m3x2_rn_relu_satfinite(1.0f, 1.0f);
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e3m2x2_rn_satfinite(1.0f, 1.0f);
- // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
- // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX86_SM100a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM101a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
+ // CHECK_PTX86_SM120a: call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff_to_e3m2x2_rn_relu_satfinite(1.0f, 1.0f);
// CHECK_PTX86_SM100a: call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 19532)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 87ae664ce4c17..bc7b6d58653d6 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2057,18 +2057,25 @@ let TargetPrefix = "nvvm" in {
foreach type = ["e4m3x2", "e5m2x2"] in {
foreach relu = ["", "_relu"] in {
- def int_nvvm_ff_to_ # type # _rn # relu : NVVMBuiltin,
- PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;
+ foreach rnd = ["rn", "rz"] in {
+ def int_nvvm_ff_to_ # type # _ # rnd # relu : NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
- def int_nvvm_f16x2_to_ # type # _rn # relu : NVVMBuiltin,
- PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;
+ 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_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_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
-
- 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]>;
@@ -2115,18 +2122,27 @@ let TargetPrefix = "nvvm" in {
// FP6 conversions.
foreach type = ["e2m3x2", "e3m2x2"] in {
foreach relu = ["", "_relu"] in {
- def int_nvvm_ff_to_ # type # _rn # relu # _satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;
+ foreach rnd = ["rn", "rz"] in {
+ def int_nvvm_ff_to_ # type # _ # rnd # relu # _satfinite :
+ NVVMBuiltin,
+ PureIntrinsic<[llvm_i16_ty],
+ [llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, 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_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_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
-
- def int_nvvm_f16x2_to_ # type # _rn # relu # _satfinite
- : PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;
- 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]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 846bcc0557464..140efeeffe55c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -741,16 +741,19 @@ let hasSideEffects = false in {
def _f32 :
BasicFlagsNVPTXInst<(outs B16:$dst),
(ins B32:$src1, B32:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f32">,
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ F8Name # "x2.f32">,
Requires<[PTX81, SM89]>;
def _f16x2 :
BasicFlagsNVPTXInst<(outs B16:$dst),
(ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f16x2">,
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ F8Name # "x2.f16x2">,
Requires<[PTX81, SM89]>;
def _bf16x2 :
BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.bf16x2">,
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ F8Name # "x2.bf16x2">,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
}
@@ -849,14 +852,17 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
def _f32_sf :
BasicFlagsNVPTXInst<(outs B16:$dst),
(ins B32:$src1, B32:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.f32">;
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ FP6Name # "x2.f32">;
def _f16x2_sf :
BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.f16x2">,
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ FP6Name # "x2.f16x2">,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
def _bf16x2_sf :
BasicFlagsNVPTXInst<(outs B16:$dst), (ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}.satfinite${mode:relu}." # FP6Name # "x2.bf16x2">,
+ "cvt${mode:base}.satfinite${mode:relu}${mode:pzo}." #
+ FP6Name # "x2.bf16x2">,
Requires<[hasFP16X2ToNarrowFPConversionSupport]>;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 220ef64732830..0ccdb9abacd48 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2838,39 +2838,42 @@ def : Pat<(int_nvvm_ull2d_rp i64:$a), (CVT_f64_u64 $a, CvtRP)>;
def : Pat<(int_nvvm_f2h_rn_ftz f32:$a), (CVT_f16_f32 $a, CvtRN_FTZ)>;
def : Pat<(int_nvvm_f2h_rn f32:$a), (CVT_f16_f32 $a, CvtRN)>;
-let Predicates = [hasFP8ConversionSupport] in {
- def : Pat<(int_nvvm_ff_to_e4m3x2_rn f32:$a, f32:$b),
- (CVT_e4m3x2_f32 $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff_to_e4m3x2_rn_relu f32:$a, f32:$b),
- (CVT_e4m3x2_f32 $a, $b, CvtRN_RELU)>;
- def : Pat<(int_nvvm_ff_to_e5m2x2_rn f32:$a, f32:$b),
- (CVT_e5m2x2_f32 $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff_to_e5m2x2_rn_relu f32:$a, f32:$b),
- (CVT_e5m2x2_f32 $a, $b, CvtRN_RELU)>;
-}
-
-let Predicates = [hasFP16X2ToNarrowFPConversionSupport] in {
- foreach dst_type = ["e4m3x2", "e5m2x2"] in {
- foreach relu = ["", "_relu"] in {
- defvar intrin = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # dst_type # "_rn" # relu # "_satfinite");
- defvar cvt_inst = !cast<NVPTXInst>("CVT_"# dst_type #"_bf16x2");
- defvar cvt_mode = !cast<PatLeaf>("CvtRN" # !toupper(relu));
- def : Pat<(intrin v2bf16:$a),
- (cvt_inst $a, cvt_mode)>;
+foreach dst = ["e4m3x2", "e5m2x2"] in {
+ foreach relu = ["", "_relu"] in {
+ foreach rnd = ["rn", "rz"] in {
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
+ defvar Suffix = dst # "_" # rnd # relu;
+ defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix);
+ defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix);
+ defvar BF16X2 =
+ !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite");
+
+ foreach pzo = [0, 1] in {
+ defvar PZO = !if(pzo, -1, 0);
+ defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
+ defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1));
+ defvar FFPreds =
+ !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP8ConversionSupport]<Predicate>);
+ defvar BF16X2Preds =
+ !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
+
+ def : Pat<(FF f32:$a, f32:$b, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f32") $a, $b, Mode)>,
+ Requires<FFPreds>;
+ def : Pat<(F16X2 v2f16:$a, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f16x2") $a, Mode)>,
+ Requires<FFPreds>;
+ def : Pat<(BF16X2 v2bf16:$a, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>,
+ Requires<BF16X2Preds>;
+ }
}
}
}
let Predicates = [hasFP8ConversionSupport] in {
- def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn v2f16:$a),
- (CVT_e4m3x2_f16x2 $a, CvtRN)>;
- def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn_relu v2f16:$a),
- (CVT_e4m3x2_f16x2 $a, CvtRN_RELU)>;
- def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn v2f16:$a),
- (CVT_e5m2x2_f16x2 $a, CvtRN)>;
- def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn_relu v2f16:$a),
- (CVT_e5m2x2_f16x2 $a, CvtRN_RELU)>;
-
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn i16:$a),
(CVT_f16x2_e4m3x2 $a, CvtRN)>;
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn_relu i16:$a),
@@ -2918,16 +2921,41 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
(CVT_bf16x2_s2f6x2_sf_scale $a, $scale, CvtRN_RELU)>;
}
-let Predicates = [hasNarrowFPConversionSupport] in {
- def : Pat<(int_nvvm_ff_to_e2m3x2_rn_satfinite f32:$a, f32:$b),
- (CVT_e2m3x2_f32_sf $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff_to_e2m3x2_rn_relu_satfinite f32:$a, f32:$b),
- (CVT_e2m3x2_f32_sf $a, $b, CvtRN_RELU)>;
- def : Pat<(int_nvvm_ff_to_e3m2x2_rn_satfinite f32:$a, f32:$b),
- (CVT_e3m2x2_f32_sf $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff_to_e3m2x2_rn_relu_satfinite f32:$a, f32:$b),
- (CVT_e3m2x2_f32_sf $a, $b, CvtRN_RELU)>;
+foreach dst = ["e2m3x2", "e3m2x2"] in {
+ foreach relu = ["", "_relu"] in {
+ foreach rnd = ["rn", "rz"] in {
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
+ defvar Suffix = dst # "_" # rnd # relu # "_satfinite";
+ defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix);
+ defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix);
+ defvar BF16X2 = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix);
+
+ foreach pzo = [0, 1] in {
+ defvar PZO = !if(pzo, -1, 0);
+ defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
+ defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1));
+ defvar FFPreds =
+ !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasNarrowFPConversionSupport]<Predicate>);
+ defvar FPX2Preds =
+ !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
+
+ def : Pat<(FF f32:$a, f32:$b, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f32_sf") $a, $b, Mode)>,
+ Requires<FFPreds>;
+ def : Pat<(F16X2 v2f16:$a, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf") $a, Mode)>,
+ Requires<FPX2Preds>;
+ def : Pat<(BF16X2 v2bf16:$a, PZO),
+ (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>,
+ Requires<FPX2Preds>;
+ }
+ }
+ }
+}
+let Predicates = [hasNarrowFPConversionSupport] in {
def : Pat<(int_nvvm_e2m3x2_to_f16x2_rn i16:$a),
(CVT_f16x2_e2m3x2 $a, CvtRN)>;
def : Pat<(int_nvvm_e2m3x2_to_f16x2_rn_relu i16:$a),
@@ -2938,19 +2966,6 @@ let Predicates = [hasNarrowFPConversionSupport] in {
(CVT_f16x2_e3m2x2 $a, CvtRN_RELU)>;
}
-let Predicates = [hasFP16X2ToNarrowFPConversionSupport] in {
- foreach src_type = ["f16x2", "bf16x2"] in {
- foreach dst_type = ["e2m3x2", "e3m2x2"] in {
- foreach relu = ["", "_relu"] in {
- defvar intrin = !cast<Intrinsic>("int_nvvm_" # src_type # "_to_" # dst_type # "_rn" # relu # "_satfinite");
- defvar cvt_inst = !cast<NVPTXInst>("CVT_" # dst_type # "_" # src_type # "_sf");
- defvar cvt_mode = !cast<PatLeaf>("CvtRN" # !toupper(relu));
- def : Pat<(intrin B32:$a), (cvt_inst $a, cvt_mode)>;
- }
- }
- }
-}
-
let Predicates = [hasNarrowFPToBF16x2ConversionSupport] in {
foreach src_type = ["e2m3x2", "e3m2x2"] in {
foreach relu = ["", "_relu"] in {
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll
new file mode 100644
index 0000000000000..01c1554e99644
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp6-pzo.ll
@@ -0,0 +1,407 @@
+; 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 conversions from f32
+
+define i16 @cvt_rn_pzo_e2m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_e2m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e2m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e2m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_e2m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e2m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rz.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e2m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e2m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; 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_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e2m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rz.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+; E3M2X2 conversions from f32
+
+define i16 @cvt_rn_pzo_e3m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_e3m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e3m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e3m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_e3m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e3m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rz.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e3m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e3m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e3m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e3m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rz.relu.satfinite(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+; E2M3X2 conversions from f16x2
+
+define i16 @cvt_rn_pzo_e2m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e2m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e2m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e2m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rz.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e2m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e2m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rz.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+; E3M2X2 conversions from f16x2
+
+define i16 @cvt_rn_pzo_e3m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e3m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e3m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e3m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rz.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e3m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e3m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rz.relu.satfinite(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+; E2M3X2 conversions from bf16x2
+
+define i16 @cvt_rn_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e2m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e2m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e2m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e2m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e2m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e2m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e2m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e2m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e2m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e2m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e2m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; 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_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e2m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+; E3M2X2 conversions from bf16x2
+
+define i16 @cvt_rn_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e3m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e3m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e3m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e3m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e3m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e3m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rz.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e3m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e3m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e3m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e3m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e3m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e3m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e3m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll
new file mode 100644
index 0000000000000..01ea380d61403
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp8-pzo.ll
@@ -0,0 +1,407 @@
+; 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 conversions from f32
+
+define i16 @cvt_rn_pzo_e4m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_e4m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e4m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e4m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_e4m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e4m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rz(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e4m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e4m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e4m3x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; 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_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e4m3x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e4m3x2.rz.relu(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+; E5M2X2 conversions from f32
+
+define i16 @cvt_rn_pzo_e5m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_e5m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_e5m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e5m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_e5m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_e5m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rz(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e5m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_e5m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e5m2x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_e5m2x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.f32 %rs1, %r1, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.ff.to.e5m2x2.rz.relu(float %f1, float %f2, i1 true)
+ ret i16 %val
+}
+
+; E4M3X2 conversions from f16x2
+
+define i16 @cvt_rn_pzo_e4m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e4m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e4m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e4m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rz(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e4m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e4m3x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e4m3x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rz.relu(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+; E5M2X2 conversions from f16x2
+
+define i16 @cvt_rn_pzo_e5m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e5m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e5m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e5m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rz(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e5m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e5m2x2_f16x2(<2 x half> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_f16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.f16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rz.relu(<2 x half> %a, i1 true)
+ ret i16 %val
+}
+
+; E4M3X2 conversions from bf16x2
+
+define i16 @cvt_rn_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e4m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e4m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e4m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e4m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e4m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e4m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e4m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e4m3x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e4m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e4m3x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e4m3x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; 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_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e4m3x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+; E5M2X2 conversions from bf16x2
+
+define i16 @cvt_rn_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_pzo_e5m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_e5m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.e5m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_pzo_e5m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_e5m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.e5m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rz.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rn_relu_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rn_relu_pzo_e5m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_e5m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.relu.pzo.e5m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
+
+define i16 @cvt_rz_relu_pzo_e5m2x2_bf16x2(<2 x bfloat> %a) {
+; CHECK-LABEL: cvt_rz_relu_pzo_e5m2x2_bf16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_e5m2x2_bf16x2_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.relu.pzo.e5m2x2.bf16x2 %rs1, %r1;
+; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %val = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rz.relu.satfinite(<2 x bfloat> %a, i1 true)
+ ret i16 %val
+}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 80a0094190383..ffad9fe6d0d04 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -2085,7 +2085,8 @@ def NVVM_ConvertF32x2ToF6x2Op : NVVM_Op<"convert.f32x2.to.f6x2"> {
string llvmBuilder = [{
auto intId = NVVM::ConvertF32x2ToF6x2Op::getIntrinsicID($dstTy, $relu);
- llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a, $b});
+ llvm::Value *packedI16 =
+ createIntrinsicCall(builder, intId, {$a, $b, builder.getInt1(false)});
if(op.getDst().getType().isInteger(16))
$dst = packedI16;
else
@@ -2127,7 +2128,8 @@ class NVVM_ConvertFPx2ToF6x2Op<string srcType>
string llvmBuilder = [{
auto intId = NVVM::Convert}]#srcType#[{x2ToF6x2Op::getIntrinsicID($dstTy, $relu);
- llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$src});
+ llvm::Value *packedI16 =
+ createIntrinsicCall(builder, intId, {$src, builder.getInt1(false)});
if(op.getDst().getType().isInteger(16))
$dst = packedI16;
else
@@ -2181,7 +2183,10 @@ def NVVM_ConvertF32x2ToF8x2Op : NVVM_Op<"convert.f32x2.to.f8x2"> {
string llvmBuilder = [{
auto intId = NVVM::ConvertF32x2ToF8x2Op::getIntrinsicID($dstTy, $rnd, $sat, $relu);
- llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a, $b});
+ llvm::SmallVector<llvm::Value *> args = {$a, $b};
+ if (!mlir::isa<mlir::Float8E8M0FNUType>($dstTy))
+ args.push_back(builder.getInt1(false));
+ llvm::Value *packedI16 = createIntrinsicCall(builder, intId, args);
if(op.getDst().getType().isInteger(16))
$dst = packedI16;
else
@@ -2227,7 +2232,8 @@ def NVVM_ConvertF16x2ToF8x2Op : NVVM_Op<"convert.f16x2.to.f8x2"> {
string llvmBuilder = [{
auto intId = NVVM::ConvertF16x2ToF8x2Op::getIntrinsicID($dstTy, $relu);
- llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$a});
+ llvm::Value *packedI16 =
+ createIntrinsicCall(builder, intId, {$a, builder.getInt1(false)});
if(op.getDst().getType().isInteger(16))
$dst = packedI16;
else
@@ -2277,7 +2283,10 @@ def NVVM_ConvertBF16x2ToF8x2Op : NVVM_Op<"convert.bf16x2.to.f8x2"> {
string llvmBuilder = [{
auto intId = NVVM::ConvertBF16x2ToF8x2Op::getIntrinsicID($dstTy, $rnd, $sat, $relu);
- llvm::Value *packedI16 = createIntrinsicCall(builder, intId, {$src});
+ llvm::SmallVector<llvm::Value *> args = {$src};
+ if (!mlir::isa<mlir::Float8E8M0FNUType>($dstTy))
+ args.push_back(builder.getInt1(false));
+ llvm::Value *packedI16 = createIntrinsicCall(builder, intId, args);
if(op.getDst().getType().isInteger(16))
$dst = packedI16;
else
diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir
index 6cf4e8d2b98fd..21f23521405b1 100644
--- a/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp6x2.mlir
@@ -2,28 +2,28 @@
// CHECK-LABEL: @convert_f32x2_to_fp6x2_e2m3
llvm.func @convert_f32x2_to_fp6x2_e2m3(%srcA : f32, %srcB : f32) {
- //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : i16 (f6E2M3FN)
- //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB relu = true : i16 (f6E2M3FN)
llvm.return
}
// CHECK-LABEL: @convert_f32x2_to_fp6x2_e3m2
llvm.func @convert_f32x2_to_fp6x2_e3m2(%srcA : f32, %srcB : f32) {
- //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : i16 (f6E3M2FN)
- //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB relu = true : i16 (f6E3M2FN)
llvm.return
}
// CHECK-LABEL: @convert_f32x2_to_fp6x2_vector
llvm.func @convert_f32x2_to_fp6x2_vector(%srcA : f32, %srcB : f32) {
- //CHECK: %[[res0:.*]] = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %[[res0:.*]] = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
//CHECK-NEXT: %{{.*}} = bitcast i16 %[[res0]] to <2 x i8>
%res1 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : vector<2xi8> (f6E2M3FN)
- //CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ //CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
//CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8>
%res2 = nvvm.convert.f32x2.to.f6x2 %srcA, %srcB : vector<2xi8> (f6E3M2FN)
llvm.return
@@ -33,28 +33,28 @@ llvm.func @convert_f32x2_to_fp6x2_vector(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f16x2_to_fp6x2_e2m3
llvm.func @convert_f16x2_to_fp6x2_e2m3(%srcA : vector<2xf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}, i1 false)
%res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> i16 (f6E2M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.relu.satfinite(<2 x half> %{{.*}}, i1 false)
%res2 = nvvm.convert.f16x2.to.f6x2 %srcA relu = true : vector<2xf16> -> i16 (f6E2M3FN)
llvm.return
}
// CHECK-LABEL: @convert_f16x2_to_fp6x2_e3m2
llvm.func @convert_f16x2_to_fp6x2_e3m2(%srcA : vector<2xf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}, i1 false)
%res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> i16 (f6E3M2FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.relu.satfinite(<2 x half> %{{.*}}, i1 false)
%res2 = nvvm.convert.f16x2.to.f6x2 %srcA relu = true : vector<2xf16> -> i16 (f6E3M2FN)
llvm.return
}
// CHECK-LABEL: @convert_f16x2_to_fp6x2_vector
llvm.func @convert_f16x2_to_fp6x2_vector(%srcA : vector<2xf16>) {
- // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}})
+ // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e2m3x2.rn.satfinite(<2 x half> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8>
%res1 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> vector<2xi8> (f6E2M3FN)
- // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}})
+ // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e3m2x2.rn.satfinite(<2 x half> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8>
%res2 = nvvm.convert.f16x2.to.f6x2 %srcA : vector<2xf16> -> vector<2xi8> (f6E3M2FN)
llvm.return
@@ -64,28 +64,28 @@ llvm.func @convert_f16x2_to_fp6x2_vector(%srcA : vector<2xf16>) {
// CHECK-LABEL: @convert_bf16x2_to_fp6x2_e2m3
llvm.func @convert_bf16x2_to_fp6x2_e2m3(%srcA : vector<2xbf16>, %scale_factor : i16) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> i16 (f6E2M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res2 = nvvm.convert.bf16x2.to.f6x2 %srcA relu = true : vector<2xbf16> -> i16 (f6E2M3FN)
llvm.return
}
// CHECK-LABEL: @convert_bf16x2_to_fp6x2_e3m2
llvm.func @convert_bf16x2_to_fp6x2_e3m2(%srcA : vector<2xbf16>, %scale_factor : i16) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> i16 (f6E3M2FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res2 = nvvm.convert.bf16x2.to.f6x2 %srcA relu = true : vector<2xbf16> -> i16 (f6E3M2FN)
llvm.return
}
// CHECK-LABEL: @convert_bf16x2_to_fp6x2_vector
llvm.func @convert_bf16x2_to_fp6x2_vector(%srcA : vector<2xbf16>, %scale_factor : i16) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e2m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %{{.*}} to <2 x i8>
%res1 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> vector<2xi8> (f6E2M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e3m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %{{.*}} to <2 x i8>
%res2 = nvvm.convert.bf16x2.to.f6x2 %srcA : vector<2xbf16> -> vector<2xi8> (f6E3M2FN)
llvm.return
diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir
index 6a8745e065902..058e28be90283 100644
--- a/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp8x2.mlir
@@ -4,18 +4,18 @@
// CHECK-LABEL: @convert_f32x2_to_f8x2_e4m3
llvm.func @convert_f32x2_to_f8x2_e4m3(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : i16 (f8E4M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : i16 (f8E4M3FN)
llvm.return
}
// CHECK-LABEL: @convert_f32x2_to_f8x2_e5m2
llvm.func @convert_f32x2_to_f8x2_e5m2(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : i16 (f8E5M2)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.ff.to.e5m2x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : i16 (f8E5M2)
llvm.return
}
@@ -35,10 +35,10 @@ llvm.func @convert_f32x2_to_f8x2_ue8m0(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_f8x2_vector_return
llvm.func @convert_f32x2_to_f8x2_vector_return(%srcA : f32, %srcB : f32) {
- // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}})
+ // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn(float %{{.*}}, float %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8>
%res1 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : vector<2xi8> (f8E4M3FN)
- // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.ff.to.e4m3x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8>
%res2 = nvvm.convert.f32x2.to.f8x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : vector<2xi8> (f8E4M3FN)
llvm.return
@@ -48,28 +48,28 @@ llvm.func @convert_f32x2_to_f8x2_vector_return(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f16x2_to_f8x2_e4m3
llvm.func @convert_f16x2_to_f8x2_e4m3(%src : vector<2xf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}, i1 false)
%res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> i16 (f8E4M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn.relu(<2 x half> %{{.*}}, i1 false)
%res2 = nvvm.convert.f16x2.to.f8x2 %src relu = true : vector<2xf16> -> i16 (f8E4M3FN)
llvm.return
}
// CHECK-LABEL: @convert_f16x2_to_f8x2_e5m2
llvm.func @convert_f16x2_to_f8x2_e5m2(%src : vector<2xf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}, i1 false)
%res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> i16 (f8E5M2)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn.relu(<2 x half> %{{.*}}, i1 false)
%res2 = nvvm.convert.f16x2.to.f8x2 %src relu = true : vector<2xf16> -> i16 (f8E5M2)
llvm.return
}
// CHECK-LABEL: @convert_f16x2_to_f8x2_vector_return
llvm.func @convert_f16x2_to_f8x2_vector_return(%src : vector<2xf16>) {
- // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}})
+ // CHECK: %[[res1:.*]] = call i16 @llvm.nvvm.f16x2.to.e4m3x2.rn(<2 x half> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res1]] to <2 x i8>
%res1 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> vector<2xi8> (f8E4M3FN)
- // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}})
+ // CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.f16x2.to.e5m2x2.rn(<2 x half> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8>
%res2 = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> vector<2xi8> (f8E5M2)
llvm.return
@@ -93,18 +93,18 @@ llvm.func @convert_bf16x2_to_f8x2_ue8m0(%src : vector<2xbf16>) {
// CHECK-LABEL: @convert_bf16x2_to_f8x2_e4m3
llvm.func @convert_bf16x2_to_f8x2_e4m3(%srcA : vector<2xbf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res1 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> : vector<2xbf16> -> i16 (f8E4M3FN)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res2 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> i16 (f8E4M3FN)
llvm.return
}
// CHECK-LABEL: @convert_bf16x2_to_f8x2_e5m2
llvm.func @convert_bf16x2_to_f8x2_e5m2(%srcA : vector<2xbf16>) {
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res1 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E5M2)
- // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %{{.*}} = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false)
%res2 = nvvm.convert.bf16x2.to.f8x2 %srcA rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> vector<2xi8> (f8E5M2)
llvm.return
}
@@ -117,10 +117,10 @@ llvm.func @convert_bf16x2_to_f8x2_vector_return(%src : vector<2xbf16>) {
// CHECK: %[[res2:.*]] = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> %{{.*}})
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res2]] to <2 x i8>
%res2 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rp> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E8M0FNU)
- // CHECK: %[[res3:.*]] = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %[[res3:.*]] = call i16 @llvm.nvvm.bf16x2.to.e4m3x2.rn.satfinite(<2 x bfloat> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res3]] to <2 x i8>
%res3 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rn> sat = <satfinite> : vector<2xbf16> -> vector<2xi8> (f8E4M3FN)
- // CHECK: %[[res4:.*]] = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}})
+ // CHECK: %[[res4:.*]] = call i16 @llvm.nvvm.bf16x2.to.e5m2x2.rn.relu.satfinite(<2 x bfloat> %{{.*}}, i1 false)
// CHECK-NEXT: %{{.*}} = bitcast i16 %[[res4]] to <2 x i8>
%res4 = nvvm.convert.bf16x2.to.f8x2 %src rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16> -> vector<2xi8> (f8E5M2)
llvm.return
>From 5f7cc47d6950cacfe74c6ed45751343b45365562 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 10 Sep 2026 10:53:14 +0000
Subject: [PATCH 2/2] Address comments
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 84 +++++++++++-------------
1 file changed, 39 insertions(+), 45 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 0ccdb9abacd48..46f9fd1727613 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2709,10 +2709,13 @@ def : Pat<(int_nvvm_ui2f_rz i32:$a), (CVT_f32_u32 $a, CvtRZ)>;
def : Pat<(int_nvvm_ui2f_rm i32:$a), (CVT_f32_u32 $a, CvtRM)>;
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
+// 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 = !cast<Intrinsic>("int_nvvm_" # suffix);
- Intrinsic sf = !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite");
+ 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"), ?);
}
foreach rnd = ["rn", "rz"] in {
@@ -2839,34 +2842,29 @@ def : Pat<(int_nvvm_f2h_rn_ftz f32:$a), (CVT_f16_f32 $a, CvtRN_FTZ)>;
def : Pat<(int_nvvm_f2h_rn f32:$a), (CVT_f16_f32 $a, CvtRN)>;
foreach dst = ["e4m3x2", "e5m2x2"] in {
- foreach relu = ["", "_relu"] in {
+ foreach relu = ["", "relu"] in {
foreach rnd = ["rn", "rz"] in {
- defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
- defvar Suffix = dst # "_" # rnd # relu;
- defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix);
- defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix);
- defvar BF16X2 =
- !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix # "_satfinite");
+ defvar Relu = !if(!empty(relu), "", "_" # relu);
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # Relu));
+ defvar Suffix = dst # "_" # rnd # Relu;
+ defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
+ defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
+ defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1));
- defvar FFPreds =
- !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
- [hasFP8ConversionSupport]<Predicate>);
- defvar BF16X2Preds =
- !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
- [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
-
- def : Pat<(FF f32:$a, f32:$b, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_f32") $a, $b, Mode)>,
- Requires<FFPreds>;
- def : Pat<(F16X2 v2f16:$a, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_f16x2") $a, Mode)>,
- Requires<FFPreds>;
- def : Pat<(BF16X2 v2bf16:$a, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>,
+ defvar F32X2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP8ConversionSupport]<Predicate>);
+ defvar BF16X2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
+
+ def : Pat<(F32X2.base f32:$a, f32:$b, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_f32") $a, $b, Mode)>,
+ Requires<F32X2Preds>;
+ def : Pat<(F16X2.base v2f16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_f16x2") $a, Mode)>,
+ Requires<F32X2Preds>;
+ def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2") $a, Mode)>,
Requires<BF16X2Preds>;
}
}
@@ -2922,33 +2920,29 @@ let Predicates = [hasS2F6X2ConversionSupport] in {
}
foreach dst = ["e2m3x2", "e3m2x2"] in {
- foreach relu = ["", "_relu"] in {
+ foreach relu = ["", "relu"] in {
foreach rnd = ["rn", "rz"] in {
- defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
- defvar Suffix = dst # "_" # rnd # relu # "_satfinite";
- defvar FF = !cast<Intrinsic>("int_nvvm_ff_to_" # Suffix);
- defvar F16X2 = !cast<Intrinsic>("int_nvvm_f16x2_to_" # Suffix);
- defvar BF16X2 = !cast<Intrinsic>("int_nvvm_bf16x2_to_" # Suffix);
+ defvar Relu = !if(!empty(relu), "", "_" # relu);
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # Relu));
+ defvar Suffix = dst # "_" # rnd # Relu;
+ defvar F32X2 = ToCvtIntrinsics<"ff_to_" # Suffix>;
+ defvar F16X2 = ToCvtIntrinsics<"f16x2_to_" # Suffix>;
+ defvar BF16X2 = ToCvtIntrinsics<"bf16x2_to_" # Suffix>;
foreach pzo = [0, 1] in {
defvar PZO = !if(pzo, -1, 0);
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar NeedsPZOSupport = !or(!eq(rnd, "rz"), !eq(pzo, 1));
- defvar FFPreds =
- !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
- [hasNarrowFPConversionSupport]<Predicate>);
- defvar FPX2Preds =
- !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
- [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
-
- def : Pat<(FF f32:$a, f32:$b, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_f32_sf") $a, $b, Mode)>,
- Requires<FFPreds>;
- def : Pat<(F16X2 v2f16:$a, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf") $a, Mode)>,
+ defvar F32X2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasNarrowFPConversionSupport]<Predicate>);
+ defvar FPX2Preds = !if(NeedsPZOSupport, [hasConvertWithPZOSupport]<Predicate>,
+ [hasFP16X2ToNarrowFPConversionSupport]<Predicate>);
+
+ def : Pat<(F32X2.sf f32:$a, f32:$b, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_f32_sf") $a, $b, Mode)>,
+ Requires<F32X2Preds>;
+ def : Pat<(F16X2.sf v2f16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_f16x2_sf") $a, Mode)>,
Requires<FPX2Preds>;
- def : Pat<(BF16X2 v2bf16:$a, PZO),
- (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>,
+ def : Pat<(BF16X2.sf v2bf16:$a, PZO), (!cast<NVPTXInst>("CVT_" # dst # "_bf16x2_sf") $a, Mode)>,
Requires<FPX2Preds>;
}
}
More information about the cfe-commits
mailing list