[clang] [llvm] [mlir] [clang][NVPTX] Add support for pzo in f32 to f16/bf16 conversions (PR #214667)
Srinivasa Ravi via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 7 01:50:32 PDT 2026
https://github.com/Wolfram70 created https://github.com/llvm/llvm-project/pull/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.
>From e27d5e72397c5a4f0d052f83ae962e11cb63802d Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Thu, 6 Aug 2026 05:48:02 +0000
Subject: [PATCH] add f32x2 to fp16x2 pzo conversion intrinsics
---
clang/include/clang/Basic/BuiltinsNVPTX.td | 34 ++-
clang/lib/CodeGen/CGBuiltin.cpp | 18 ++
clang/lib/CodeGen/CGBuiltin.h | 5 +
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 49 +++-
clang/test/CodeGen/builtins-nvptx.c | 173 ++++++++++----
llvm/include/llvm/IR/IntrinsicsNVVM.td | 12 +-
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 5 +
llvm/lib/Target/NVPTX/NVPTX.h | 3 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 19 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 134 ++++++-----
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 4 +
llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll | 221 ++++++++++++++++++
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 6 +
.../Target/LLVMIR/nvvm/convert_fp16x2.mlir | 48 ++--
14 files changed, 582 insertions(+), 149 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index bcfe1f9bf8572..3475b721e95e0 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -80,7 +80,7 @@ multiclass SM_Instantiate<list<int> gpu_list> {
}
}
-defm SM : SM_Instantiate<[121, 120, 110, 103, 101, 100, 90, 89, 88, 87, 86, 80, 75, 72, 70, 62, 61, 60, 53]>;
+defm SM : SM_Instantiate<[121, 120, 110, 107, 103, 101, 100, 90, 89, 88, 87, 86, 80, 75, 72, 70, 62, 61, 60, 53]>;
class PTXFeatures {
string Features;
@@ -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 4c1318f6543c1..8af299e8e3f1d 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -253,6 +253,22 @@ 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();
+ if (Args.size() == FTy->getNumParams())
+ return;
+
+ auto [FirstDefault, Defaults] =
+ Intrinsic::getAllDefaultArgValues(F->getIntrinsicID());
+ for (unsigned I = Args.size(), E = FTy->getNumParams(); I != E; ++I) {
+ if (I < FirstDefault || I - FirstDefault >= Defaults.size())
+ break;
+ Args.push_back(llvm::ConstantInt::get(FTy->getParamType(I),
+ Defaults[I - FirstDefault]));
+ }
+}
+
/// Emit the conversions required to turn the given value into an
/// integer of the given size.
Value *EmitToInt(CodeGenFunction &CGF, llvm::Value *V,
@@ -7114,6 +7130,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..3a344f190712f 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,43 @@ 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 87be7b46aad8e..0ff01f1a82b8c 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -55,6 +55,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 \
@@ -1025,79 +1028,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
@@ -1111,6 +1114,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)
@@ -1295,51 +1372,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 fc04a5dc64e78..515fd35ed274e 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1724,16 +1724,16 @@ 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>>]>;
}
}
}
@@ -1743,10 +1743,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 fe9e2fa260811..f3f50eacb7c33 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -99,6 +99,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 365ac038b79cc..2f5b8fbc9497f 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -308,7 +308,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 83060bb906157..dbbeb2c5c5f65 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -59,6 +59,13 @@ def CvtNONE_SATFINITE : PatLeaf<(i32 0x80)>;
def CvtRN_SATFINITE : PatLeaf<(i32 0x85)>;
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";
}
@@ -567,7 +574,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.
[hasPTX<70>, SM80],
@@ -595,7 +602,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>;
@@ -613,13 +620,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<[hasPTX<70>, 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>;
@@ -630,13 +637,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 baa4be7c0d708..22ba237191ae8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2117,68 +2117,82 @@ 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 = [hasPTX<81>, 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)>;
-}
-let Predicates = [callSubtarget<"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 = [hasPTX<81>, 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, [callSubtarget<"hasConvertWithPZOSupport">]<Predicate>,
+ []<Predicate>);
+
+ defvar FF2BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_" # rnd # relu);
+ defvar FF2F16X2 = !cast<Intrinsic>("int_nvvm_ff2f16x2_" # rnd # relu);
+ defvar F2BF16 = !cast<Intrinsic>("int_nvvm_f2bf16_" # rnd # relu);
+ defvar F2F16 = !cast<Intrinsic>("int_nvvm_f2f16_" # rnd # relu);
+
+ def : Pat<(FF2BF16X2 f32:$a, f32:$b, PZO),
+ (CVT_bf16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
+ def : Pat<(FF2F16X2 f32:$a, f32:$b, PZO),
+ (CVT_f16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
+ def : Pat<(F2BF16 f32:$a, PZO), (CVT_bf16_f32 $a, Mode)>,
+ Requires<PZOPreds>;
+ def : Pat<(F2F16 f32:$a, PZO), (CVT_f16_f32 $a, Mode)>,
+ Requires<PZOPreds>;
+
+ defvar FF2BF16X2_SF =
+ !cast<Intrinsic>("int_nvvm_ff2bf16x2_" # rnd # relu # "_satfinite");
+ defvar FF2F16X2_SF =
+ !cast<Intrinsic>("int_nvvm_ff2f16x2_" # rnd # relu # "_satfinite");
+ defvar F2BF16_SF =
+ !cast<Intrinsic>("int_nvvm_f2bf16_" # rnd # relu # "_satfinite");
+ defvar F2F16_SF =
+ !cast<Intrinsic>("int_nvvm_f2f16_" # rnd # relu # "_satfinite");
+ defvar SFPreds = !listconcat([hasPTX<81>, SM80], 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 = [callSubtarget<"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 = [hasPTX<81>, 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 = [hasPTX<81>, 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(
+ [callSubtarget<"hasConvertWithStochasticRounding">]<Predicate>,
+ !if(pzo, [callSubtarget<"hasConvertWithPZOSupport">]<Predicate>,
+ []<Predicate>));
+
+ defvar BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_rs" # relu);
+ defvar BF16X2_SF =
+ !cast<Intrinsic>("int_nvvm_ff2bf16x2_rs" # relu # "_satfinite");
+ defvar F16X2 = !cast<Intrinsic>("int_nvvm_ff2f16x2_rs" # relu);
+ defvar F16X2_SF =
+ !cast<Intrinsic>("int_nvvm_ff2f16x2_rs" # relu # "_satfinite");
+
+ def : Pat<(BF16X2 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 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/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 21b4cf6169f4f..acf9a0e604977 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -294,6 +294,10 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
return hasAnyFeature({NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f});
}
+ bool hasConvertWithPZOSupport() const {
+ return hasPTXWithFamilySMs(94, {107});
+ }
+
// Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
// terminates a basic block. Instead, it would assume that control flow
// continued to the next instruction. The next instruction could be in the
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
new file mode 100644
index 0000000000000..6dc4dd10e6352
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
@@ -0,0 +1,221 @@
+; 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 %}
+
+; 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.b32 %r1, [cvt_rn_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.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.b32 %r1, [cvt_rz_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_pzo_f16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.pzo.f16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.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.b32 %r1, [cvt_rn_relu_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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.b32 %r1, [cvt_rn_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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.b32 %r1, [cvt_rn_relu_sf_pzo_f16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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
+}
+
+; 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.b32 %r1, [cvt_rn_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rn.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.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.b32 %r1, [cvt_rz_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_pzo_bf16x2_f32_param_1];
+; CHECK-NEXT: cvt.rz.pzo.bf16x2.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.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.b32 %r1, [cvt_rn_relu_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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.b32 %r1, [cvt_rn_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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.b32 %r1, [cvt_rn_relu_sf_pzo_bf16x2_f32_param_0];
+; CHECK-NEXT: ld.param.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.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
+}
+
+; 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.b32 %r1, [cvt_rn_pzo_f16_f32_param_0];
+; CHECK-NEXT: cvt.rn.pzo.f16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param.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.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.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call half @llvm.nvvm.f2f16.rn.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.b32 %r1, [cvt_rn_pzo_bf16_f32_param_0];
+; CHECK-NEXT: cvt.rn.pzo.bf16.f32 %rs1, %r1;
+; CHECK-NEXT: st.param.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.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.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
+ %val = call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float %f1, i1 true)
+ ret bfloat %val
+}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 67c4be9670c91..21ea548173e35 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5085,6 +5085,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)};
@@ -5133,6 +5136,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 a4bece83f832a..5a0bb19198a77 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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rn>, relu = true, sat = #nvvm.sat_mode<satfinite>} : 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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rz>, relu = true, sat = #nvvm.sat_mode<satfinite>} : 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, %rbits {rnd = #nvvm.fp_rnd_mode<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, %rbits {relu = true, rnd = #nvvm.fp_rnd_mode<rs>} : 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, %rbits {rnd = #nvvm.fp_rnd_mode<rs>, sat = #nvvm.sat_mode<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, %rbits {relu = true, rnd = #nvvm.fp_rnd_mode<rs>, sat = #nvvm.sat_mode<satfinite>} : 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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rn>, relu = true, sat = #nvvm.sat_mode<satfinite>} : 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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<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 = #nvvm.fp_rnd_mode<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 = #nvvm.fp_rnd_mode<rz>, relu = true, sat = #nvvm.sat_mode<satfinite>} : 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, %rbits {rnd = #nvvm.fp_rnd_mode<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, %rbits {relu = true, rnd = #nvvm.fp_rnd_mode<rs>} : 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, %rbits {rnd = #nvvm.fp_rnd_mode<rs>, sat = #nvvm.sat_mode<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, %rbits {relu = true, rnd = #nvvm.fp_rnd_mode<rs>, sat = #nvvm.sat_mode<satfinite>} : vector<2xbf16>
llvm.return
More information about the llvm-commits
mailing list