llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-mlir Author: Srinivasa Ravi (Wolfram70) <details> <summary>Changes</summary> 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. --- Patch is 64.67 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/214667.diff 14 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsNVPTX.td (+33-1) - (modified) clang/lib/CodeGen/CGBuiltin.cpp (+18) - (modified) clang/lib/CodeGen/CGBuiltin.h (+5) - (modified) clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp (+46-3) - (modified) clang/test/CodeGen/builtins-nvptx.c (+125-48) - (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+6-6) - (modified) llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp (+5) - (modified) llvm/lib/Target/NVPTX/NVPTX.h (+2-1) - (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+13-6) - (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+74-60) - (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.h (+4) - (added) llvm/test/CodeGen/NVPTX/convert-sm107f-pzo.ll (+221) - (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+6) - (modified) mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir (+24-24) ``````````diff 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 ... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/214667 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
