[llvm] [NVPTX] support packed f32 instructions for sm_100+ (PR #126337)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 27 09:08:55 PDT 2025
================
@@ -0,0 +1,1699 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; ## Full FP32x2 support enabled by default.
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_100 \
+; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
+; RUN: | FileCheck --check-prefixes=CHECK %s
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_100 \
+; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_100 \
+; RUN: %}
+
+target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
+target triple = "nvptx64-nvidia-cuda"
+
+define <2 x float> @test_ret_const() #0 {
+; CHECK-LABEL: test_ret_const(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.b32 %r1, 0f40000000;
+; CHECK-NEXT: mov.b32 %r2, 0f3F800000;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
+; CHECK-NEXT: ret;
+ ret <2 x float> <float 1.0, float 2.0>
+}
+
+define float @test_extract_0(<2 x float> %a) #0 {
+; CHECK-LABEL: test_extract_0(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_0_param_0];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %e = extractelement <2 x float> %a, i32 0
+ ret float %e
+}
+
+define float @test_extract_1(<2 x float> %a) #0 {
+; CHECK-LABEL: test_extract_1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_1_param_0];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %e = extractelement <2 x float> %a, i32 1
+ ret float %e
+}
+
+; NOTE: disabled as -O3 miscompiles this into pointer arithmetic on
+; test_extract_i_param_0 where the symbol's address is not taken first (that
+; is, moved to a temporary)
+; define float @test_extract_i(<2 x float> %a, i64 %idx) #0 {
+; %e = extractelement <2 x float> %a, i64 %idx
+; ret float %e
+; }
+
+define <2 x float> @test_fadd(<2 x float> %a, <2 x float> %b) #0 {
+; CHECK-LABEL: test_fadd(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fadd_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_param_0];
+; CHECK-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fadd_imm_0(<2 x float> %a) #0 {
+; CHECK-LABEL: test_fadd_imm_0(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40000000;
+; CHECK-NEXT: mov.b32 %r2, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd2, {%r2, %r1};
+; CHECK-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> <float 1.0, float 2.0>, %a
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fadd_imm_1(<2 x float> %a) #0 {
+; CHECK-LABEL: test_fadd_imm_1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40000000;
+; CHECK-NEXT: mov.b32 %r2, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd2, {%r2, %r1};
+; CHECK-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> %a, <float 1.0, float 2.0>
+ ret <2 x float> %r
+}
+
+define <4 x float> @test_fadd_v4(<4 x float> %a, <4 x float> %b) #0 {
+; CHECK-LABEL: test_fadd_v4(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_param_1];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_param_0];
+; CHECK-NEXT: add.rn.f32x2 %rd5, %rd2, %rd4;
+; CHECK-NEXT: add.rn.f32x2 %rd6, %rd1, %rd3;
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};
+; CHECK-NEXT: ret;
+ %r = fadd <4 x float> %a, %b
+ ret <4 x float> %r
+}
+
+define <4 x float> @test_fadd_imm_0_v4(<4 x float> %a) #0 {
+; CHECK-LABEL: test_fadd_imm_0_v4(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40800000;
+; CHECK-NEXT: mov.b32 %r2, 0f40400000;
+; CHECK-NEXT: mov.b64 %rd3, {%r2, %r1};
+; CHECK-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;
+; CHECK-NEXT: mov.b32 %r3, 0f40000000;
+; CHECK-NEXT: mov.b32 %r4, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd5, {%r4, %r3};
+; CHECK-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
+; CHECK-NEXT: ret;
+ %r = fadd <4 x float> <float 1.0, float 2.0, float 3.0, float 4.0>, %a
+ ret <4 x float> %r
+}
+
+define <4 x float> @test_fadd_imm_1_v4(<4 x float> %a) #0 {
+; CHECK-LABEL: test_fadd_imm_1_v4(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40800000;
+; CHECK-NEXT: mov.b32 %r2, 0f40400000;
+; CHECK-NEXT: mov.b64 %rd3, {%r2, %r1};
+; CHECK-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;
+; CHECK-NEXT: mov.b32 %r3, 0f40000000;
+; CHECK-NEXT: mov.b32 %r4, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd5, {%r4, %r3};
+; CHECK-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
+; CHECK-NEXT: ret;
+ %r = fadd <4 x float> %a, <float 1.0, float 2.0, float 3.0, float 4.0>
+ ret <4 x float> %r
+}
+
+define <2 x float> @test_fsub(<2 x float> %a, <2 x float> %b) #0 {
+; CHECK-LABEL: test_fsub(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fsub_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fsub_param_0];
+; CHECK-NEXT: sub.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fsub <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fneg(<2 x float> %a) #0 {
+; CHECK-LABEL: test_fneg(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_param_0];
+; CHECK-NEXT: neg.f32 %r3, %r2;
+; CHECK-NEXT: neg.f32 %r4, %r1;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
+; CHECK-NEXT: ret;
+ %r = fneg <2 x float> %a
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fmul(<2 x float> %a, <2 x float> %b) #0 {
+; CHECK-LABEL: test_fmul(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fmul_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fmul_param_0];
+; CHECK-NEXT: mul.rn.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fmul <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fdiv(<2 x float> %a, <2 x float> %b) #0 {
+; CHECK-LABEL: test_fdiv(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<7>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_param_0];
+; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_param_1];
+; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
+; CHECK-NEXT: div.rn.f32 %r6, %r1, %r3;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
+; CHECK-NEXT: ret;
+ %r = fdiv <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_frem(<2 x float> %a, <2 x float> %b) #0 {
+; CHECK-LABEL: test_frem(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<3>;
+; CHECK-NEXT: .reg .b32 %r<15>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_param_0];
+; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_param_1];
+; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
+; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;
+; CHECK-NEXT: neg.f32 %r7, %r6;
+; CHECK-NEXT: fma.rn.f32 %r8, %r7, %r4, %r2;
+; CHECK-NEXT: testp.infinite.f32 %p1, %r4;
+; CHECK-NEXT: selp.f32 %r9, %r2, %r8, %p1;
+; CHECK-NEXT: div.rn.f32 %r10, %r1, %r3;
+; CHECK-NEXT: cvt.rzi.f32.f32 %r11, %r10;
+; CHECK-NEXT: neg.f32 %r12, %r11;
+; CHECK-NEXT: fma.rn.f32 %r13, %r12, %r3, %r1;
+; CHECK-NEXT: testp.infinite.f32 %p2, %r3;
+; CHECK-NEXT: selp.f32 %r14, %r1, %r13, %p2;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};
+; CHECK-NEXT: ret;
+ %r = frem <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fadd_ftz(<2 x float> %a, <2 x float> %b) #2 {
+; CHECK-LABEL: test_fadd_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fadd_ftz_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_ftz_param_0];
+; CHECK-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> %a, %b
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fadd_imm_0_ftz(<2 x float> %a) #2 {
+; CHECK-LABEL: test_fadd_imm_0_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_ftz_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40000000;
+; CHECK-NEXT: mov.b32 %r2, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd2, {%r2, %r1};
+; CHECK-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> <float 1.0, float 2.0>, %a
+ ret <2 x float> %r
+}
+
+define <2 x float> @test_fadd_imm_1_ftz(<2 x float> %a) #2 {
+; CHECK-LABEL: test_fadd_imm_1_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_ftz_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40000000;
+; CHECK-NEXT: mov.b32 %r2, 0f3F800000;
+; CHECK-NEXT: mov.b64 %rd2, {%r2, %r1};
+; CHECK-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %r = fadd <2 x float> %a, <float 1.0, float 2.0>
+ ret <2 x float> %r
+}
+
+define <4 x float> @test_fadd_v4_ftz(<4 x float> %a, <4 x float> %b) #2 {
+; CHECK-LABEL: test_fadd_v4_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_ftz_param_1];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_ftz_param_0];
+; CHECK-NEXT: add.rn.ftz.f32x2 %rd5, %rd2, %rd4;
+; CHECK-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd3;
+; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};
+; CHECK-NEXT: ret;
+ %r = fadd <4 x float> %a, %b
+ ret <4 x float> %r
+}
+
+define <4 x float> @test_fadd_imm_0_v4_ftz(<4 x float> %a) #2 {
+; CHECK-LABEL: test_fadd_imm_0_v4_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_ftz_param_0];
+; CHECK-NEXT: mov.b32 %r1, 0f40800000;
+; CHECK-NEXT: mov.b32 %r2, 0f40400000;
+; CHECK-NEXT: mov.b64 %rd3, {%r2, %r1};
----------------
AlexMaclean wrote:
This seems pretty redundant. I would hope the SDAG would simplify this to a single mov of an immediate.
https://github.com/llvm/llvm-project/pull/126337
More information about the llvm-commits
mailing list