Author: Alex MacLean Date: 2025-04-17T07:37:24-07:00 New Revision: 7daa5010ab6a2ae77ac88ffd84e2cc37a2e11faa
URL: https://github.com/llvm/llvm-project/commit/7daa5010ab6a2ae77ac88ffd84e2cc37a2e11faa DIFF: https://github.com/llvm/llvm-project/commit/7daa5010ab6a2ae77ac88ffd84e2cc37a2e11faa.diff LOG: [NVPTX] Cleanup and document nvvm.fabs intrinsics, adding f16 support (#135644) This change unifies the NVVM intrinsics for floating point absolute value into two new overloaded intrinsics "llvm.nvvm.fabs.*" and "llvm.nvvm.fabs.ftz.*". Documentation has been added specifying the semantics of these intrinsics to clarify how they differ from "llvm.fabs.*". In addition, support for these new intrinsics is extended to cover the f16 variants. Added: llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll Modified: clang/include/clang/Basic/BuiltinsNVPTX.td clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp clang/test/CodeGen/builtins-nvptx-native-half-type.c clang/test/CodeGen/builtins-nvptx.c llvm/docs/NVPTXUsage.rst llvm/include/llvm/IR/IntrinsicsNVVM.td llvm/lib/IR/AutoUpgrade.cpp llvm/lib/Target/NVPTX/NVPTXInstrInfo.td llvm/lib/Target/NVPTX/NVPTXIntrinsics.td llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-autoupgrade.ll llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll Removed: ################################################################################ diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td index bdbdfa2cea6c6..f797e29fe66a3 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.td +++ b/clang/include/clang/Basic/BuiltinsNVPTX.td @@ -321,6 +321,11 @@ def __nvvm_fabs_ftz_f : NVPTXBuiltin<"float(float)">; def __nvvm_fabs_f : NVPTXBuiltin<"float(float)">; def __nvvm_fabs_d : NVPTXBuiltin<"double(double)">; +def __nvvm_fabs_f16 : NVPTXBuiltinSMAndPTX<"__fp16(__fp16)", SM_53, PTX65>; +def __nvvm_fabs_f16x2 : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(_Vector<2, __fp16>)", SM_53, PTX65>; +def __nvvm_fabs_ftz_f16 : NVPTXBuiltinSMAndPTX<"__fp16(__fp16)", SM_53, PTX65>; +def __nvvm_fabs_ftz_f16x2 : NVPTXBuiltinSMAndPTX<"_Vector<2, __fp16>(_Vector<2, __fp16>)", SM_53, PTX65>; + // Round def __nvvm_round_ftz_f : NVPTXBuiltin<"float(float)">; diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp index 0f7ab9fd3b099..002af4f931c09 100644 --- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp @@ -1034,6 +1034,21 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, case NVPTX::BI__nvvm_fmin_xorsign_abs_f16x2: return MakeHalfType(Intrinsic::nvvm_fmin_xorsign_abs_f16x2, BuiltinID, E, *this); + case NVPTX::BI__nvvm_fabs_f: + case NVPTX::BI__nvvm_abs_bf16: + case NVPTX::BI__nvvm_abs_bf16x2: + case NVPTX::BI__nvvm_fabs_f16: + case NVPTX::BI__nvvm_fabs_f16x2: + return Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, + EmitScalarExpr(E->getArg(0))); + case NVPTX::BI__nvvm_fabs_ftz_f: + case NVPTX::BI__nvvm_fabs_ftz_f16: + case NVPTX::BI__nvvm_fabs_ftz_f16x2: + return Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs_ftz, + EmitScalarExpr(E->getArg(0))); + case NVPTX::BI__nvvm_fabs_d: + return Builder.CreateUnaryIntrinsic(Intrinsic::fabs, + EmitScalarExpr(E->getArg(0))); case NVPTX::BI__nvvm_ldg_h: case NVPTX::BI__nvvm_ldg_h2: return MakeHalfType(Intrinsic::not_intrinsic, BuiltinID, E, *this); diff --git a/clang/test/CodeGen/builtins-nvptx-native-half-type.c b/clang/test/CodeGen/builtins-nvptx-native-half-type.c index 511497702ff7f..01a004efd71e4 100644 --- a/clang/test/CodeGen/builtins-nvptx-native-half-type.c +++ b/clang/test/CodeGen/builtins-nvptx-native-half-type.c @@ -26,14 +26,14 @@ // RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 %s // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \ -// RUN: sm_53 -target-feature +ptx42 -fcuda-is-device -fnative-half-type \ +// RUN: sm_53 -target-feature +ptx65 -fcuda-is-device -fnative-half-type \ // RUN: -emit-llvm -o - -x cuda %s \ -// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s +// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX65_SM53 %s // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \ -// RUN: -target-cpu sm_53 -target-feature +ptx42 -fcuda-is-device \ +// RUN: -target-cpu sm_53 -target-feature +ptx65 -fcuda-is-device \ // RUN: -fnative-half-type -emit-llvm -o - -x cuda %s \ -// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s +// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX65_SM53 %s #define __device__ __attribute__((device)) @@ -108,25 +108,25 @@ __device__ void nvvm_fma_f16_f16x2_sm80() { // CHECK-LABEL: nvvm_fma_f16_f16x2_sm53 __device__ void nvvm_fma_f16_f16x2_sm53() { #if __CUDA_ARCH__ >= 530 - // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.f16 + // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.f16 __nvvm_fma_rn_f16(0.1f16, 0.1f16, 0.1f16); - // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.f16 + // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.ftz.f16 __nvvm_fma_rn_ftz_f16(0.1f16, 0.1f16, 0.1f16); - // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.sat.f16 + // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.sat.f16 __nvvm_fma_rn_sat_f16(0.1f16, 0.1f16, 0.1f16); - // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.sat.f16 + // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.ftz.sat.f16 __nvvm_fma_rn_ftz_sat_f16(0.1f16, 0.1f16, 0.1f16); - // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.f16x2 + // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.f16x2 __nvvm_fma_rn_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16}, {0.1f16, 0.7f16}); - // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2 + // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2 __nvvm_fma_rn_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16}, {0.1f16, 0.7f16}); - // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2 + // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2 __nvvm_fma_rn_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16}, {0.1f16, 0.7f16}); - // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2 + // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2 __nvvm_fma_rn_ftz_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16}, {0.1f16, 0.7f16}); #endif @@ -173,6 +173,23 @@ __device__ void nvvm_min_max_sm86() { // CHECK: ret void } +// CHECK-LABEL: nvvm_fabs_f16 +__device__ void nvvm_fabs_f16() { +#if __CUDA_ARCH__ >= 530 + // CHECK: call half @llvm.nvvm.fabs.f16 + __nvvm_fabs_f16(0.1f16); + // CHECK: call half @llvm.nvvm.fabs.ftz.f16 + __nvvm_fabs_ftz_f16(0.1f16); + // CHECK: call <2 x half> @llvm.nvvm.fabs.v2f16 + __nvvm_fabs_f16x2({0.1f16, 0.7f16}); + // CHECK: call <2 x half> @llvm.nvvm.fabs.ftz.v2f16 + __nvvm_fabs_ftz_f16x2({0.1f16, 0.7f16}); +#endif + // CHECK: ret void +} + + + typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2))); // CHECK-LABEL: nvvm_ldg_native_half_types diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 7404ce01c535c..639c18190f436 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -245,6 +245,14 @@ __device__ void nvvm_math(float f1, float f2, double d1, double d2) { // CHECK: call double @llvm.nvvm.rcp.rn.d double td4 = __nvvm_rcp_rn_d(d2); +// CHECK: call float @llvm.nvvm.fabs.f32 + float t6 = __nvvm_fabs_f(f1); +// CHECK: call float @llvm.nvvm.fabs.ftz.f32 + float t7 = __nvvm_fabs_ftz_f(f2); + +// CHECK: call double @llvm.fabs.f64 + double td5 = __nvvm_fabs_d(d1); + // CHECK: call void @llvm.nvvm.membar.cta() __nvvm_membar_cta(); // CHECK: call void @llvm.nvvm.membar.gl() @@ -1181,9 +1189,9 @@ __device__ void nvvm_cvt_sm100a_sm101a_sm120a() { __device__ void nvvm_abs_neg_bf16_bf16x2_sm80() { #if __CUDA_ARCH__ >= 800 - // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.abs.bf16(bfloat 0xR3DCD) + // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.fabs.bf16(bfloat 0xR3DCD) __nvvm_abs_bf16(BF16); - // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.abs.bf16x2(<2 x bfloat> splat (bfloat 0xR3DCD)) + // CHECK_PTX70_SM80: call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> splat (bfloat 0xR3DCD)) __nvvm_abs_bf16x2(BF16X2); // CHECK_PTX70_SM80: call bfloat @llvm.nvvm.neg.bf16(bfloat 0xR3DCD) diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst index 621879fc5648b..ef6e888286def 100644 --- a/llvm/docs/NVPTXUsage.rst +++ b/llvm/docs/NVPTXUsage.rst @@ -309,6 +309,59 @@ space casted to this space), 1 is returned, otherwise 0 is returned. Arithmetic Intrinsics --------------------- +'``llvm.nvvm.fabs.*``' Intrinsic +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare float @llvm.nvvm.fabs.f32(float %a) + declare double @llvm.nvvm.fabs.f64(double %a) + declare half @llvm.nvvm.fabs.f16(half %a) + declare <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half> %a) + declare bfloat @llvm.nvvm.fabs.bf16(bfloat %a) + declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> %a) + +Overview: +""""""""" + +The '``llvm.nvvm.fabs.*``' intrinsics return the absolute value of the operand. + +Semantics: +"""""""""" + +Unlike, '``llvm.fabs.*``', these intrinsics do not perfectly preserve NaN +values. Instead, a NaN input yeilds an unspecified NaN output. + + +'``llvm.nvvm.fabs.ftz.*``' Intrinsic +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare float @llvm.nvvm.fabs.ftz.f32(float %a) + declare half @llvm.nvvm.fabs.ftz.f16(half %a) + declare <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half> %a) + +Overview: +""""""""" + +The '``llvm.nvvm.fabs.ftz.*``' intrinsics return the absolute value of the +operand, flushing subnormals to sign preserving zero. + +Semantics: +"""""""""" + +Before the absolute value is taken, the input is flushed to sign preserving +zero if it is a subnormal. In addtion, unlike '``llvm.fabs.*``', a NaN input +yields an unspecified NaN output. + + '``llvm.nvvm.idp2a.[us].[us]``' Intrinsics ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 7af051d22f9b7..93895048faaf3 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -1039,18 +1039,18 @@ let TargetPrefix = "nvvm" in { // Abs // - def int_nvvm_fabs_ftz_f : ClangBuiltin<"__nvvm_fabs_ftz_f">, - DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>; - def int_nvvm_fabs_f : ClangBuiltin<"__nvvm_fabs_f">, - DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]>; - def int_nvvm_fabs_d : ClangBuiltin<"__nvvm_fabs_d">, - DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem, IntrSpeculatable]>; + def int_nvvm_fabs_ftz : + DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], + [IntrNoMem, IntrSpeculatable]>; + def int_nvvm_fabs : + DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>], + [IntrNoMem, IntrSpeculatable]>; // // Abs, Neg bf16, bf16x2 // - foreach unary = ["abs", "neg"] in { + foreach unary = ["neg"] in { def int_nvvm_ # unary # _bf16 : ClangBuiltin<!strconcat("__nvvm_", unary, "_bf16")>, DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty], [IntrNoMem]>; diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp index 699191ceb8d4d..7a91da5f9f7eb 100644 --- a/llvm/lib/IR/AutoUpgrade.cpp +++ b/llvm/lib/IR/AutoUpgrade.cpp @@ -939,12 +939,6 @@ static bool upgradeArmOrAarch64IntrinsicFunction(bool IsArm, Function *F, } static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name) { - if (Name.consume_front("abs.")) - return StringSwitch<Intrinsic::ID>(Name) - .Case("bf16", Intrinsic::nvvm_abs_bf16) - .Case("bf16x2", Intrinsic::nvvm_abs_bf16x2) - .Default(Intrinsic::not_intrinsic); - if (Name.consume_front("fma.rn.")) return StringSwitch<Intrinsic::ID>(Name) .Case("bf16", Intrinsic::nvvm_fma_rn_bf16) @@ -1291,7 +1285,8 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn, bool Expand = false; if (Name.consume_front("abs.")) // nvvm.abs.{i,ii} - Expand = Name == "i" || Name == "ll"; + Expand = + Name == "i" || Name == "ll" || Name == "bf16" || Name == "bf16x2"; else if (Name == "clz.ll" || Name == "popc.ll" || Name == "h2f" || Name == "swap.lo.hi.b64") Expand = true; @@ -2316,6 +2311,13 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI, Value *Cmp = Builder.CreateICmpSGE( Arg, llvm::Constant::getNullValue(Arg->getType()), "abs.cond"); Rep = Builder.CreateSelect(Cmp, Arg, Neg, "abs"); + } else if (Name == "abs.bf16" || Name == "abs.bf16x2") { + Type *Ty = (Name == "abs.bf16") + ? Builder.getBFloatTy() + : FixedVectorType::get(Builder.getBFloatTy(), 2); + Value *Arg = Builder.CreateBitCast(CI->getArgOperand(0), Ty); + Value *Abs = Builder.CreateUnaryIntrinsic(Intrinsic::nvvm_fabs, Arg); + Rep = Builder.CreateBitCast(Abs, CI->getType()); } else if (Name.starts_with("atomic.load.add.f32.p") || Name.starts_with("atomic.load.add.f64.p")) { Value *Ptr = CI->getArgOperand(0); diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 46549d850cf0a..ee6380a8a89c4 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -226,14 +226,17 @@ class RegTyInfo<ValueType ty, NVPTXRegClass rc, Operand imm, SDNode imm_node, int Size = ty.Size; } -def I16RT : RegTyInfo<i16, Int16Regs, i16imm, imm>; -def I32RT : RegTyInfo<i32, Int32Regs, i32imm, imm>; -def I64RT : RegTyInfo<i64, Int64Regs, i64imm, imm>; - -def F32RT : RegTyInfo<f32, Float32Regs, f32imm, fpimm>; -def F64RT : RegTyInfo<f64, Float64Regs, f64imm, fpimm>; -def F16RT : RegTyInfo<f16, Int16Regs, f16imm, fpimm, supports_imm = 0>; -def BF16RT : RegTyInfo<bf16, Int16Regs, bf16imm, fpimm, supports_imm = 0>; +def I16RT : RegTyInfo<i16, Int16Regs, i16imm, imm>; +def I32RT : RegTyInfo<i32, Int32Regs, i32imm, imm>; +def I64RT : RegTyInfo<i64, Int64Regs, i64imm, imm>; + +def F32RT : RegTyInfo<f32, Float32Regs, f32imm, fpimm>; +def F64RT : RegTyInfo<f64, Float64Regs, f64imm, fpimm>; +def F16RT : RegTyInfo<f16, Int16Regs, f16imm, fpimm, supports_imm = 0>; +def BF16RT : RegTyInfo<bf16, Int16Regs, bf16imm, fpimm, supports_imm = 0>; + +def F16X2RT : RegTyInfo<v2f16, Int32Regs, ?, ?, supports_imm = 0>; +def BF16X2RT : RegTyInfo<v2bf16, Int32Regs, ?, ?, supports_imm = 0>; // Template for instructions which take three int64, int32, or int16 args. diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index b5d3c9a05ec34..9634ad8eb8da2 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -983,12 +983,13 @@ def : Pat<(int_nvvm_fmin_d // We need a full string for OpcStr here because we need to deal with case like // INT_PTX_RECIP. -class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass, - NVPTXRegClass src_regclass, Intrinsic IntOP, list<Predicate> Preds = []> - : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0), - OpcStr, - [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>, - Requires<Preds>; +class F_MATH_1<string OpcStr, RegTyInfo dst, RegTyInfo src, Intrinsic IntOP, + list<Predicate> Preds = []> + : NVPTXInst<(outs dst.RC:$dst), + (ins src.RC:$src0), + OpcStr, + [(set dst.Ty:$dst, (IntOP src.Ty:$src0))]>, + Requires<Preds>; // We need a full string for OpcStr here because we need to deal with the case // like INT_PTX_NATIVE_POWR_F. @@ -1307,13 +1308,20 @@ def : Pat<(int_nvvm_ceil_d f64:$a), // Abs // -def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_fabs_ftz_f>; -def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_fabs_f>; +multiclass F_ABS<string suffix, RegTyInfo RT, bit support_ftz, list<Predicate> preds = []> { + def "" : F_MATH_1<"abs." # suffix # " \t$dst, $src0;", RT, RT, int_nvvm_fabs, preds>; + if support_ftz then + def _FTZ : F_MATH_1<"abs.ftz." # suffix # " \t$dst, $src0;", RT, RT, int_nvvm_fabs_ftz, preds>; +} + +defm ABS_F16 : F_ABS<"f16", F16RT, support_ftz = true, preds = [hasPTX<65>, hasSM<53>]>; +defm ABS_F16X2 : F_ABS<"f16x2", F16X2RT, support_ftz = true, preds = [hasPTX<65>, hasSM<53>]>; + +defm ABS_BF16 : F_ABS<"bf16", BF16RT, support_ftz = false, preds = [hasPTX<70>, hasSM<80>]>; +defm ABS_BF16X2 : F_ABS<"bf16x2", BF16X2RT, support_ftz = false, preds = [hasPTX<70>, hasSM<80>]>; -def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_fabs_d>; +defm ABS_F32 : F_ABS<"f32", F32RT, support_ftz = true>; +defm ABS_F64 : F_ABS<"f64", F64RT, support_ftz = false>; // // copysign @@ -1332,17 +1340,13 @@ def COPYSIGN_D : [(set f64:$dst, (fcopysign_nvptx f64:$src1, f64:$src0))]>; // -// Abs, Neg bf16, bf16x2 +// Neg bf16, bf16x2 // -def INT_NVVM_ABS_BF16 : F_MATH_1<"abs.bf16 \t$dst, $src0;", Int16Regs, - Int16Regs, int_nvvm_abs_bf16, [hasPTX<70>, hasSM<80>]>; -def INT_NVVM_ABS_BF16X2 : F_MATH_1<"abs.bf16x2 \t$dst, $src0;", Int32Regs, - Int32Regs, int_nvvm_abs_bf16x2, [hasPTX<70>, hasSM<80>]>; -def INT_NVVM_NEG_BF16 : F_MATH_1<"neg.bf16 \t$dst, $src0;", Int16Regs, - Int16Regs, int_nvvm_neg_bf16, [hasPTX<70>, hasSM<80>]>; -def INT_NVVM_NEG_BF16X2 : F_MATH_1<"neg.bf16x2 \t$dst, $src0;", Int32Regs, - Int32Regs, int_nvvm_neg_bf16x2, [hasPTX<70>, hasSM<80>]>; +def INT_NVVM_NEG_BF16 : F_MATH_1<"neg.bf16 \t$dst, $src0;", BF16RT, + BF16RT, int_nvvm_neg_bf16, [hasPTX<70>, hasSM<80>]>; +def INT_NVVM_NEG_BF16X2 : F_MATH_1<"neg.bf16x2 \t$dst, $src0;", BF16X2RT, + BF16X2RT, int_nvvm_neg_bf16x2, [hasPTX<70>, hasSM<80>]>; // // Round @@ -1382,16 +1386,16 @@ def : Pat<(int_nvvm_saturate_d f64:$a), // def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>; + F32RT, F32RT, int_nvvm_ex2_approx_ftz_f>; def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>; + F32RT, F32RT, int_nvvm_ex2_approx_f>; def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;", - Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>; + F64RT, F64RT, int_nvvm_ex2_approx_d>; def INT_NVVM_EX2_APPROX_F16 : F_MATH_1<"ex2.approx.f16 \t$dst, $src0;", - Int16Regs, Int16Regs, int_nvvm_ex2_approx_f16, [hasPTX<70>, hasSM<75>]>; + F16RT, F16RT, int_nvvm_ex2_approx_f16, [hasPTX<70>, hasSM<75>]>; def INT_NVVM_EX2_APPROX_F16X2 : F_MATH_1<"ex2.approx.f16x2 \t$dst, $src0;", - Int32Regs, Int32Regs, int_nvvm_ex2_approx_f16x2, [hasPTX<70>, hasSM<75>]>; + F16X2RT, F16X2RT, int_nvvm_ex2_approx_f16x2, [hasPTX<70>, hasSM<75>]>; def : Pat<(fexp2 f32:$a), (INT_NVVM_EX2_APPROX_FTZ_F $a)>, Requires<[doF32FTZ]>; @@ -1403,11 +1407,11 @@ def : Pat<(fexp2 v2f16:$a), (INT_NVVM_EX2_APPROX_F16X2 $a)>, Requires<[useFP16Math]>; def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>; + F32RT, F32RT, int_nvvm_lg2_approx_ftz_f>; def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>; + F32RT, F32RT, int_nvvm_lg2_approx_f>; def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;", - Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>; + F64RT, F64RT, int_nvvm_lg2_approx_d>; def : Pat<(flog2 f32:$a), (INT_NVVM_LG2_APPROX_FTZ_F $a)>, Requires<[doF32FTZ]>; @@ -1419,14 +1423,14 @@ def : Pat<(flog2 f32:$a), (INT_NVVM_LG2_APPROX_F $a)>, // def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>; + F32RT, F32RT, int_nvvm_sin_approx_ftz_f>; def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sin_approx_f>; + F32RT, F32RT, int_nvvm_sin_approx_f>; def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>; + F32RT, F32RT, int_nvvm_cos_approx_ftz_f>; def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_cos_approx_f>; + F32RT, F32RT, int_nvvm_cos_approx_f>; // // Fma @@ -1511,69 +1515,69 @@ defm INT_NVVM_FMA : FMA_INST; // def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>; + F32RT, F32RT, int_nvvm_rcp_rn_ftz_f>; def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>; + F32RT, F32RT, int_nvvm_rcp_rn_f>; def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>; + F32RT, F32RT, int_nvvm_rcp_rz_ftz_f>; def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>; + F32RT, F32RT, int_nvvm_rcp_rz_f>; def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>; + F32RT, F32RT, int_nvvm_rcp_rm_ftz_f>; def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>; + F32RT, F32RT, int_nvvm_rcp_rm_f>; def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>; + F32RT, F32RT, int_nvvm_rcp_rp_ftz_f>; def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>; + F32RT, F32RT, int_nvvm_rcp_rp_f>; -def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_rcp_rn_d>; -def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_rcp_rz_d>; -def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_rcp_rm_d>; -def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_rcp_rp_d>; +def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_rcp_rn_d>; +def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_rcp_rz_d>; +def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_rcp_rm_d>; +def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_rcp_rp_d>; def INT_NVVM_RCP_APPROX_FTZ_F : F_MATH_1<"rcp.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rcp_approx_ftz_f>; + F32RT, F32RT, int_nvvm_rcp_approx_ftz_f>; def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;", - Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>; + F64RT, F64RT, int_nvvm_rcp_approx_ftz_d>; // // Sqrt // def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>; -def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_sqrt_rn_f>; + F32RT, F32RT, int_nvvm_sqrt_rn_ftz_f>; +def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", F32RT, + F32RT, int_nvvm_sqrt_rn_f>; def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>; -def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_sqrt_rz_f>; + F32RT, F32RT, int_nvvm_sqrt_rz_ftz_f>; +def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", F32RT, + F32RT, int_nvvm_sqrt_rz_f>; def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>; -def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_sqrt_rm_f>; + F32RT, F32RT, int_nvvm_sqrt_rm_ftz_f>; +def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", F32RT, + F32RT, int_nvvm_sqrt_rm_f>; def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>; -def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs, - Float32Regs, int_nvvm_sqrt_rp_f>; + F32RT, F32RT, int_nvvm_sqrt_rp_ftz_f>; +def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", F32RT, + F32RT, int_nvvm_sqrt_rp_f>; def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>; + F32RT, F32RT, int_nvvm_sqrt_approx_ftz_f>; def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>; + F32RT, F32RT, int_nvvm_sqrt_approx_f>; -def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_sqrt_rn_d>; -def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_sqrt_rz_d>; -def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_sqrt_rm_d>; -def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs, - Float64Regs, int_nvvm_sqrt_rp_d>; +def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_sqrt_rn_d>; +def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_sqrt_rz_d>; +def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_sqrt_rm_d>; +def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", F64RT, + F64RT, int_nvvm_sqrt_rp_d>; // nvvm_sqrt intrinsic def : Pat<(int_nvvm_sqrt_f f32:$a), @@ -1590,16 +1594,16 @@ def : Pat<(int_nvvm_sqrt_f f32:$a), // def INT_NVVM_RSQRT_APPROX_FTZ_F - : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs, + : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", F32RT, F32RT, int_nvvm_rsqrt_approx_ftz_f>; def INT_NVVM_RSQRT_APPROX_FTZ_D - : F_MATH_1<"rsqrt.approx.ftz.f64 \t$dst, $src0;", Float64Regs, Float64Regs, + : F_MATH_1<"rsqrt.approx.ftz.f64 \t$dst, $src0;", F64RT, F64RT, int_nvvm_rsqrt_approx_ftz_d>; def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;", - Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>; + F32RT, F32RT, int_nvvm_rsqrt_approx_f>; def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;", - Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>; + F64RT, F64RT, int_nvvm_rsqrt_approx_d>; // 1.0f / sqrt_approx -> rsqrt_approx def: Pat<(fdiv FloatConst1, (int_nvvm_sqrt_approx_f f32:$a)), @@ -1815,13 +1819,13 @@ def INT_NVVM_D2I_LO : F_MATH_1< ".reg .b32 %temp; \n\t", "mov.b64 \t{$dst, %temp}, $src0;\n\t", "}}"), - Int32Regs, Float64Regs, int_nvvm_d2i_lo>; + I32RT, F64RT, int_nvvm_d2i_lo>; def INT_NVVM_D2I_HI : F_MATH_1< !strconcat("{{\n\t", ".reg .b32 %temp; \n\t", "mov.b64 \t{%temp, $dst}, $src0;\n\t", "}}"), - Int32Regs, Float64Regs, int_nvvm_d2i_hi>; + I32RT, F64RT, int_nvvm_d2i_hi>; def : Pat<(int_nvvm_f2ll_rn_ftz f32:$a), (CVT_s64_f32 $a, CvtRNI_FTZ)>; diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp index 81ad01bea8867..72245fe83491d 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp @@ -185,8 +185,6 @@ static Instruction *convertNvvmIntrinsicToLlvm(InstCombiner &IC, return {Intrinsic::ceil, FTZ_MustBeOff}; case Intrinsic::nvvm_ceil_ftz_f: return {Intrinsic::ceil, FTZ_MustBeOn}; - case Intrinsic::nvvm_fabs_d: - return {Intrinsic::fabs, FTZ_Any}; case Intrinsic::nvvm_floor_d: return {Intrinsic::floor, FTZ_Any}; case Intrinsic::nvvm_floor_f: diff --git a/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll b/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll new file mode 100644 index 0000000000000..dd9ef220a9b47 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll @@ -0,0 +1,142 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc < %s -mcpu=sm_80 -mattr=+ptx70 | FileCheck --check-prefixes=CHECK %s +; RUN: %if ptxas-11.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %} + +target triple = "nvptx-nvidia-cuda" + +declare float @llvm.nvvm.fabs.f32(float) +declare float @llvm.nvvm.fabs.ftz.f32(float) +declare double @llvm.nvvm.fabs.f64(double) +declare half @llvm.nvvm.fabs.f16(half) +declare half @llvm.nvvm.fabs.ftz.f16(half) +declare <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half>) +declare <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half>) +declare bfloat @llvm.nvvm.fabs.bf16(bfloat) +declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat>) + + +define float @fabs_float(float %a) { +; CHECK-LABEL: fabs_float( +; CHECK: { +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [fabs_float_param_0]; +; CHECK-NEXT: abs.f32 %f2, %f1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %ret = call float @llvm.nvvm.fabs.f32(float %a) + ret float %ret +} + +define float @fabs_float_ftz(float %a) { +; CHECK-LABEL: fabs_float_ftz( +; CHECK: { +; CHECK-NEXT: .reg .f32 %f<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f32 %f1, [fabs_float_ftz_param_0]; +; CHECK-NEXT: abs.ftz.f32 %f2, %f1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f2; +; CHECK-NEXT: ret; + %ret = call float @llvm.nvvm.fabs.ftz.f32(float %a) + ret float %ret +} + +define double @fabs_double(double %a) { +; CHECK-LABEL: fabs_double( +; CHECK: { +; CHECK-NEXT: .reg .f64 %fd<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.f64 %fd1, [fabs_double_param_0]; +; CHECK-NEXT: abs.f64 %fd2, %fd1; +; CHECK-NEXT: st.param.f64 [func_retval0], %fd2; +; CHECK-NEXT: ret; + %ret = call double @llvm.nvvm.fabs.f64(double %a) + ret double %ret +} + +define half @fabs_half(half %a) { +; CHECK-LABEL: fabs_half( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_param_0]; +; CHECK-NEXT: abs.f16 %rs2, %rs1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs2; +; CHECK-NEXT: ret; + %ret = call half @llvm.nvvm.fabs.f16(half %a) + ret half %ret +} + +define half @fabs_half_ftz(half %a) { +; CHECK-LABEL: fabs_half_ftz( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_ftz_param_0]; +; CHECK-NEXT: abs.ftz.f16 %rs2, %rs1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs2; +; CHECK-NEXT: ret; + %ret = call half @llvm.nvvm.fabs.ftz.f16(half %a) + ret half %ret +} + +define <2 x half> @fabs_v2half(<2 x half> %a) { +; CHECK-LABEL: fabs_v2half( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_param_0]; +; CHECK-NEXT: abs.f16x2 %r2, %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %ret = call <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half> %a) + ret <2 x half> %ret +} + +define <2 x half> @fabs_v2half_ftz(<2 x half> %a) { +; CHECK-LABEL: fabs_v2half_ftz( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_ftz_param_0]; +; CHECK-NEXT: abs.ftz.f16x2 %r2, %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %ret = call <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half> %a) + ret <2 x half> %ret +} + +define bfloat @fabs_bf16(bfloat %a) { +; CHECK-LABEL: fabs_bf16( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b16 %rs1, [fabs_bf16_param_0]; +; CHECK-NEXT: abs.bf16 %rs2, %rs1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs2; +; CHECK-NEXT: ret; + %ret = call bfloat @llvm.nvvm.fabs.bf16(bfloat %a) + ret bfloat %ret +} + +define <2 x bfloat> @fabs_v2bf16(<2 x bfloat> %a) { +; CHECK-LABEL: fabs_v2bf16( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2bf16_param_0]; +; CHECK-NEXT: abs.bf16x2 %r2, %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %ret = call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> %a) + ret <2 x bfloat> %ret +} diff --git a/llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-autoupgrade.ll b/llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-autoupgrade.ll index 7faac51ff27ca..c04fd07ec5da1 100644 --- a/llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-autoupgrade.ll +++ b/llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70-autoupgrade.ll @@ -1,6 +1,9 @@ ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck %s ; RUN: %if ptxas-11.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %} +declare bfloat @llvm.nvvm.abs.bf16(bfloat) +declare <2 x bfloat> @llvm.nvvm.abs.bf16x2(<2 x bfloat>) + ; CHECK-LABEL: abs_bf16 define bfloat @abs_bf16(bfloat %0) { ; CHECK-NOT: call diff --git a/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll b/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll index a1517409e1ee1..1819f4ed181c0 100644 --- a/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll +++ b/llvm/test/Transforms/InstCombine/NVPTX/nvvm-intrins.ll @@ -43,7 +43,7 @@ define float @ceil_float_ftz(float %a) #0 { ; CHECK-LABEL: @fabs_double define double @fabs_double(double %a) #0 { -; CHECK: call double @llvm.fabs.f64 +; CHECK: call double @llvm.nvvm.fabs.f64 %ret = call double @llvm.nvvm.fabs.d(double %a) ret double %ret } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
