shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
llvm-project/llvm/test/CodeGen/NVPTX/fexp2.ll
Alex MacLean f03782dd67
[NVPTX] Fixup v2i8 parameter and return lowering (#145585) 
This change fixes v2i8 lowering for parameters and returned values. As
part of this work, I move the lowering for return values to use generic
ISD::STORE nodes as these are more flexible and have existing
legalization handling.

Note that calling a function with v2i8 arguments or returns is still not
working but this is left for a subsequent change as this MR is already
fairly large.

Partially addresses #128853
2025-06-27 09:26:10 -07:00

393 lines
14 KiB
LLVM
Raw Permalink Blame History

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mcpu=sm_50 -mattr=+ptx32 | FileCheck --check-prefixes=CHECK %s
; RUN: llc < %s -mcpu=sm_75 -mattr=+ptx70 | FileCheck --check-prefixes=CHECK-FP16 %s
; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK-BF16 %s
; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_50 -mattr=+ptx32 | %ptxas-verify -arch=sm_50 %}
; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_75 -mattr=+ptx70 | %ptxas-verify -arch=sm_75 %}
; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}
target triple = "nvptx64-nvidia-cuda"
; --- f32 ---
; CHECK-LABEL: exp2_test
define float @exp2_test(float %in) {
; CHECK-LABEL: exp2_test(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [exp2_test_param_0];
; CHECK-NEXT: ex2.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_test(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b32 %r<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_test_param_0];
; CHECK-FP16-NEXT: ex2.approx.f32 %r2, %r1;
; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_test(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b32 %r<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_test_param_0];
; CHECK-BF16-NEXT: ex2.approx.f32 %r2, %r1;
; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call float @llvm.exp2.f32(float %in)
ret float %exp2
}
; CHECK-LABEL: exp2_ftz_test
define float @exp2_ftz_test(float %in) #0 {
; CHECK-LABEL: exp2_ftz_test(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];
; CHECK-NEXT: ex2.approx.ftz.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_ftz_test(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b32 %r<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];
; CHECK-FP16-NEXT: ex2.approx.ftz.f32 %r2, %r1;
; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_ftz_test(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b32 %r<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];
; CHECK-BF16-NEXT: ex2.approx.ftz.f32 %r2, %r1;
; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call float @llvm.exp2.f32(float %in)
ret float %exp2
}
; CHECK-LABEL: exp2_test_v
define <2 x float> @exp2_test_v(<2 x float> %in) {
; CHECK-LABEL: exp2_test_v(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];
; CHECK-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-NEXT: ex2.approx.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_test_v(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b32 %r<5>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];
; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-FP16-NEXT: ex2.approx.f32 %r4, %r1;
; CHECK-FP16-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_test_v(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b32 %r<5>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];
; CHECK-BF16-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-BF16-NEXT: ex2.approx.f32 %r4, %r1;
; CHECK-BF16-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call <2 x float> @llvm.exp2.v2f32(<2 x float> %in)
ret <2 x float> %exp2
}
; --- f16 ---
; CHECK-LABEL: exp2_f16_test
define half @exp2_f16_test(half %in) {
; CHECK-LABEL: exp2_f16_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];
; CHECK-NEXT: cvt.f32.f16 %r1, %rs1;
; CHECK-NEXT: ex2.approx.f32 %r2, %r1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_f16_test(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b16 %rs<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];
; CHECK-FP16-NEXT: ex2.approx.f16 %rs2, %rs1;
; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_f16_test(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b16 %rs<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];
; CHECK-BF16-NEXT: ex2.approx.f16 %rs2, %rs1;
; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call half @llvm.exp2.f16(half %in)
ret half %exp2
}
; COM: we should never have .ftz for f16
; CHECK-LABEL: exp2_f16_ftz_test
define half @exp2_f16_ftz_test(half %in) #0 {
; CHECK-LABEL: exp2_f16_ftz_test(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];
; CHECK-NEXT: cvt.ftz.f32.f16 %r1, %rs1;
; CHECK-NEXT: ex2.approx.ftz.f32 %r2, %r1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_f16_ftz_test(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b16 %rs<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];
; CHECK-FP16-NEXT: ex2.approx.f16 %rs2, %rs1;
; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_f16_ftz_test(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b16 %rs<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];
; CHECK-BF16-NEXT: ex2.approx.f16 %rs2, %rs1;
; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call half @llvm.exp2.f16(half %in)
ret half %exp2
}
; CHECK-LABEL: exp2_f16_test_v
define <2 x half> @exp2_f16_test_v(<2 x half> %in) {
; CHECK-LABEL: exp2_f16_test_v(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_f16_test_v_param_0];
; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;
; CHECK-NEXT: ex2.approx.f32 %r2, %r1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r2;
; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;
; CHECK-NEXT: ex2.approx.f32 %r4, %r3;
; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r4;
; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};
; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_f16_test_v(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .b32 %r<3>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_f16_test_v_param_0];
; CHECK-FP16-NEXT: ex2.approx.f16x2 %r2, %r1;
; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_f16_test_v(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b32 %r<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_f16_test_v_param_0];
; CHECK-BF16-NEXT: ex2.approx.f16x2 %r2, %r1;
; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call <2 x half> @llvm.exp2.v2f16(<2 x half> %in)
ret <2 x half> %exp2
}
; --- bf16 ---
; COM: we should always have .ftz for bf16
; CHECK-LABEL: exp2_bf16_test
define bfloat @exp2_bf16_test(bfloat %in) {
; CHECK-LABEL: exp2_bf16_test(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.b16 %r1, [exp2_bf16_test_param_0];
; CHECK-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;
; CHECK-NEXT: add.s32 %r5, %r4, %r3;
; CHECK-NEXT: add.s32 %r6, %r5, 32767;
; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;
; CHECK-NEXT: or.b32 %r7, %r3, 4194304;
; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;
; CHECK-NEXT: shr.u32 %r9, %r8, 16;
; CHECK-NEXT: st.param.b16 [func_retval0], %r9;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_bf16_test(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .pred %p<2>;
; CHECK-FP16-NEXT: .reg .b32 %r<10>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.b16 %r1, [exp2_bf16_test_param_0];
; CHECK-FP16-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-FP16-NEXT: bfe.u32 %r4, %r3, 16, 1;
; CHECK-FP16-NEXT: add.s32 %r5, %r4, %r3;
; CHECK-FP16-NEXT: add.s32 %r6, %r5, 32767;
; CHECK-FP16-NEXT: setp.nan.f32 %p1, %r3, %r3;
; CHECK-FP16-NEXT: or.b32 %r7, %r3, 4194304;
; CHECK-FP16-NEXT: selp.b32 %r8, %r7, %r6, %p1;
; CHECK-FP16-NEXT: shr.u32 %r9, %r8, 16;
; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %r9;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_bf16_test(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b16 %rs<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_bf16_test_param_0];
; CHECK-BF16-NEXT: ex2.approx.ftz.bf16 %rs2, %rs1;
; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call bfloat @llvm.exp2.bf16(bfloat %in)
ret bfloat %exp2
}
; CHECK-LABEL: exp2_bf16_test_v
define <2 x bfloat> @exp2_bf16_test_v(<2 x bfloat> %in) {
; CHECK-LABEL: exp2_bf16_test_v(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<18>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_bf16_test_v_param_0];
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
; CHECK-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;
; CHECK-NEXT: add.s32 %r5, %r4, %r3;
; CHECK-NEXT: add.s32 %r6, %r5, 32767;
; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;
; CHECK-NEXT: or.b32 %r7, %r3, 4194304;
; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;
; CHECK-NEXT: cvt.u32.u16 %r9, %rs1;
; CHECK-NEXT: shl.b32 %r10, %r9, 16;
; CHECK-NEXT: ex2.approx.f32 %r11, %r10;
; CHECK-NEXT: bfe.u32 %r12, %r11, 16, 1;
; CHECK-NEXT: add.s32 %r13, %r12, %r11;
; CHECK-NEXT: add.s32 %r14, %r13, 32767;
; CHECK-NEXT: setp.nan.f32 %p2, %r11, %r11;
; CHECK-NEXT: or.b32 %r15, %r11, 4194304;
; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p2;
; CHECK-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;
; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
; CHECK-NEXT: ret;
;
; CHECK-FP16-LABEL: exp2_bf16_test_v(
; CHECK-FP16: {
; CHECK-FP16-NEXT: .reg .pred %p<3>;
; CHECK-FP16-NEXT: .reg .b16 %rs<3>;
; CHECK-FP16-NEXT: .reg .b32 %r<18>;
; CHECK-FP16-EMPTY:
; CHECK-FP16-NEXT: // %bb.0: // %entry
; CHECK-FP16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_bf16_test_v_param_0];
; CHECK-FP16-NEXT: cvt.u32.u16 %r1, %rs2;
; CHECK-FP16-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;
; CHECK-FP16-NEXT: bfe.u32 %r4, %r3, 16, 1;
; CHECK-FP16-NEXT: add.s32 %r5, %r4, %r3;
; CHECK-FP16-NEXT: add.s32 %r6, %r5, 32767;
; CHECK-FP16-NEXT: setp.nan.f32 %p1, %r3, %r3;
; CHECK-FP16-NEXT: or.b32 %r7, %r3, 4194304;
; CHECK-FP16-NEXT: selp.b32 %r8, %r7, %r6, %p1;
; CHECK-FP16-NEXT: cvt.u32.u16 %r9, %rs1;
; CHECK-FP16-NEXT: shl.b32 %r10, %r9, 16;
; CHECK-FP16-NEXT: ex2.approx.f32 %r11, %r10;
; CHECK-FP16-NEXT: bfe.u32 %r12, %r11, 16, 1;
; CHECK-FP16-NEXT: add.s32 %r13, %r12, %r11;
; CHECK-FP16-NEXT: add.s32 %r14, %r13, 32767;
; CHECK-FP16-NEXT: setp.nan.f32 %p2, %r11, %r11;
; CHECK-FP16-NEXT: or.b32 %r15, %r11, 4194304;
; CHECK-FP16-NEXT: selp.b32 %r16, %r15, %r14, %p2;
; CHECK-FP16-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;
; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r17;
; CHECK-FP16-NEXT: ret;
;
; CHECK-BF16-LABEL: exp2_bf16_test_v(
; CHECK-BF16: {
; CHECK-BF16-NEXT: .reg .b32 %r<3>;
; CHECK-BF16-EMPTY:
; CHECK-BF16-NEXT: // %bb.0: // %entry
; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_bf16_test_v_param_0];
; CHECK-BF16-NEXT: ex2.approx.ftz.bf16x2 %r2, %r1;
; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-BF16-NEXT: ret;
entry:
%exp2 = call <2 x bfloat> @llvm.exp2.v2bf16(<2 x bfloat> %in)
ret <2 x bfloat> %exp2
}
declare float @llvm.exp2.f32(float %val)
declare <2 x float> @llvm.exp2.v2f32(<2 x float> %val)
declare half @llvm.exp2.f16(half %val)
declare <2 x half> @llvm.exp2.v2f16(<2 x half> %val)
declare bfloat @llvm.exp2.bf16(bfloat %val)
declare <2 x bfloat> @llvm.exp2.v2bf16(<2 x bfloat> %val)
attributes #0 = {"denormal-fp-math"="preserve-sign"}
Reference in New Issue View Git Blame Copy Permalink