[llvm] 5587627 - [NVPTX] Add support for f16 fabs (#116107)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 18 13:47:11 PST 2024
Author: Alex MacLean
Date: 2024-11-18T13:47:08-08:00
New Revision: 55876278d362020503db5f0e66313829c40ff640
URL: https://github.com/llvm/llvm-project/commit/55876278d362020503db5f0e66313829c40ff640
DIFF: https://github.com/llvm/llvm-project/commit/55876278d362020503db5f0e66313829c40ff640.diff
LOG: [NVPTX] Add support for f16 fabs (#116107)
Add support for f16 and f16x2 support for abs. See PTX ISA 9.7.4.6. Half
Precision Floating Point Instructions: abs
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-abs
Added:
llvm/test/CodeGen/NVPTX/f16-abs.ll
Modified:
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 4ad0200ca5cf83..e93430a27dc32e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -862,16 +862,19 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setOperationAction(Op, MVT::bf16, Promote);
AddPromotedToType(Op, MVT::bf16, MVT::f32);
}
- for (const auto &Op : {ISD::FABS}) {
- setOperationAction(Op, MVT::f16, Promote);
- setOperationAction(Op, MVT::f32, Legal);
- setOperationAction(Op, MVT::f64, Legal);
- setOperationAction(Op, MVT::v2f16, Expand);
- setBF16OperationAction(Op, MVT::v2bf16, Legal, Expand);
- setBF16OperationAction(Op, MVT::bf16, Legal, Promote);
- if (getOperationAction(Op, MVT::bf16) == Promote)
- AddPromotedToType(Op, MVT::bf16, MVT::f32);
+
+ setOperationAction(ISD::FABS, {MVT::f32, MVT::f64}, Legal);
+ if (STI.getPTXVersion() >= 65) {
+ setFP16OperationAction(ISD::FABS, MVT::f16, Legal, Promote);
+ setFP16OperationAction(ISD::FABS, MVT::v2f16, Legal, Expand);
+ } else {
+ setOperationAction(ISD::FABS, MVT::f16, Promote);
+ setOperationAction(ISD::FABS, MVT::v2f16, Expand);
}
+ setBF16OperationAction(ISD::FABS, MVT::v2bf16, Legal, Expand);
+ setBF16OperationAction(ISD::FABS, MVT::bf16, Legal, Promote);
+ if (getOperationAction(ISD::FABS, MVT::bf16) == Promote)
+ AddPromotedToType(ISD::FABS, MVT::bf16, MVT::f32);
for (const auto &Op : {ISD::FMINNUM, ISD::FMAXNUM}) {
setOperationAction(Op, MVT::f32, Legal);
diff --git a/llvm/test/CodeGen/NVPTX/f16-abs.ll b/llvm/test/CodeGen/NVPTX/f16-abs.ll
new file mode 100644
index 00000000000000..d12653e813bd10
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/f16-abs.ll
@@ -0,0 +1,98 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+
+; ## FP16 abs is not supported by PTX version (PTX < 65).
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx60 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | FileCheck -check-prefix CHECK-NOF16 %s
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx60 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_53 \
+; RUN: %}
+
+; ## FP16 support explicitly disabled (--nvptx-no-f16-math).
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx65 --nvptx-no-f16-math \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | FileCheck -check-prefix CHECK-NOF16 %s
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx65 --nvptx-no-f16-math \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_53 \
+; RUN: %}
+
+; ## FP16 is not supported by hardware (SM < 53).
+; RUN: llc < %s -mcpu=sm_52 -mattr=+ptx65 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | FileCheck -check-prefix CHECK-NOF16 %s
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_52 -mattr=+ptx65 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_52 \
+; RUN: %}
+
+; ## Full FP16 abs support.
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx65 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | FileCheck -check-prefix CHECK-F16-ABS %s
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_53 -mattr=+ptx65 \
+; RUN: -O0 -disable-post-ra -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_53 \
+; RUN: %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare half @llvm.fabs.f16(half %a)
+declare <2 x half> @llvm.fabs.v2f16(<2 x half> %a)
+
+define half @test_fabs(half %a) {
+; CHECK-NOF16-LABEL: test_fabs(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [test_fabs_param_0];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs1;
+; CHECK-NOF16-NEXT: abs.f32 %f2, %f1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs2, %f2;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs2;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-ABS-LABEL: test_fabs(
+; CHECK-F16-ABS: {
+; CHECK-F16-ABS-NEXT: .reg .b16 %rs<3>;
+; CHECK-F16-ABS-EMPTY:
+; CHECK-F16-ABS-NEXT: // %bb.0:
+; CHECK-F16-ABS-NEXT: ld.param.b16 %rs1, [test_fabs_param_0];
+; CHECK-F16-ABS-NEXT: abs.f16 %rs2, %rs1;
+; CHECK-F16-ABS-NEXT: st.param.b16 [func_retval0], %rs2;
+; CHECK-F16-ABS-NEXT: ret;
+ %r = call half @llvm.fabs.f16(half %a)
+ ret half %r
+}
+
+define <2 x half> @test_fabs_2(<2 x half> %a) #0 {
+; CHECK-F16-LABEL: test_fabs_2(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .b32 %r<5>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fabs_2_param_0];
+; CHECK-F16-NEXT: and.b32 %r3, %r1, 2147450879;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-F16-ABS-LABEL: test_fabs_2(
+; CHECK-F16-ABS: {
+; CHECK-F16-ABS-NEXT: .reg .b32 %r<3>;
+; CHECK-F16-ABS-EMPTY:
+; CHECK-F16-ABS-NEXT: // %bb.0:
+; CHECK-F16-ABS-NEXT: ld.param.b32 %r1, [test_fabs_2_param_0];
+; CHECK-F16-ABS-NEXT: abs.f16x2 %r2, %r1;
+; CHECK-F16-ABS-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-F16-ABS-NEXT: ret;
+ %r = call <2 x half> @llvm.fabs.v2f16(<2 x half> %a)
+ ret <2 x half> %r
+}
+
More information about the llvm-commits
mailing list