[clang] db41750 - [clang][NVPTX] Add support for pzo in f32 to f16/bf16 conversions (#214667)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Sep 1 23:28:29 PDT 2026
Author: Srinivasa Ravi
Date: 2026-09-02T11:58:22+05:30
New Revision: db41750388a51f21e00f8017a6db148f34058109
URL: https://github.com/llvm/llvm-project/commit/db41750388a51f21e00f8017a6db148f34058109
DIFF: https://github.com/llvm/llvm-project/commit/db41750388a51f21e00f8017a6db148f34058109.diff
LOG: [clang][NVPTX] Add support for pzo in f32 to f16/bf16 conversions (#214667)
Add support for `pzo` variants to existing `f32` to `f16/bf16`
conversion
intrinsics through the use of a default argument (defaulting to
`false`).
Also adds clang builtins for the new variants and support for lowering
clang builtins to intrinsics with default arguments.
Added:
llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
Modified:
clang/include/clang/Basic/BuiltinsNVPTX.td
clang/lib/CodeGen/CGBuiltin.cpp
clang/lib/CodeGen/CGBuiltin.h
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
clang/test/CodeGen/builtins-nvptx.c
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
llvm/lib/Target/NVPTX/NVPTX.h
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index b91b356b3d295..3475b721e95e0 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -649,6 +649,14 @@ def __nvvm_ff2bf16x2_rs_satfinite :
def __nvvm_ff2bf16x2_rs_relu_satfinite :
NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float, uint32_t)",
SMa<[100, 103]>, PTX87>;
+def __nvvm_ff2bf16x2_rn_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rz_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2bf16x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __bf16>(float, float)", SM_107f, PTX94>;
def __nvvm_ff2f16x2_rn : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_80, PTX70>;
def __nvvm_ff2f16x2_rn_relu : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_80, PTX70>;
@@ -670,6 +678,14 @@ def __nvvm_ff2f16x2_rs_satfinite :
def __nvvm_ff2f16x2_rs_relu_satfinite :
NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float, uint32_t)",
SMa<[100, 103]>, PTX87>;
+def __nvvm_ff2f16x2_rn_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rz_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
+def __nvvm_ff2f16x2_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(float, float)", SM_107f, PTX94>;
def __nvvm_f2bf16_rn : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PTX70>;
def __nvvm_f2bf16_rn_relu : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PTX70>;
@@ -679,6 +695,14 @@ def __nvvm_f2bf16_rn_satfinite : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PT
def __nvvm_f2bf16_rn_relu_satfinite : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PTX81>;
def __nvvm_f2bf16_rz_satfinite : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PTX81>;
def __nvvm_f2bf16_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_80, PTX81>;
+def __nvvm_f2bf16_rn_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rz_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
+def __nvvm_f2bf16_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__bf16(float)", SM_107f, PTX94>;
def __nvvm_f2f16_rn : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX70>;
def __nvvm_f2f16_rn_relu : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX70>;
@@ -688,6 +712,14 @@ def __nvvm_f2f16_rn_satfinite : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX
def __nvvm_f2f16_rn_relu_satfinite : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX81>;
def __nvvm_f2f16_rz_satfinite : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX81>;
def __nvvm_f2f16_rz_relu_satfinite : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_80, PTX81>;
+def __nvvm_f2f16_rn_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rn_relu_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rz_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rz_relu_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rn_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rn_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rz_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
+def __nvvm_f2f16_rz_relu_satfinite_pzo : NVPTXBuiltinSMAndPTX<"__fp16(float)", SM_107f, PTX94>;
def __nvvm_f2tf32_rna : NVPTXBuiltinSMAndPTX<"int32_t(float)", SM_80, PTX70>;
def __nvvm_f2tf32_rna_satfinite : NVPTXBuiltinSMAndPTX<"int32_t(float)", SM_80, PTX81>;
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 07379056f2f0a..16c271d8fd82c 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -253,6 +253,27 @@ llvm::Constant *CodeGenModule::getBuiltinLibFunction(const FunctionDecl *FD,
return GetOrCreateLLVMFunction(Name, Ty, D, /*ForVTable=*/false);
}
+void appendDefaultIntrinsicArgs(SmallVectorImpl<llvm::Value *> &Args,
+ llvm::Function *F) {
+ llvm::FunctionType *FTy = F->getFunctionType();
+ unsigned NumParams = FTy->getNumParams();
+ if (Args.size() >= NumParams)
+ return;
+
+ auto [FirstDefault, Defaults] =
+ Intrinsic::getAllDefaultArgValues(F->getIntrinsicID());
+ assert(Args.size() >= FirstDefault &&
+ "builtin passes fewer arguments than the intrinsic requires");
+
+ for (unsigned I = Args.size(); I != NumParams; ++I) {
+ llvm::Type *ParamTy = FTy->getParamType(I);
+ unsigned DefaultIdx = I - FirstDefault;
+ assert(ParamTy->isIntegerTy() &&
+ "intrinsic default arguments must be integer-typed");
+ Args.push_back(llvm::ConstantInt::get(ParamTy, Defaults[DefaultIdx]));
+ }
+}
+
/// Emit the conversions required to turn the given value into an
/// integer of the given size.
Value *EmitToInt(CodeGenFunction &CGF, llvm::Value *V,
@@ -6908,6 +6929,8 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
Args.push_back(ArgValue);
}
+ appendDefaultIntrinsicArgs(Args, F);
+
Value *V = Builder.CreateCall(F, Args);
QualType BuiltinRetType = E->getType();
diff --git a/clang/lib/CodeGen/CGBuiltin.h b/clang/lib/CodeGen/CGBuiltin.h
index df71e46629884..394b397fde6a9 100644
--- a/clang/lib/CodeGen/CGBuiltin.h
+++ b/clang/lib/CodeGen/CGBuiltin.h
@@ -72,6 +72,11 @@ llvm::Value *emitBuiltinWithOneOverloadedType(clang::CodeGen::CodeGenFunction &C
return CGF.Builder.CreateCall(F, Args, Name);
}
+// Fills in the trailing parameters of an intrinsic that the builtin does not
+// expose, using the values declared via ImmArg<..., DefaultValue<...>>.
+void appendDefaultIntrinsicArgs(llvm::SmallVectorImpl<llvm::Value *> &Args,
+ llvm::Function *F);
+
llvm::Value *emitUnaryMaybeConstrainedFPBuiltin(clang::CodeGen::CodeGenFunction &CGF,
const clang::CallExpr *E,
unsigned IntrinsicID,
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index 64fdae9d8934d..e3ef3ae488a0c 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -395,7 +395,8 @@ static Value *MakeCpAsync(unsigned IntrinsicID, unsigned IntrinsicIDS,
}
static Value *MakeHalfType(Function *Intrinsic, unsigned BuiltinID,
- const CallExpr *E, CodeGenFunction &CGF) {
+ const CallExpr *E, CodeGenFunction &CGF,
+ ArrayRef<Value *> TrailingArgs = {}) {
SmallVector<Value *, 16> Args;
auto *FTy = Intrinsic->getFunctionType();
unsigned ICEArguments = 0;
@@ -411,12 +412,17 @@ static Value *MakeHalfType(Function *Intrinsic, unsigned BuiltinID,
Args.push_back(ArgValue);
}
+ llvm::append_range(Args, TrailingArgs);
+ appendDefaultIntrinsicArgs(Args, Intrinsic);
+
return CGF.Builder.CreateCall(Intrinsic, Args);
}
static Value *MakeHalfType(unsigned IntrinsicID, unsigned BuiltinID,
- const CallExpr *E, CodeGenFunction &CGF) {
- return MakeHalfType(CGF.CGM.getIntrinsic(IntrinsicID), BuiltinID, E, CGF);
+ const CallExpr *E, CodeGenFunction &CGF,
+ ArrayRef<Value *> TrailingArgs = {}) {
+ return MakeHalfType(CGF.CGM.getIntrinsic(IntrinsicID), BuiltinID, E, CGF,
+ TrailingArgs);
}
static Value *MakeFMAOOB(unsigned IntrinsicID, llvm::Type *Ty,
@@ -975,6 +981,47 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
return MakeHalfType(Intrinsic::nvvm_ff2f16x2_rz, BuiltinID, E, *this);
case NVPTX::BI__nvvm_ff2f16x2_rz_relu:
return MakeHalfType(Intrinsic::nvvm_ff2f16x2_rz_relu, BuiltinID, E, *this);
+
+#define PZO_CVT(cvt) \
+ case NVPTX::BI__nvvm_##cvt##_pzo: \
+ return MakeHalfType(Intrinsic::nvvm_##cvt, BuiltinID, E, *this, \
+ {Builder.getTrue()})
+
+ PZO_CVT(ff2f16x2_rn);
+ PZO_CVT(ff2f16x2_rn_relu);
+ PZO_CVT(ff2f16x2_rz);
+ PZO_CVT(ff2f16x2_rz_relu);
+ PZO_CVT(ff2f16x2_rn_satfinite);
+ PZO_CVT(ff2f16x2_rn_relu_satfinite);
+ PZO_CVT(ff2f16x2_rz_satfinite);
+ PZO_CVT(ff2f16x2_rz_relu_satfinite);
+ PZO_CVT(ff2bf16x2_rn);
+ PZO_CVT(ff2bf16x2_rn_relu);
+ PZO_CVT(ff2bf16x2_rz);
+ PZO_CVT(ff2bf16x2_rz_relu);
+ PZO_CVT(ff2bf16x2_rn_satfinite);
+ PZO_CVT(ff2bf16x2_rn_relu_satfinite);
+ PZO_CVT(ff2bf16x2_rz_satfinite);
+ PZO_CVT(ff2bf16x2_rz_relu_satfinite);
+ PZO_CVT(f2f16_rn);
+ PZO_CVT(f2f16_rn_relu);
+ PZO_CVT(f2f16_rz);
+ PZO_CVT(f2f16_rz_relu);
+ PZO_CVT(f2f16_rn_satfinite);
+ PZO_CVT(f2f16_rn_relu_satfinite);
+ PZO_CVT(f2f16_rz_satfinite);
+ PZO_CVT(f2f16_rz_relu_satfinite);
+ PZO_CVT(f2bf16_rn);
+ PZO_CVT(f2bf16_rn_relu);
+ PZO_CVT(f2bf16_rz);
+ PZO_CVT(f2bf16_rz_relu);
+ PZO_CVT(f2bf16_rn_satfinite);
+ PZO_CVT(f2bf16_rn_relu_satfinite);
+ PZO_CVT(f2bf16_rz_satfinite);
+ PZO_CVT(f2bf16_rz_relu_satfinite);
+
+#undef PZO_CVT
+
case NVPTX::BI__nvvm_fma_rn_f16:
return MakeHalfType(Intrinsic::nvvm_fma_rn_f16, BuiltinID, E, *this);
case NVPTX::BI__nvvm_fma_rn_f16x2:
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index 169df03fe63ee..c72634025dbc7 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -58,6 +58,9 @@
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx87 -DPTX=87 \
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX87_SM100a %s
+// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_107f -target-feature +ptx94 -DPTX=94 \
+// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
+// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX94_SM107f %s
// ### The last run to check with the highest SM and PTX version available
// ### to make sure target builtins are still accepted.
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_120a -target-feature +ptx87 -DPTX=87 \
@@ -1028,79 +1031,79 @@ __device__ void nvvm_async_copy(__attribute__((address_space(3))) void* dst, __a
// CHECK-LABEL: nvvm_cvt_sm80
__device__ void nvvm_cvt_sm80() {
#if __CUDA_ARCH__ >= 800
- // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rn(1, 1);
- // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rn_relu(1, 1);
- // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rz(1, 1);
- // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rz_relu(1, 1);
#if PTX >= 81
- // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rn_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rn_relu_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rz_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2bf16x2_rz_relu_satfinite(1, 1);
#endif
- // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rn(1, 1);
- // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rn_relu(1, 1);
- // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rz(1, 1);
- // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rz_relu(1, 1);
#if PTX >= 81
- // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rn_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rn_relu_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rz_satfinite(1, 1);
- // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00)
+ // CHECK_PTX81_SM80: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 false)
__nvvm_ff2f16x2_rz_relu_satfinite(1, 1);
#endif
- // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rn(1);
- // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rn_relu(1);
- // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rz(1);
- // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rz_relu(1);
#if PTX >= 81
- // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rn_satfinite(1);
- // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rn_relu_satfinite(1);
- // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rz_satfinite(1);
- // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2bf16_rz_relu_satfinite(1);
#endif
- // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn(float 1.000000e+00, i1 false)
__nvvm_f2f16_rn(1);
- // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn.relu(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rn.relu(float 1.000000e+00, i1 false)
__nvvm_f2f16_rn_relu(1);
- // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz(float 1.000000e+00, i1 false)
__nvvm_f2f16_rz(1);
- // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz.relu(float 1.000000e+00)
+ // CHECK_PTX70_SM80: call half @llvm.nvvm.f2f16.rz.relu(float 1.000000e+00, i1 false)
__nvvm_f2f16_rz_relu(1);
#if PTX >= 81
- // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2f16_rn_satfinite(1);
- // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.relu.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rn.relu.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2f16_rn_relu_satfinite(1);
- // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2f16_rz_satfinite(1);
- // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.relu.satfinite(float 1.000000e+00)
+ // CHECK_PTX81_SM80: call half @llvm.nvvm.f2f16.rz.relu.satfinite(float 1.000000e+00, i1 false)
__nvvm_f2f16_rz_relu_satfinite(1);
#endif
@@ -1114,6 +1117,80 @@ __device__ void nvvm_cvt_sm80() {
// CHECK: ret void
}
+// CHECK-LABEL: nvvm_cvt_pzo_sm107f
+__device__ void nvvm_cvt_pzo_sm107f() {
+#if (PTX >= 94) && (__CUDA_ARCH__ >= 1070)
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rn_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rn_relu_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rz_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rz_relu_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rn_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rn_relu_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rz_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2bf16x2_rz_relu_satfinite_pzo(1, 1);
+
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rn(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rn_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rn_relu_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rz(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rz_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rz_relu_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rn_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rn_relu_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rz_satfinite_pzo(1, 1);
+ // CHECK_PTX94_SM107f: call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i1 true)
+ __nvvm_ff2f16x2_rz_relu_satfinite_pzo(1, 1);
+
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rn(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rn_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rn.relu(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rn_relu_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rz(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rz_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rz.relu(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rz_relu_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rn_satfinite_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rn_relu_satfinite_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rz_satfinite_pzo(1);
+ // CHECK_PTX94_SM107f: call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2bf16_rz_relu_satfinite_pzo(1);
+
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rn(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rn_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rn.relu(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rn_relu_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rz(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rz_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rz.relu(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rz_relu_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rn.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rn_satfinite_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rn.relu.satfinite(float 1.000000e+00, i1 true)
+ __nvvm_f2f16_rn_relu_satfinite_pzo(1);
+ // CHECK_PTX94_SM107f: call half @llvm.nvvm.f2f16.rz.satfinite(float 1.000000e+00, i1 true)
+ __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);
+#endif
+ // CHECK: ret void
+}
+
// CHECK-LABEL: nvvm_cvt_sm89
__device__ void nvvm_cvt_sm89() {
#if (PTX >= 81) && (__CUDA_ARCH__ >= 890)
@@ -1298,51 +1375,51 @@ __device__ void nvvm_cvt_sm100a_sm103a() {
typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));
typedef char uint8x4 __attribute__((ext_vector_type(4)));
-// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x half> %[[R1]], ptr %r1
-// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x half> %[[R1]], ptr %r1
f16x2 r1 = __nvvm_ff2f16x2_rs(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x half> %[[R2]], ptr %r2
-// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x half> %[[R2]], ptr %r2
f16x2 r2 = __nvvm_ff2f16x2_rs_relu(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x half> %[[R3]], ptr %r3
-// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x half> %[[R3]], ptr %r3
f16x2 r3 = __nvvm_ff2f16x2_rs_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x half> %[[R4]], ptr %r4
-// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x half> %[[R4]], ptr %r4
f16x2 r4 = __nvvm_ff2f16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R5]], ptr %r5
-// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R5]], ptr %r5
bf16x2 r5 = __nvvm_ff2bf16x2_rs(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R6]], ptr %r6
-// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R6]], ptr %r6
bf16x2 r6 = __nvvm_ff2bf16x2_rs_relu(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R7]], ptr %r7
-// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R7]], ptr %r7
bf16x2 r7 = __nvvm_ff2bf16x2_rs_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R8]], ptr %r8
-// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
+// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R8]], ptr %r8
bf16x2 r8 = __nvvm_ff2bf16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 5fe8051d2b039..a85b7de7693d6 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1951,16 +1951,20 @@ let TargetPrefix = "nvvm" in {
foreach relu = ["", "_relu"] in {
foreach satfinite = ["", "_satfinite"] in {
def int_nvvm_ff2bf16x2_ # rnd # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty]>;
+ PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
def int_nvvm_ff2f16x2_ # rnd # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty]>;
+ PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
def int_nvvm_f2bf16_ # rnd # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_bfloat_ty], [llvm_float_ty]>;
+ PureIntrinsic<[llvm_bfloat_ty], [llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
def int_nvvm_f2f16_ # rnd # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_half_ty], [llvm_float_ty]>;
+ PureIntrinsic<[llvm_half_ty], [llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
}
}
}
@@ -1970,10 +1974,10 @@ let TargetPrefix = "nvvm" in {
foreach relu = ["", "_relu"] in {
foreach satfinite = ["", "_satfinite"] in {
def int_nvvm_ff2f16x2_rs # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;
+ PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
def int_nvvm_ff2bf16x2_rs # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;
+ PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
}
}
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 6f3d554907364..0875875611f2c 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -98,6 +98,11 @@ void NVPTXInstPrinter::printCvtMode(const MCInst *MI, int OpNum,
if (Imm & NVPTX::PTXCvtMode::SATFINITE_FLAG)
O << ".satfinite";
return;
+ } else if (Modifier == "pzo") {
+ // PZO flag
+ if (Imm & NVPTX::PTXCvtMode::PZO_FLAG)
+ O << ".pzo";
+ return;
} else if (Modifier == "relu") {
// RELU flag
if (Imm & NVPTX::PTXCvtMode::RELU_FLAG)
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index f7899e72abb20..6897baaede9c4 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -424,7 +424,8 @@ enum CvtMode {
FTZ_FLAG = 0x10,
SAT_FLAG = 0x20,
RELU_FLAG = 0x40,
- SATFINITE_FLAG = 0x80
+ SATFINITE_FLAG = 0x80,
+ PZO_FLAG = 0x100
};
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index ecd5ebea5016a..53d682468f7f5 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -61,6 +61,13 @@ def CvtRZ_SATFINITE : PatLeaf<(i32 0x86)>;
def CvtRP_SATFINITE : PatLeaf<(i32 0x88)>;
def CvtRN_RELU_SATFINITE : PatLeaf<(i32 0xC5)>;
+// Derives the mode constant for one of the modes above with the PZO flag set,
+// so that the PZO variants do not need a PatLeaf of their own.
+class CvtModeWithPZO<PatLeaf mode, int pzo> {
+ int Value = !add(!getdagarg<int>(!head(mode.Fragments), 0),
+ !mul(pzo, 0x100)); // PTXCvtMode::PZO_FLAG
+}
+
def CvtMode : Operand<i32> {
let PrintMethod = "printCvtMode";
}
@@ -236,7 +243,9 @@ def hasSetMaxNRegSupport : PredOr<[SM90a, SM100f, SM110f, SM120f]>;
def hasLdStmatrixBlackwellSupport : PredOr<[SM100f, SM110f, SM120f]>;
def hasConvertWithStochasticRounding
- : PredAnd<[PTX87, PredOr<[SM100a, SM103a]>]>;
+ : PredAnd<[PTX87, PredOr<[SM100a, SM103a, SM107a]>]>;
+
+def hasConvertWithPZOSupport : PredOr<[hasRubinFamilySupport]>;
//===----------------------------------------------------------------------===//
@@ -648,7 +657,7 @@ let hasSideEffects = false in {
def _f32 :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:ftz}${mode:relu}${mode:sat}." # ToType # ".f32">,
+ "cvt${mode:base}${mode:ftz}${mode:relu}${mode:sat}${mode:pzo}." # ToType # ".f32">,
Requires<!if(!eq(ToType, "bf16"),
// f32->bf16 was introduced early.
[SM80],
@@ -676,7 +685,7 @@ let hasSideEffects = false in {
def _f32_sf :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.satfinite." # ToName # ".f32">;
+ "cvt${mode:base}${mode:relu}.satfinite${mode:pzo}." # ToName # ".f32">;
}
defm CVT_bf16 : CVT_FROM_FLOAT_SATFINITE<"bf16", B16>;
defm CVT_f16 : CVT_FROM_FLOAT_SATFINITE<"f16", B16>;
@@ -694,13 +703,13 @@ let hasSideEffects = false in {
def _f32 :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}." # FromName # ".f32">,
+ "cvt${mode:base}${mode:relu}${mode:pzo}." # FromName # ".f32">,
Requires<[SM80]>;
def _f32_sf :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2), (ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">;
+ "cvt${mode:base}${mode:relu}.satfinite${mode:pzo}." # FromName # ".f32">;
}
defm CVT_f16x2 : CVT_FROM_FLOAT_V2_SM80<"f16x2", B32>;
@@ -711,13 +720,13 @@ let hasSideEffects = false in {
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2, B32:$src3),
(ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}." # FromName # ".f32">;
+ "cvt${mode:base}${mode:relu}${mode:pzo}." # FromName # ".f32">;
def _f32_rs_sf :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2, B32:$src3),
(ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">;
+ "cvt${mode:base}${mode:relu}.satfinite${mode:pzo}." # FromName # ".f32">;
}
defm CVT_f16x2 : CVT_FROM_FLOAT_V2_RS<"f16x2", B32>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 356cd9bf88dd1..b7801b3496bcb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2473,68 +2473,74 @@ 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)>;
-def : Pat<(int_nvvm_ff2bf16x2_rn f32:$a, f32:$b), (CVT_bf16x2_f32 $a, $b, CvtRN)>;
-def : Pat<(int_nvvm_ff2bf16x2_rn_relu f32:$a, f32:$b), (CVT_bf16x2_f32 $a, $b, CvtRN_RELU)>;
-def : Pat<(int_nvvm_ff2bf16x2_rz f32:$a, f32:$b), (CVT_bf16x2_f32 $a, $b, CvtRZ)>;
-def : Pat<(int_nvvm_ff2bf16x2_rz_relu f32:$a, f32:$b), (CVT_bf16x2_f32 $a, $b, CvtRZ_RELU)>;
-let Predicates = [PTX81, SM80] in {
- def : Pat<(int_nvvm_ff2bf16x2_rn_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff2bf16x2_rn_relu_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRN_RELU)>;
- def : Pat<(int_nvvm_ff2bf16x2_rz_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRZ)>;
- def : Pat<(int_nvvm_ff2bf16x2_rz_relu_satfinite f32:$a, f32:$b), (CVT_bf16x2_f32_sf $a, $b, CvtRZ_RELU)>;
+// Utility to look up base and satfinite variants of a conversion intrinsic
+class ToCvtIntrinsics<string suffix> {
+ Intrinsic base = !cast<Intrinsic>("int_nvvm_" # suffix);
+ Intrinsic sf = !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite");
}
-let Predicates = [hasConvertWithStochasticRounding] in {
-def : Pat<(int_nvvm_ff2bf16x2_rs f32:$a, f32:$b, i32:$c),
- (CVT_bf16x2_f32_rs $a, $b, $c, CvtRS)>;
-def : Pat<(int_nvvm_ff2bf16x2_rs_relu f32:$a, f32:$b, i32:$c),
- (CVT_bf16x2_f32_rs $a, $b, $c, CvtRS_RELU)>;
-def : Pat<(int_nvvm_ff2bf16x2_rs_satfinite f32:$a, f32:$b, i32:$c),
- (CVT_bf16x2_f32_rs_sf $a, $b, $c, CvtRS)>;
-def : Pat<(int_nvvm_ff2bf16x2_rs_relu_satfinite f32:$a, f32:$b, i32:$c),
- (CVT_bf16x2_f32_rs_sf $a, $b, $c, CvtRS_RELU)>;
-}
-
-def : Pat<(int_nvvm_ff2f16x2_rn f32:$a, f32:$b), (CVT_f16x2_f32 $a, $b, CvtRN)>;
-def : Pat<(int_nvvm_ff2f16x2_rn_relu f32:$a, f32:$b), (CVT_f16x2_f32 $a, $b, CvtRN_RELU)>;
-def : Pat<(int_nvvm_ff2f16x2_rz f32:$a, f32:$b), (CVT_f16x2_f32 $a, $b, CvtRZ)>;
-def : Pat<(int_nvvm_ff2f16x2_rz_relu f32:$a, f32:$b), (CVT_f16x2_f32 $a, $b, CvtRZ_RELU)>;
-let Predicates = [PTX81, SM80] in {
- def : Pat<(int_nvvm_ff2f16x2_rn_satfinite f32:$a, f32:$b), (CVT_f16x2_f32_sf $a, $b, CvtRN)>;
- def : Pat<(int_nvvm_ff2f16x2_rn_relu_satfinite f32:$a, f32:$b), (CVT_f16x2_f32_sf $a, $b, CvtRN_RELU)>;
- def : Pat<(int_nvvm_ff2f16x2_rz_satfinite f32:$a, f32:$b), (CVT_f16x2_f32_sf $a, $b, CvtRZ)>;
- def : Pat<(int_nvvm_ff2f16x2_rz_relu_satfinite f32:$a, f32:$b), (CVT_f16x2_f32_sf $a, $b, CvtRZ_RELU)>;
+
+foreach rnd = ["rn", "rz"] in {
+ foreach relu = ["", "_relu"] in {
+ defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
+
+ foreach pzo = [0, 1] in {
+ // An i1 immediate of one is matched as the sign-extended value -1.
+ defvar PZO = !if(pzo, -1, 0);
+ defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
+ defvar PZOPreds =
+ !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>);
+ defvar SFPreds = !if(pzo, [hasConvertWithPZOSupport]<Predicate>,
+ [PTX81, SM80]<Predicate>);
+
+ defvar FF2BF16X2 = ToCvtIntrinsics<"ff2bf16x2_" # rnd # relu>;
+ defvar FF2F16X2 = ToCvtIntrinsics<"ff2f16x2_" # rnd # relu>;
+ defvar F2BF16 = ToCvtIntrinsics<"f2bf16_" # rnd # relu>;
+ defvar F2F16 = ToCvtIntrinsics<"f2f16_" # rnd # relu>;
+
+ def : Pat<(FF2BF16X2.base f32:$a, f32:$b, PZO),
+ (CVT_bf16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
+ def : Pat<(FF2F16X2.base f32:$a, f32:$b, PZO),
+ (CVT_f16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
+ def : Pat<(F2BF16.base f32:$a, PZO), (CVT_bf16_f32 $a, Mode)>,
+ Requires<PZOPreds>;
+ def : Pat<(F2F16.base f32:$a, PZO), (CVT_f16_f32 $a, Mode)>,
+ Requires<PZOPreds>;
+
+ def : Pat<(FF2BF16X2.sf f32:$a, f32:$b, PZO),
+ (CVT_bf16x2_f32_sf $a, $b, Mode)>, Requires<SFPreds>;
+ def : Pat<(FF2F16X2.sf f32:$a, f32:$b, PZO),
+ (CVT_f16x2_f32_sf $a, $b, Mode)>, Requires<SFPreds>;
+ def : Pat<(F2BF16.sf f32:$a, PZO), (CVT_bf16_f32_sf $a, Mode)>,
+ Requires<SFPreds>;
+ def : Pat<(F2F16.sf f32:$a, PZO), (CVT_f16_f32_sf $a, Mode)>,
+ Requires<SFPreds>;
+ }
+ }
}
-let Predicates = [hasConvertWithStochasticRounding] in {
-def : Pat<(int_nvvm_ff2f16x2_rs f32:$a, f32:$b, i32:$c),
- (CVT_f16x2_f32_rs $a, $b, $c, CvtRS)>;
-def : Pat<(int_nvvm_ff2f16x2_rs_relu f32:$a, f32:$b, i32:$c),
- (CVT_f16x2_f32_rs $a, $b, $c, CvtRS_RELU)>;
-def : Pat<(int_nvvm_ff2f16x2_rs_satfinite f32:$a, f32:$b, i32:$c),
- (CVT_f16x2_f32_rs_sf $a, $b, $c, CvtRS)>;
-def : Pat<(int_nvvm_ff2f16x2_rs_relu_satfinite f32:$a, f32:$b, i32:$c),
- (CVT_f16x2_f32_rs_sf $a, $b, $c, CvtRS_RELU)>;
-}
-def : Pat<(int_nvvm_f2bf16_rn f32:$a), (CVT_bf16_f32 $a, CvtRN)>;
-def : Pat<(int_nvvm_f2bf16_rn_relu f32:$a), (CVT_bf16_f32 $a, CvtRN_RELU)>;
-def : Pat<(int_nvvm_f2bf16_rz f32:$a), (CVT_bf16_f32 $a, CvtRZ)>;
-def : Pat<(int_nvvm_f2bf16_rz_relu f32:$a), (CVT_bf16_f32 $a, CvtRZ_RELU)>;
-let Predicates = [PTX81, SM80] in {
- def : Pat<(int_nvvm_f2bf16_rz_satfinite f32:$a), (CVT_bf16_f32_sf $a, CvtRZ)>;
- def : Pat<(int_nvvm_f2bf16_rz_relu_satfinite f32:$a), (CVT_bf16_f32_sf $a, CvtRZ_RELU)>;
- def : Pat<(int_nvvm_f2bf16_rn_satfinite f32:$a), (CVT_bf16_f32_sf $a, CvtRN)>;
- def : Pat<(int_nvvm_f2bf16_rn_relu_satfinite f32:$a), (CVT_bf16_f32_sf $a, CvtRN_RELU)>;
-}
-
-def : Pat<(int_nvvm_f2f16_rn f32:$a), (CVT_f16_f32 $a, CvtRN)>;
-def : Pat<(int_nvvm_f2f16_rn_relu f32:$a), (CVT_f16_f32 $a, CvtRN_RELU)>;
-def : Pat<(int_nvvm_f2f16_rz f32:$a), (CVT_f16_f32 $a, CvtRZ)>;
-def : Pat<(int_nvvm_f2f16_rz_relu f32:$a), (CVT_f16_f32 $a, CvtRZ_RELU)>;
-let Predicates = [PTX81, SM80] in {
- def : Pat<(int_nvvm_f2f16_rz_satfinite f32:$a), (CVT_f16_f32_sf $a, CvtRZ)>;
- def : Pat<(int_nvvm_f2f16_rz_relu_satfinite f32:$a), (CVT_f16_f32_sf $a, CvtRZ_RELU)>;
- def : Pat<(int_nvvm_f2f16_rn_satfinite f32:$a), (CVT_f16_f32_sf $a, CvtRN)>;
- def : Pat<(int_nvvm_f2f16_rn_relu_satfinite f32:$a), (CVT_f16_f32_sf $a, CvtRN_RELU)>;
+foreach relu = ["", "_relu"] in {
+ defvar BaseMode = !cast<PatLeaf>("CvtRS" # !toupper(relu));
+
+ foreach pzo = [0, 1] in {
+ defvar PZO = !if(pzo, -1, 0);
+ defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
+ defvar Preds =
+ !listconcat(
+ [hasConvertWithStochasticRounding]<Predicate>,
+ !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>));
+
+ defvar BF16X2 = ToCvtIntrinsics<"ff2bf16x2_rs" # relu>;
+ defvar F16X2 = ToCvtIntrinsics<"ff2f16x2_rs" # relu>;
+
+ def : Pat<(BF16X2.base f32:$a, f32:$b, i32:$c, PZO),
+ (CVT_bf16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>;
+ def : Pat<(BF16X2.sf f32:$a, f32:$b, i32:$c, PZO),
+ (CVT_bf16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>;
+ def : Pat<(F16X2.base f32:$a, f32:$b, i32:$c, PZO),
+ (CVT_f16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>;
+ def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c, PZO),
+ (CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>;
+ }
}
def : Pat<(int_nvvm_lohi_i2d i32:$a, i32:$b), (V2I32toI64 $a, $b)>;
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
new file mode 100644
index 0000000000000..6861daf25aacf
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
@@ -0,0 +1,494 @@
+; 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 %}
+
+; Conversions to fp16 types with the preserve-zero-order modifier. The RS
+; (stochastic rounding) variants need sm_107a and live in convert-fp16-rs-pzo.ll.
+
+; F16X2 conversions
+
+define <2 x half> @cvt_rn_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rz_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rn_relu_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.relu.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rn_sf_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rn_relu_sf_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rz_relu_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.relu.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rz_sf_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rz_relu_sf_pzo_f16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.relu.satfinite.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x half> %val
+}
+
+; BF16X2 conversions
+
+define <2 x bfloat> @cvt_rn_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rz_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rn_relu_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.relu.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rn_sf_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rn_relu_sf_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rn_relu_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rn_relu_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.relu.satfinite.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rz_relu_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.relu.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rz_sf_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rz_relu_sf_pzo_bf16x2_f32(float %f1, float %f2) {
+; CHECK-LABEL: cvt_rz_relu_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rz_relu_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.relu.satfinite.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float %f1, float %f2, i1 true)
+ ret <2 x bfloat> %val
+}
+
+; F16 conversions
+
+define half @cvt_rn_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rn.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rn(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rn_relu_sf_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_relu_sf_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_sf_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rn.relu.satfinite.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rn.relu.satfinite(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rn_relu_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_relu_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rn.relu.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rn.relu(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rn_sf_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_sf_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_sf_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rn.satfinite(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rz_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rz.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rz(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rz_relu_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_relu_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rz.relu.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rz.relu(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rz_sf_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_sf_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_sf_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rz.satfinite(float %f1, i1 true)
+ ret half %val
+}
+
+define half @cvt_rz_relu_sf_pzo_f16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_relu_sf_pzo_f16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_sf_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rz.relu.satfinite.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rz.relu.satfinite(float %f1, i1 true)
+ ret half %val
+}
+
+; BF16 conversions
+
+define bfloat @cvt_rn_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rn.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rn(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rn_relu_sf_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_relu_sf_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_sf_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rn.relu.satfinite.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rn_relu_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_relu_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_relu_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rn.relu.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rn.relu(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rn_sf_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rn_sf_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rn_sf_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rn.satfinite.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rz_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rz.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rz(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rz_relu_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_relu_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rz.relu.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rz.relu(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rz_sf_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_sf_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_sf_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rz.satfinite.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float %f1, i1 true)
+ ret bfloat %val
+}
+
+define bfloat @cvt_rz_relu_sf_pzo_bf16_f32(float %f1) {
+; CHECK-LABEL: cvt_rz_relu_sf_pzo_bf16_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rz_relu_sf_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rz.relu.satfinite.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float %f1, i1 true)
+ ret bfloat %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
new file mode 100644
index 0000000000000..58eacc36bcf55
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
@@ -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_107a -mattr=+ptx94 | FileCheck %s
+; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %}
+
+; F16X2 conversions
+
+define <2 x half> @cvt_rs_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_f16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rs_relu_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_relu_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_f16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.relu.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rs_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_f16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x half> %val
+}
+
+define <2 x half> @cvt_rs_relu_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_relu_sf_pzo_f16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_f16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x half> %val
+}
+
+; BF16X2 conversions
+
+define <2 x bfloat> @cvt_rs_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_bf16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rs_relu_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_relu_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_bf16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.relu.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rs_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_bf16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x bfloat> %val
+}
+
+define <2 x bfloat> @cvt_rs_relu_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
+; CHECK-LABEL: cvt_rs_relu_sf_pzo_bf16x2_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_2];
+; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
+; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
+ ret <2 x bfloat> %val
+}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index aa45bba341240..0a9ba990a41a8 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5625,6 +5625,9 @@ ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
if (op.getRandomBits())
args.push_back(mt.lookupValue(op.getRandomBits()));
+ // TODO: Add support for PZO modifier
+ args.push_back(builder.getInt1(false));
+
switch (op.getRnd()) {
case FPRoundingMode::RN:
return {rndRNIds[idx], std::move(args)};
@@ -5673,6 +5676,9 @@ ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
if (op.getRandomBits())
args.push_back(mt.lookupValue(op.getRandomBits()));
+ // TODO: Add support for PZO modifier
+ args.push_back(builder.getInt1(false));
+
switch (op.getRnd()) {
case FPRoundingMode::RN:
return {rndRNIds[idx], std::move(args)};
diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
index 0e51c921855fb..ddb6defe270bb 100644
--- a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
@@ -2,13 +2,13 @@
// CHECK-LABEL: @convert_f32x2_to_f16x2_rn
llvm.func @convert_f32x2_to_f16x2_rn(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rn> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rn> relu = true : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : vector<2xf16>
llvm.return
@@ -16,13 +16,13 @@ llvm.func @convert_f32x2_to_f16x2_rn(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_f16x2_rz
llvm.func @convert_f32x2_to_f16x2_rz(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rz> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rz> sat = <satfinite> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rz> relu = true : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB rnd = <rz> sat = <satfinite> relu = true : vector<2xf16>
llvm.return
@@ -30,13 +30,13 @@ llvm.func @convert_f32x2_to_f16x2_rz(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_f16x2_rs_stochastic
llvm.func @convert_f32x2_to_f16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits : i32) {
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> relu = true : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> relu = true : vector<2xf16>
llvm.return
@@ -46,13 +46,13 @@ llvm.func @convert_f32x2_to_f16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits
// CHECK-LABEL: @convert_f32x2_to_bf16x2_rn
llvm.func @convert_f32x2_to_bf16x2_rn(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rn> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rn> sat = <satfinite> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rn> relu = true : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rn> sat = <satfinite> relu = true : vector<2xbf16>
llvm.return
@@ -60,13 +60,13 @@ llvm.func @convert_f32x2_to_bf16x2_rn(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_bf16x2_rz
llvm.func @convert_f32x2_to_bf16x2_rz(%srcA : f32, %srcB : f32) {
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float %{{.*}}, float %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rz> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rz> sat = <satfinite> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float %{{.*}}, float %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rz> relu = true : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float %{{.*}}, float %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float %{{.*}}, float %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB rnd = <rz> sat = <satfinite> relu = true : vector<2xbf16>
llvm.return
@@ -74,13 +74,13 @@ llvm.func @convert_f32x2_to_bf16x2_rz(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_bf16x2_rs_stochastic
llvm.func @convert_f32x2_to_bf16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits : i32) {
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res1 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res2 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> relu = true : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res3 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
%res4 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> relu = true : vector<2xbf16>
llvm.return
More information about the cfe-commits
mailing list