[Mlir-commits] [clang] [llvm] [mlir] [clang][NVPTX] Add support for pzo in f32 to f16/bf16 conversions (PR #214667)
Srinivasa Ravi
llvmlistbot at llvm.org
Wed Aug 26 02:20:43 PDT 2026
https://github.com/Wolfram70 updated https://github.com/llvm/llvm-project/pull/214667
>From 3dac2d7525f136010833c7322f9d8499571c7ce2 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 1/6] 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 07379056f2f0a..d552e0b638e09 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,
@@ -6908,6 +6924,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 69905899f3007..e8976337106bd 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1849,16 +1849,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>>]>;
}
}
}
@@ -1868,10 +1868,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 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 4c3c7792f277b..23a71f291dbc6 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";
}
@@ -648,7 +655,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 +683,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 +701,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 +718,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 5e01b251ff9fd..91110b2f5038b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2280,68 +2280,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 = [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)>;
-}
-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, [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 = [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(
+ [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 b2f24c1d5d2b3..1bd111c8d02ec 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -177,6 +177,10 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
return hasTensormapReplaceSupport();
}
+ 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 4dddb33fd050e..146f1400a25bb 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5414,6 +5414,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)};
@@ -5462,6 +5465,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
>From aef93cdac4f743c80d24729c21568a22e53751de Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Fri, 7 Aug 2026 10:53:10 +0000
Subject: [PATCH 2/6] fix formatting
---
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 68 ++++++++++++----------
1 file changed, 36 insertions(+), 32 deletions(-)
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index 3a344f190712f..e3ef3ae488a0c 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -981,43 +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);
+
+ 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:
>From d057d9b85977301d073b22f16117075adbbf233e Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Fri, 7 Aug 2026 11:32:38 +0000
Subject: [PATCH 3/6] update after rebase
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 6 ++----
llvm/lib/Target/NVPTX/NVPTXSubtarget.h | 4 ----
2 files changed, 2 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 91110b2f5038b..df11e4cc560e2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2289,8 +2289,7 @@ foreach rnd = ["rn", "rz"] in {
defvar PZO = !if(pzo, -1, 0);
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar PZOPreds =
- !if(pzo, [callSubtarget<"hasConvertWithPZOSupport">]<Predicate>,
- []<Predicate>);
+ !if(pzo, [SM107f]<Predicate>, []<Predicate>);
defvar FF2BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_" # rnd # relu);
defvar FF2F16X2 = !cast<Intrinsic>("int_nvvm_ff2f16x2_" # rnd # relu);
@@ -2337,8 +2336,7 @@ foreach relu = ["", "_relu"] in {
defvar Preds =
!listconcat(
[callSubtarget<"hasConvertWithStochasticRounding">]<Predicate>,
- !if(pzo, [callSubtarget<"hasConvertWithPZOSupport">]<Predicate>,
- []<Predicate>));
+ !if(pzo, [SM107f]<Predicate>, []<Predicate>));
defvar BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_rs" # relu);
defvar BF16X2_SF =
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 1bd111c8d02ec..b2f24c1d5d2b3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -177,10 +177,6 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
return hasTensormapReplaceSupport();
}
- 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
>From 415f621f52b73d8ea8ece7366d4bee8e207d662d Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Tue, 25 Aug 2026 07:07:21 +0000
Subject: [PATCH 4/6] update predicates and tests
---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 2 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 9 ++-
llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll | 76 +++++++++----------
3 files changed, 45 insertions(+), 42 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 23a71f291dbc6..5123239f10ea3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -245,6 +245,8 @@ def hasLdStmatrixBlackwellSupport : PredOr<[SM100f, SM110f, SM120f]>;
def hasConvertWithStochasticRounding
: PredAnd<[PTX87, PredOr<[SM100a, SM103a]>]>;
+def hasConvertWithPZOSupport : PredOr<[SM107f]>;
+
//===----------------------------------------------------------------------===//
// Some Common Instruction Class Templates
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index df11e4cc560e2..dd7cb81a227c4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2289,7 +2289,7 @@ foreach rnd = ["rn", "rz"] in {
defvar PZO = !if(pzo, -1, 0);
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar PZOPreds =
- !if(pzo, [SM107f]<Predicate>, []<Predicate>);
+ !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>);
defvar FF2BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_" # rnd # relu);
defvar FF2F16X2 = !cast<Intrinsic>("int_nvvm_ff2f16x2_" # rnd # relu);
@@ -2313,7 +2313,8 @@ foreach rnd = ["rn", "rz"] in {
!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);
+ defvar SFPreds = !if(pzo, [hasConvertWithPZOSupport]<Predicate>,
+ [PTX81, SM80]<Predicate>);
def : Pat<(FF2BF16X2_SF f32:$a, f32:$b, PZO),
(CVT_bf16x2_f32_sf $a, $b, Mode)>, Requires<SFPreds>;
@@ -2335,8 +2336,8 @@ foreach relu = ["", "_relu"] in {
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar Preds =
!listconcat(
- [callSubtarget<"hasConvertWithStochasticRounding">]<Predicate>,
- !if(pzo, [SM107f]<Predicate>, []<Predicate>));
+ [hasConvertWithStochasticRounding]<Predicate>,
+ !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>));
defvar BF16X2 = !cast<Intrinsic>("int_nvvm_ff2bf16x2_rs" # relu);
defvar BF16X2_SF =
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
index 6dc4dd10e6352..fcb1a7d06d554 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll
@@ -10,10 +10,10 @@ define <2 x half> @cvt_rn_pzo_f16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -25,10 +25,10 @@ define <2 x half> @cvt_rz_pzo_f16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -40,10 +40,10 @@ define <2 x half> @cvt_rn_relu_pzo_f16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -55,10 +55,10 @@ define <2 x half> @cvt_rn_sf_pzo_f16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -70,10 +70,10 @@ define <2 x half> @cvt_rn_relu_sf_pzo_f16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -87,10 +87,10 @@ define <2 x bfloat> @cvt_rn_pzo_bf16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -102,10 +102,10 @@ define <2 x bfloat> @cvt_rz_pzo_bf16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -117,10 +117,10 @@ define <2 x bfloat> @cvt_rn_relu_pzo_bf16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -132,10 +132,10 @@ define <2 x bfloat> @cvt_rn_sf_pzo_bf16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -147,10 +147,10 @@ define <2 x bfloat> @cvt_rn_relu_sf_pzo_bf16x2_f32(float %f1, float %f2) {
; 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: 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.b32 [func_retval0], %r3;
+; 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
@@ -165,9 +165,9 @@ define half @cvt_rn_pzo_f16_f32(float %f1) {
; 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: 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.b16 [func_retval0], %rs1;
+; 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
@@ -180,9 +180,9 @@ define half @cvt_rn_relu_sf_pzo_f16_f32(float %f1) {
; 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: 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.b16 [func_retval0], %rs1;
+; 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
@@ -197,9 +197,9 @@ define bfloat @cvt_rn_pzo_bf16_f32(float %f1) {
; 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: 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.b16 [func_retval0], %rs1;
+; 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
@@ -212,9 +212,9 @@ define bfloat @cvt_rn_relu_sf_pzo_bf16_f32(float %f1) {
; 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: 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.b16 [func_retval0], %rs1;
+; 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
>From 3c0854efeda91e187648eea2ca59b02157833398 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Tue, 25 Aug 2026 15:03:28 +0000
Subject: [PATCH 5/6] address comments
---
clang/lib/CodeGen/CGBuiltin.cpp | 7 +--
llvm/include/llvm/IR/IntrinsicsNVVM.td | 12 +++--
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 59 +++++++++++-------------
3 files changed, 38 insertions(+), 40 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index d552e0b638e09..784071505ee2f 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -262,10 +262,11 @@ void appendDefaultIntrinsicArgs(SmallVectorImpl<llvm::Value *> &Args,
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())
+ unsigned DefaultIdx = I - FirstDefault;
+ if (DefaultIdx < 0 || DefaultIdx >= Defaults.size())
break;
- Args.push_back(llvm::ConstantInt::get(FTy->getParamType(I),
- Defaults[I - FirstDefault]));
+ Args.push_back(
+ llvm::ConstantInt::get(FTy->getParamType(I), Defaults[DefaultIdx]));
}
}
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index e8976337106bd..83e706de37af4 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1849,16 +1849,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, llvm_i1_ty], [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ 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, llvm_i1_ty], [ImmArg<ArgIndex<2>, DefaultValue<0>>]>;
+ 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, llvm_i1_ty], [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
+ 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, llvm_i1_ty], [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
+ PureIntrinsic<[llvm_half_ty], [llvm_float_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>]>;
}
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index dd7cb81a227c4..06e2e11be9da9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2280,6 +2280,12 @@ def : Pat<(int_nvvm_ui2f_rz i32:$a), (CVT_f32_u32 $a, CvtRZ)>;
def : Pat<(int_nvvm_ui2f_rm i32:$a), (CVT_f32_u32 $a, CvtRM)>;
def : Pat<(int_nvvm_ui2f_rp i32:$a), (CVT_f32_u32 $a, CvtRP)>;
+// Utility to look up base and satfinite variants of a conversion intrinsic
+class CvtIntrinsics<string suffix> {
+ Intrinsic base = !cast<Intrinsic>("int_nvvm_" # suffix);
+ Intrinsic sf = !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite");
+}
+
foreach rnd = ["rn", "rz"] in {
foreach relu = ["", "_relu"] in {
defvar BaseMode = !cast<PatLeaf>("Cvt" # !toupper(rnd # relu));
@@ -2290,39 +2296,30 @@ foreach rnd = ["rn", "rz"] in {
defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
defvar PZOPreds =
!if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>);
+ defvar SFPreds = !if(pzo, [hasConvertWithPZOSupport]<Predicate>,
+ [PTX81, SM80]<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);
+ defvar FF2BF16X2 = CvtIntrinsics<"ff2bf16x2_" # rnd # relu>;
+ defvar FF2F16X2 = CvtIntrinsics<"ff2f16x2_" # rnd # relu>;
+ defvar F2BF16 = CvtIntrinsics<"f2bf16_" # rnd # relu>;
+ defvar F2F16 = CvtIntrinsics<"f2f16_" # rnd # relu>;
- def : Pat<(FF2BF16X2 f32:$a, f32:$b, PZO),
+ def : Pat<(FF2BF16X2.base f32:$a, f32:$b, PZO),
(CVT_bf16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
- def : Pat<(FF2F16X2 f32:$a, f32:$b, PZO),
+ def : Pat<(FF2F16X2.base f32:$a, f32:$b, PZO),
(CVT_f16x2_f32 $a, $b, Mode)>, Requires<PZOPreds>;
- def : Pat<(F2BF16 f32:$a, PZO), (CVT_bf16_f32 $a, Mode)>,
+ def : Pat<(F2BF16.base f32:$a, PZO), (CVT_bf16_f32 $a, Mode)>,
Requires<PZOPreds>;
- def : Pat<(F2F16 f32:$a, PZO), (CVT_f16_f32 $a, Mode)>,
+ def : Pat<(F2F16.base 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 = !if(pzo, [hasConvertWithPZOSupport]<Predicate>,
- [PTX81, SM80]<Predicate>);
-
- def : Pat<(FF2BF16X2_SF f32:$a, f32:$b, PZO),
+ 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),
+ 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)>,
+ 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)>,
+ def : Pat<(F2F16.sf f32:$a, PZO), (CVT_f16_f32_sf $a, Mode)>,
Requires<SFPreds>;
}
}
@@ -2339,20 +2336,16 @@ foreach relu = ["", "_relu"] in {
[hasConvertWithStochasticRounding]<Predicate>,
!if(pzo, [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");
+ defvar BF16X2 = CvtIntrinsics<"ff2bf16x2_rs" # relu>;
+ defvar F16X2 = CvtIntrinsics<"ff2f16x2_rs" # relu>;
- def : Pat<(BF16X2 f32:$a, f32:$b, i32:$c, PZO),
+ 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),
+ 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),
+ 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),
+ def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c, PZO),
(CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>;
}
}
>From e78b3b4ced712a9d89c07da65b75a7019d707b2c Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Wed, 26 Aug 2026 09:20:14 +0000
Subject: [PATCH 6/6] address comment
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 14 +++++++-------
1 file changed, 7 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 06e2e11be9da9..57524a316db4e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2281,7 +2281,7 @@ def : Pat<(int_nvvm_ui2f_rm i32:$a), (CVT_f32_u32 $a, CvtRM)>;
def : Pat<(int_nvvm_ui2f_rp i32:$a), (CVT_f32_u32 $a, CvtRP)>;
// Utility to look up base and satfinite variants of a conversion intrinsic
-class CvtIntrinsics<string suffix> {
+class ToCvtIntrinsics<string suffix> {
Intrinsic base = !cast<Intrinsic>("int_nvvm_" # suffix);
Intrinsic sf = !cast<Intrinsic>("int_nvvm_" # suffix # "_satfinite");
}
@@ -2299,10 +2299,10 @@ foreach rnd = ["rn", "rz"] in {
defvar SFPreds = !if(pzo, [hasConvertWithPZOSupport]<Predicate>,
[PTX81, SM80]<Predicate>);
- defvar FF2BF16X2 = CvtIntrinsics<"ff2bf16x2_" # rnd # relu>;
- defvar FF2F16X2 = CvtIntrinsics<"ff2f16x2_" # rnd # relu>;
- defvar F2BF16 = CvtIntrinsics<"f2bf16_" # rnd # relu>;
- defvar F2F16 = CvtIntrinsics<"f2f16_" # rnd # relu>;
+ 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>;
@@ -2336,8 +2336,8 @@ foreach relu = ["", "_relu"] in {
[hasConvertWithStochasticRounding]<Predicate>,
!if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>));
- defvar BF16X2 = CvtIntrinsics<"ff2bf16x2_rs" # relu>;
- defvar F16X2 = CvtIntrinsics<"ff2f16x2_rs" # relu>;
+ 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>;
More information about the Mlir-commits
mailing list