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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
