[llvm] 82113a4 - [LLVM][NVPTX] Remove nonexistent ftz ops (#106100)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 28 09:10:22 PDT 2024
Author: Billy Zhu
Date: 2024-08-28T09:10:17-07:00
New Revision: 82113a432c5bffe026682ea117a3e2cd67a2fed0
URL: https://github.com/llvm/llvm-project/commit/82113a432c5bffe026682ea117a3e2cd67a2fed0
DIFF: https://github.com/llvm/llvm-project/commit/82113a432c5bffe026682ea117a3e2cd67a2fed0.diff
LOG: [LLVM][NVPTX] Remove nonexistent ftz ops (#106100)
According to the PTX
[spec](https://docs.nvidia.com/cuda/parallel-thread-execution/#half-precision-floating-point-instructions-max),
max & min instructions do not support the `ftz` modifier for `bf16` &
`bf16x2` types. This PR removes them from instr info, and the non-ftz
legal versions will be emitted instead.
Added:
Modified:
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/test/CodeGen/NVPTX/bf16-instructions.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index a6dfb704e38d2e..b7e210805db904 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -334,25 +334,12 @@ multiclass FMINIMUMMAXIMUM<string OpcStr, bit NaN, SDNode OpNode> {
!strconcat(OpcStr, ".f16x2 \t$dst, $a, $b;"),
[(set Int32Regs:$dst, (OpNode (v2f16 Int32Regs:$a), (v2f16 Int32Regs:$b)))]>,
Requires<[useFP16Math, hasSM<80>, hasPTX<70>]>;
- def bf16rr_ftz :
- NVPTXInst<(outs Int16Regs:$dst),
- (ins Int16Regs:$a, Int16Regs:$b),
- !strconcat(OpcStr, ".ftz.bf16 \t$dst, $a, $b;"),
- [(set Int16Regs:$dst, (OpNode (bf16 Int16Regs:$a), (bf16 Int16Regs:$b)))]>,
- Requires<[hasBF16Math, doF32FTZ, hasSM<80>, hasPTX<70>]>;
def bf16rr :
NVPTXInst<(outs Int16Regs:$dst),
(ins Int16Regs:$a, Int16Regs:$b),
!strconcat(OpcStr, ".bf16 \t$dst, $a, $b;"),
[(set Int16Regs:$dst, (OpNode (bf16 Int16Regs:$a), (bf16 Int16Regs:$b)))]>,
Requires<[hasBF16Math, hasSM<80>, hasPTX<70>]>;
-
- def bf16x2rr_ftz :
- NVPTXInst<(outs Int32Regs:$dst),
- (ins Int32Regs:$a, Int32Regs:$b),
- !strconcat(OpcStr, ".ftz.bf16x2 \t$dst, $a, $b;"),
- [(set Int32Regs:$dst, (OpNode (v2bf16 Int32Regs:$a), (v2bf16 Int32Regs:$b)))]>,
- Requires<[hasBF16Math, hasSM<80>, hasPTX<70>, doF32FTZ]>;
def bf16x2rr :
NVPTXInst<(outs Int32Regs:$dst),
(ins Int32Regs:$a, Int32Regs:$b),
diff --git a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
index f029e5f322968b..95bca39c73ad73 100644
--- a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll
@@ -1,8 +1,10 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK,SM70 %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK,SM80 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx71 -denormal-fp-math-f32=preserve-sign | FileCheck --check-prefixes=CHECK,SM80-FTZ %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK,SM90 %s
; RUN: %if ptxas-11.8 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | %ptxas-verify -arch=sm_80 %}
+; RUN: %if ptxas-11.8 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx71 -denormal-fp-math-f32=preserve-sign | %ptxas-verify -arch=sm_80 %}
; RUN: %if ptxas-11.8 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}
target triple = "nvptx64-nvidia-cuda"
@@ -53,6 +55,21 @@ define bfloat @test_fadd(bfloat %0, bfloat %1) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fadd(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];
+; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_fadd_param_1];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs1;
+; SM80-FTZ-NEXT: add.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %f3;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fadd(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<4>;
@@ -109,6 +126,21 @@ define bfloat @test_fsub(bfloat %0, bfloat %1) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fsub(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];
+; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_fsub_param_1];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs1;
+; SM80-FTZ-NEXT: sub.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %f3;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fsub(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<4>;
@@ -193,6 +225,29 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_faddx2(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_faddx2_param_1];
+; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
+; SM80-FTZ-NEXT: add.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
+; SM80-FTZ-NEXT: add.rn.ftz.f32 %f6, %f5, %f4;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
+; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_faddx2(
; SM90: {
; SM90-NEXT: .reg .b32 %r<4>;
@@ -277,6 +332,29 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fsubx2(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_fsubx2_param_1];
+; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
+; SM80-FTZ-NEXT: sub.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
+; SM80-FTZ-NEXT: sub.rn.ftz.f32 %f6, %f5, %f4;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
+; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fsubx2(
; SM90: {
; SM90-NEXT: .reg .b32 %r<4>;
@@ -361,6 +439,29 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fmulx2(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_fmulx2_param_1];
+; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
+; SM80-FTZ-NEXT: mul.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
+; SM80-FTZ-NEXT: mul.rn.ftz.f32 %f6, %f5, %f4;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
+; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fmulx2(
; SM90: {
; SM90-NEXT: .reg .b32 %r<4>;
@@ -445,6 +546,29 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fdiv(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-NEXT: .reg .f32 %f<7>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_fdiv_param_1];
+; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
+; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
+; SM80-FTZ-NEXT: div.rn.ftz.f32 %f3, %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
+; SM80-FTZ-NEXT: div.rn.ftz.f32 %f6, %f5, %f4;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
+; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fdiv(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<7>;
@@ -521,6 +645,17 @@ define float @test_fpext_float(bfloat %a) #0 {
; SM80-NEXT: st.param.f32 [func_retval0+0], %f1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fpext_float(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs1;
+; SM80-FTZ-NEXT: st.param.f32 [func_retval0+0], %f1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fpext_float(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
@@ -567,6 +702,17 @@ define bfloat @test_fptrunc_float(float %a) #0 {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fptrunc_float(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.f32 %f1, [test_fptrunc_float_param_0];
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fptrunc_float(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
@@ -618,6 +764,19 @@ define bfloat @test_fadd_imm_1(bfloat %a) #0 {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs2;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fadd_imm_1(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .f32 %f<3>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs1;
+; SM80-FTZ-NEXT: add.rn.ftz.f32 %f2, %f1, 0f3F800000;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %f2;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs2;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fadd_imm_1(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<4>;
@@ -722,6 +881,32 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {
; SM80-NEXT: st.param.v4.f32 [func_retval0+16], {%f4, %f3, %f2, %f1};
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_extload_bf16x8(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<9>;
+; SM80-FTZ-NEXT: .reg .b32 %r<5>;
+; SM80-FTZ-NEXT: .reg .f32 %f<9>;
+; SM80-FTZ-NEXT: .reg .b64 %rd<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u64 %rd1, [test_extload_bf16x8_param_0];
+; SM80-FTZ-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r1;
+; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r2;
+; SM80-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r3;
+; SM80-FTZ-NEXT: mov.b32 {%rs7, %rs8}, %r4;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs8;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs7;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f3, %rs6;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs5;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs4;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f6, %rs3;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f7, %rs2;
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f8, %rs1;
+; SM80-FTZ-NEXT: st.param.v4.f32 [func_retval0+0], {%f8, %f7, %f6, %f5};
+; SM80-FTZ-NEXT: st.param.v4.f32 [func_retval0+16], {%f4, %f3, %f2, %f1};
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_extload_bf16x8(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<9>;
@@ -782,6 +967,20 @@ define i16 @test_fptosi_i16(bfloat %a) {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fptosi_i16(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .b32 %r<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rzi.ftz.s16.f32 %rs2, %f1;
+; SM80-FTZ-NEXT: cvt.u32.u16 %r1, %rs2;
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fptosi_i16(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -827,6 +1026,20 @@ define i16 @test_fptoui_i16(bfloat %a) {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_fptoui_i16(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .b32 %r<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rzi.ftz.u16.f32 %rs2, %f1;
+; SM80-FTZ-NEXT: cvt.u32.u16 %r1, %rs2;
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_fptoui_i16(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -876,6 +1089,18 @@ define bfloat @test_sitofp_i16(i16 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs2;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_sitofp_i16(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u16 %rs1, [test_sitofp_i16_param_0];
+; SM80-FTZ-NEXT: cvt.rn.f32.s16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs2;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_sitofp_i16(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -923,6 +1148,18 @@ define bfloat @test_uitofp_i8(i8 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs2;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_uitofp_i8(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u8 %rs1, [test_uitofp_i8_param_0];
+; SM80-FTZ-NEXT: cvt.rn.f32.u16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs2;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_uitofp_i8(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -978,6 +1215,23 @@ define bfloat @test_uitofp_i1(i1 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_uitofp_i1(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .pred %p<2>;
+; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
+; SM80-FTZ-NEXT: .reg .b32 %r<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u8 %rs1, [test_uitofp_i1_param_0];
+; SM80-FTZ-NEXT: and.b16 %rs2, %rs1, 1;
+; SM80-FTZ-NEXT: setp.eq.b16 %p1, %rs2, 1;
+; SM80-FTZ-NEXT: selp.u32 %r1, 1, 0, %p1;
+; SM80-FTZ-NEXT: cvt.rn.f32.u32 %f1, %r1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_uitofp_i1(
; SM90: {
; SM90-NEXT: .reg .pred %p<2>;
@@ -1030,6 +1284,18 @@ define bfloat @test_uitofp_i16(i16 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs2;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_uitofp_i16(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u16 %rs1, [test_uitofp_i16_param_0];
+; SM80-FTZ-NEXT: cvt.rn.f32.u16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs2;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_uitofp_i16(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -1078,6 +1344,19 @@ define bfloat @test_uitofp_i32(i32 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_uitofp_i32(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
+; SM80-FTZ-NEXT: .reg .b32 %r<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u32 %r1, [test_uitofp_i32_param_0];
+; SM80-FTZ-NEXT: cvt.rn.f32.u32 %f1, %r1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_uitofp_i32(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
@@ -1128,6 +1407,19 @@ define bfloat @test_uitofp_i64(i64 %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs1;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_uitofp_i64(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<2>;
+; SM80-FTZ-NEXT: .reg .f32 %f<2>;
+; SM80-FTZ-NEXT: .reg .b64 %rd<2>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.u64 %rd1, [test_uitofp_i64_param_0];
+; SM80-FTZ-NEXT: cvt.rn.f32.u64 %f1, %rd1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %f1;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs1;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_uitofp_i64(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<2>;
@@ -1179,6 +1471,19 @@ define bfloat @test_roundeven(bfloat %a) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs2;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_roundeven(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<3>;
+; SM80-FTZ-NEXT: .reg .f32 %f<3>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];
+; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs1;
+; SM80-FTZ-NEXT: cvt.rni.ftz.f32.f32 %f2, %f1;
+; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %f2;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs2;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_roundeven(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<3>;
@@ -1236,6 +1541,17 @@ define bfloat @test_maximum(bfloat %a, bfloat %b) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_maximum(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];
+; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];
+; SM80-FTZ-NEXT: max.NaN.bf16 %rs3, %rs1, %rs2;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_maximum(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<4>;
@@ -1288,6 +1604,17 @@ define bfloat @test_maxnum(bfloat %a, bfloat %b) {
; SM80-NEXT: st.param.b16 [func_retval0+0], %rs3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_maxnum(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b16 %rs<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_maxnum_param_0];
+; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_maxnum_param_1];
+; SM80-FTZ-NEXT: max.bf16 %rs3, %rs1, %rs2;
+; SM80-FTZ-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_maxnum(
; SM90: {
; SM90-NEXT: .reg .b16 %rs<4>;
@@ -1368,6 +1695,17 @@ define <2 x bfloat> @test_maximum_v2(<2 x bfloat> %a, <2 x bfloat> %b) {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_maximum_v2(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_1];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_maximum_v2_param_0];
+; SM80-FTZ-NEXT: max.NaN.bf16x2 %r3, %r2, %r1;
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_maximum_v2(
; SM90: {
; SM90-NEXT: .reg .b32 %r<4>;
@@ -1440,6 +1778,17 @@ define <2 x bfloat> @test_maxnum_v2(<2 x bfloat> %a, <2 x bfloat> %b) {
; SM80-NEXT: st.param.b32 [func_retval0+0], %r3;
; SM80-NEXT: ret;
;
+; SM80-FTZ-LABEL: test_maxnum_v2(
+; SM80-FTZ: {
+; SM80-FTZ-NEXT: .reg .b32 %r<4>;
+; SM80-FTZ-EMPTY:
+; SM80-FTZ-NEXT: // %bb.0:
+; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_1];
+; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_maxnum_v2_param_0];
+; SM80-FTZ-NEXT: max.bf16x2 %r3, %r2, %r1;
+; SM80-FTZ-NEXT: st.param.b32 [func_retval0+0], %r3;
+; SM80-FTZ-NEXT: ret;
+;
; SM90-LABEL: test_maxnum_v2(
; SM90: {
; SM90-NEXT: .reg .b32 %r<4>;
More information about the llvm-commits
mailing list