393 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_50 -mattr=+ptx32 | FileCheck --check-prefixes=CHECK %s3; RUN: llc < %s -mcpu=sm_75 -mattr=+ptx70 | FileCheck --check-prefixes=CHECK-FP16 %s4; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK-BF16 %s5; RUN: %if ptxas-sm_50 && ptxas-isa-3.2 %{ llc < %s -mcpu=sm_50 -mattr=+ptx32 | %ptxas-verify -arch=sm_50 %}6; RUN: %if ptxas-sm_75 && ptxas-isa-7.0 %{ llc < %s -mcpu=sm_75 -mattr=+ptx70 | %ptxas-verify -arch=sm_75 %}7; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}8target triple = "nvptx64-nvidia-cuda"9 10; --- f32 ---11 12; CHECK-LABEL: exp2_test13define float @exp2_test(float %in) {14; CHECK-LABEL: exp2_test(15; CHECK: {16; CHECK-NEXT: .reg .b32 %r<3>;17; CHECK-EMPTY:18; CHECK-NEXT: // %bb.0: // %entry19; CHECK-NEXT: ld.param.b32 %r1, [exp2_test_param_0];20; CHECK-NEXT: ex2.approx.f32 %r2, %r1;21; CHECK-NEXT: st.param.b32 [func_retval0], %r2;22; CHECK-NEXT: ret;23;24; CHECK-FP16-LABEL: exp2_test(25; CHECK-FP16: {26; CHECK-FP16-NEXT: .reg .b32 %r<3>;27; CHECK-FP16-EMPTY:28; CHECK-FP16-NEXT: // %bb.0: // %entry29; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_test_param_0];30; CHECK-FP16-NEXT: ex2.approx.f32 %r2, %r1;31; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;32; CHECK-FP16-NEXT: ret;33;34; CHECK-BF16-LABEL: exp2_test(35; CHECK-BF16: {36; CHECK-BF16-NEXT: .reg .b32 %r<3>;37; CHECK-BF16-EMPTY:38; CHECK-BF16-NEXT: // %bb.0: // %entry39; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_test_param_0];40; CHECK-BF16-NEXT: ex2.approx.f32 %r2, %r1;41; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;42; CHECK-BF16-NEXT: ret;43entry:44 %exp2 = call float @llvm.exp2.f32(float %in)45 ret float %exp246}47 48; CHECK-LABEL: exp2_ftz_test49define float @exp2_ftz_test(float %in) #0 {50; CHECK-LABEL: exp2_ftz_test(51; CHECK: {52; CHECK-NEXT: .reg .b32 %r<3>;53; CHECK-EMPTY:54; CHECK-NEXT: // %bb.0: // %entry55; CHECK-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];56; CHECK-NEXT: ex2.approx.ftz.f32 %r2, %r1;57; CHECK-NEXT: st.param.b32 [func_retval0], %r2;58; CHECK-NEXT: ret;59;60; CHECK-FP16-LABEL: exp2_ftz_test(61; CHECK-FP16: {62; CHECK-FP16-NEXT: .reg .b32 %r<3>;63; CHECK-FP16-EMPTY:64; CHECK-FP16-NEXT: // %bb.0: // %entry65; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];66; CHECK-FP16-NEXT: ex2.approx.ftz.f32 %r2, %r1;67; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;68; CHECK-FP16-NEXT: ret;69;70; CHECK-BF16-LABEL: exp2_ftz_test(71; CHECK-BF16: {72; CHECK-BF16-NEXT: .reg .b32 %r<3>;73; CHECK-BF16-EMPTY:74; CHECK-BF16-NEXT: // %bb.0: // %entry75; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_ftz_test_param_0];76; CHECK-BF16-NEXT: ex2.approx.ftz.f32 %r2, %r1;77; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;78; CHECK-BF16-NEXT: ret;79entry:80 %exp2 = call float @llvm.exp2.f32(float %in)81 ret float %exp282}83 84; CHECK-LABEL: exp2_test_v85define <2 x float> @exp2_test_v(<2 x float> %in) {86; CHECK-LABEL: exp2_test_v(87; CHECK: {88; CHECK-NEXT: .reg .b32 %r<5>;89; CHECK-EMPTY:90; CHECK-NEXT: // %bb.0: // %entry91; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];92; CHECK-NEXT: ex2.approx.f32 %r3, %r2;93; CHECK-NEXT: ex2.approx.f32 %r4, %r1;94; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};95; CHECK-NEXT: ret;96;97; CHECK-FP16-LABEL: exp2_test_v(98; CHECK-FP16: {99; CHECK-FP16-NEXT: .reg .b32 %r<5>;100; CHECK-FP16-EMPTY:101; CHECK-FP16-NEXT: // %bb.0: // %entry102; CHECK-FP16-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];103; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;104; CHECK-FP16-NEXT: ex2.approx.f32 %r4, %r1;105; CHECK-FP16-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};106; CHECK-FP16-NEXT: ret;107;108; CHECK-BF16-LABEL: exp2_test_v(109; CHECK-BF16: {110; CHECK-BF16-NEXT: .reg .b32 %r<5>;111; CHECK-BF16-EMPTY:112; CHECK-BF16-NEXT: // %bb.0: // %entry113; CHECK-BF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [exp2_test_v_param_0];114; CHECK-BF16-NEXT: ex2.approx.f32 %r3, %r2;115; CHECK-BF16-NEXT: ex2.approx.f32 %r4, %r1;116; CHECK-BF16-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};117; CHECK-BF16-NEXT: ret;118entry:119 %exp2 = call <2 x float> @llvm.exp2.v2f32(<2 x float> %in)120 ret <2 x float> %exp2121}122 123; --- f16 ---124 125; CHECK-LABEL: exp2_f16_test126define half @exp2_f16_test(half %in) {127; CHECK-LABEL: exp2_f16_test(128; CHECK: {129; CHECK-NEXT: .reg .b16 %rs<3>;130; CHECK-NEXT: .reg .b32 %r<3>;131; CHECK-EMPTY:132; CHECK-NEXT: // %bb.0: // %entry133; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];134; CHECK-NEXT: cvt.f32.f16 %r1, %rs1;135; CHECK-NEXT: ex2.approx.f32 %r2, %r1;136; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;137; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;138; CHECK-NEXT: ret;139;140; CHECK-FP16-LABEL: exp2_f16_test(141; CHECK-FP16: {142; CHECK-FP16-NEXT: .reg .b16 %rs<3>;143; CHECK-FP16-EMPTY:144; CHECK-FP16-NEXT: // %bb.0: // %entry145; CHECK-FP16-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];146; CHECK-FP16-NEXT: ex2.approx.f16 %rs2, %rs1;147; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;148; CHECK-FP16-NEXT: ret;149;150; CHECK-BF16-LABEL: exp2_f16_test(151; CHECK-BF16: {152; CHECK-BF16-NEXT: .reg .b16 %rs<3>;153; CHECK-BF16-EMPTY:154; CHECK-BF16-NEXT: // %bb.0: // %entry155; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_f16_test_param_0];156; CHECK-BF16-NEXT: ex2.approx.f16 %rs2, %rs1;157; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;158; CHECK-BF16-NEXT: ret;159entry:160 %exp2 = call half @llvm.exp2.f16(half %in)161 ret half %exp2162}163 164; COM: we should never have .ftz for f16165; CHECK-LABEL: exp2_f16_ftz_test166define half @exp2_f16_ftz_test(half %in) #0 {167; CHECK-LABEL: exp2_f16_ftz_test(168; CHECK: {169; CHECK-NEXT: .reg .b16 %rs<3>;170; CHECK-NEXT: .reg .b32 %r<3>;171; CHECK-EMPTY:172; CHECK-NEXT: // %bb.0: // %entry173; CHECK-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];174; CHECK-NEXT: cvt.ftz.f32.f16 %r1, %rs1;175; CHECK-NEXT: ex2.approx.ftz.f32 %r2, %r1;176; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;177; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;178; CHECK-NEXT: ret;179;180; CHECK-FP16-LABEL: exp2_f16_ftz_test(181; CHECK-FP16: {182; CHECK-FP16-NEXT: .reg .b16 %rs<3>;183; CHECK-FP16-EMPTY:184; CHECK-FP16-NEXT: // %bb.0: // %entry185; CHECK-FP16-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];186; CHECK-FP16-NEXT: ex2.approx.f16 %rs2, %rs1;187; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;188; CHECK-FP16-NEXT: ret;189;190; CHECK-BF16-LABEL: exp2_f16_ftz_test(191; CHECK-BF16: {192; CHECK-BF16-NEXT: .reg .b16 %rs<3>;193; CHECK-BF16-EMPTY:194; CHECK-BF16-NEXT: // %bb.0: // %entry195; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_f16_ftz_test_param_0];196; CHECK-BF16-NEXT: ex2.approx.f16 %rs2, %rs1;197; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;198; CHECK-BF16-NEXT: ret;199entry:200 %exp2 = call half @llvm.exp2.f16(half %in)201 ret half %exp2202}203 204; CHECK-LABEL: exp2_f16_test_v205define <2 x half> @exp2_f16_test_v(<2 x half> %in) {206; CHECK-LABEL: exp2_f16_test_v(207; CHECK: {208; CHECK-NEXT: .reg .b16 %rs<5>;209; CHECK-NEXT: .reg .b32 %r<6>;210; CHECK-EMPTY:211; CHECK-NEXT: // %bb.0: // %entry212; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_f16_test_v_param_0];213; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;214; CHECK-NEXT: ex2.approx.f32 %r2, %r1;215; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r2;216; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;217; CHECK-NEXT: ex2.approx.f32 %r4, %r3;218; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r4;219; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};220; CHECK-NEXT: st.param.b32 [func_retval0], %r5;221; CHECK-NEXT: ret;222;223; CHECK-FP16-LABEL: exp2_f16_test_v(224; CHECK-FP16: {225; CHECK-FP16-NEXT: .reg .b32 %r<3>;226; CHECK-FP16-EMPTY:227; CHECK-FP16-NEXT: // %bb.0: // %entry228; CHECK-FP16-NEXT: ld.param.b32 %r1, [exp2_f16_test_v_param_0];229; CHECK-FP16-NEXT: ex2.approx.f16x2 %r2, %r1;230; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;231; CHECK-FP16-NEXT: ret;232;233; CHECK-BF16-LABEL: exp2_f16_test_v(234; CHECK-BF16: {235; CHECK-BF16-NEXT: .reg .b32 %r<3>;236; CHECK-BF16-EMPTY:237; CHECK-BF16-NEXT: // %bb.0: // %entry238; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_f16_test_v_param_0];239; CHECK-BF16-NEXT: ex2.approx.f16x2 %r2, %r1;240; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;241; CHECK-BF16-NEXT: ret;242entry:243 %exp2 = call <2 x half> @llvm.exp2.v2f16(<2 x half> %in)244 ret <2 x half> %exp2245}246 247; --- bf16 ---248 249; COM: we should always have .ftz for bf16250; CHECK-LABEL: exp2_bf16_test251define bfloat @exp2_bf16_test(bfloat %in) {252; CHECK-LABEL: exp2_bf16_test(253; CHECK: {254; CHECK-NEXT: .reg .pred %p<2>;255; CHECK-NEXT: .reg .b32 %r<10>;256; CHECK-EMPTY:257; CHECK-NEXT: // %bb.0: // %entry258; CHECK-NEXT: ld.param.b16 %r1, [exp2_bf16_test_param_0];259; CHECK-NEXT: shl.b32 %r2, %r1, 16;260; CHECK-NEXT: ex2.approx.f32 %r3, %r2;261; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;262; CHECK-NEXT: add.s32 %r5, %r4, %r3;263; CHECK-NEXT: add.s32 %r6, %r5, 32767;264; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;265; CHECK-NEXT: or.b32 %r7, %r3, 4194304;266; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;267; CHECK-NEXT: shr.u32 %r9, %r8, 16;268; CHECK-NEXT: st.param.b16 [func_retval0], %r9;269; CHECK-NEXT: ret;270;271; CHECK-FP16-LABEL: exp2_bf16_test(272; CHECK-FP16: {273; CHECK-FP16-NEXT: .reg .pred %p<2>;274; CHECK-FP16-NEXT: .reg .b32 %r<10>;275; CHECK-FP16-EMPTY:276; CHECK-FP16-NEXT: // %bb.0: // %entry277; CHECK-FP16-NEXT: ld.param.b16 %r1, [exp2_bf16_test_param_0];278; CHECK-FP16-NEXT: shl.b32 %r2, %r1, 16;279; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;280; CHECK-FP16-NEXT: bfe.u32 %r4, %r3, 16, 1;281; CHECK-FP16-NEXT: add.s32 %r5, %r4, %r3;282; CHECK-FP16-NEXT: add.s32 %r6, %r5, 32767;283; CHECK-FP16-NEXT: setp.nan.f32 %p1, %r3, %r3;284; CHECK-FP16-NEXT: or.b32 %r7, %r3, 4194304;285; CHECK-FP16-NEXT: selp.b32 %r8, %r7, %r6, %p1;286; CHECK-FP16-NEXT: shr.u32 %r9, %r8, 16;287; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %r9;288; CHECK-FP16-NEXT: ret;289;290; CHECK-BF16-LABEL: exp2_bf16_test(291; CHECK-BF16: {292; CHECK-BF16-NEXT: .reg .b16 %rs<3>;293; CHECK-BF16-EMPTY:294; CHECK-BF16-NEXT: // %bb.0: // %entry295; CHECK-BF16-NEXT: ld.param.b16 %rs1, [exp2_bf16_test_param_0];296; CHECK-BF16-NEXT: ex2.approx.ftz.bf16 %rs2, %rs1;297; CHECK-BF16-NEXT: st.param.b16 [func_retval0], %rs2;298; CHECK-BF16-NEXT: ret;299entry:300 %exp2 = call bfloat @llvm.exp2.bf16(bfloat %in)301 ret bfloat %exp2302}303 304; CHECK-LABEL: exp2_bf16_test_v305define <2 x bfloat> @exp2_bf16_test_v(<2 x bfloat> %in) {306; CHECK-LABEL: exp2_bf16_test_v(307; CHECK: {308; CHECK-NEXT: .reg .pred %p<3>;309; CHECK-NEXT: .reg .b16 %rs<3>;310; CHECK-NEXT: .reg .b32 %r<18>;311; CHECK-EMPTY:312; CHECK-NEXT: // %bb.0: // %entry313; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_bf16_test_v_param_0];314; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;315; CHECK-NEXT: shl.b32 %r2, %r1, 16;316; CHECK-NEXT: ex2.approx.f32 %r3, %r2;317; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;318; CHECK-NEXT: add.s32 %r5, %r4, %r3;319; CHECK-NEXT: add.s32 %r6, %r5, 32767;320; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;321; CHECK-NEXT: or.b32 %r7, %r3, 4194304;322; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;323; CHECK-NEXT: cvt.u32.u16 %r9, %rs1;324; CHECK-NEXT: shl.b32 %r10, %r9, 16;325; CHECK-NEXT: ex2.approx.f32 %r11, %r10;326; CHECK-NEXT: bfe.u32 %r12, %r11, 16, 1;327; CHECK-NEXT: add.s32 %r13, %r12, %r11;328; CHECK-NEXT: add.s32 %r14, %r13, 32767;329; CHECK-NEXT: setp.nan.f32 %p2, %r11, %r11;330; CHECK-NEXT: or.b32 %r15, %r11, 4194304;331; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p2;332; CHECK-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;333; CHECK-NEXT: st.param.b32 [func_retval0], %r17;334; CHECK-NEXT: ret;335;336; CHECK-FP16-LABEL: exp2_bf16_test_v(337; CHECK-FP16: {338; CHECK-FP16-NEXT: .reg .pred %p<3>;339; CHECK-FP16-NEXT: .reg .b16 %rs<3>;340; CHECK-FP16-NEXT: .reg .b32 %r<18>;341; CHECK-FP16-EMPTY:342; CHECK-FP16-NEXT: // %bb.0: // %entry343; CHECK-FP16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [exp2_bf16_test_v_param_0];344; CHECK-FP16-NEXT: cvt.u32.u16 %r1, %rs2;345; CHECK-FP16-NEXT: shl.b32 %r2, %r1, 16;346; CHECK-FP16-NEXT: ex2.approx.f32 %r3, %r2;347; CHECK-FP16-NEXT: bfe.u32 %r4, %r3, 16, 1;348; CHECK-FP16-NEXT: add.s32 %r5, %r4, %r3;349; CHECK-FP16-NEXT: add.s32 %r6, %r5, 32767;350; CHECK-FP16-NEXT: setp.nan.f32 %p1, %r3, %r3;351; CHECK-FP16-NEXT: or.b32 %r7, %r3, 4194304;352; CHECK-FP16-NEXT: selp.b32 %r8, %r7, %r6, %p1;353; CHECK-FP16-NEXT: cvt.u32.u16 %r9, %rs1;354; CHECK-FP16-NEXT: shl.b32 %r10, %r9, 16;355; CHECK-FP16-NEXT: ex2.approx.f32 %r11, %r10;356; CHECK-FP16-NEXT: bfe.u32 %r12, %r11, 16, 1;357; CHECK-FP16-NEXT: add.s32 %r13, %r12, %r11;358; CHECK-FP16-NEXT: add.s32 %r14, %r13, 32767;359; CHECK-FP16-NEXT: setp.nan.f32 %p2, %r11, %r11;360; CHECK-FP16-NEXT: or.b32 %r15, %r11, 4194304;361; CHECK-FP16-NEXT: selp.b32 %r16, %r15, %r14, %p2;362; CHECK-FP16-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;363; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r17;364; CHECK-FP16-NEXT: ret;365;366; CHECK-BF16-LABEL: exp2_bf16_test_v(367; CHECK-BF16: {368; CHECK-BF16-NEXT: .reg .b32 %r<3>;369; CHECK-BF16-EMPTY:370; CHECK-BF16-NEXT: // %bb.0: // %entry371; CHECK-BF16-NEXT: ld.param.b32 %r1, [exp2_bf16_test_v_param_0];372; CHECK-BF16-NEXT: ex2.approx.ftz.bf16x2 %r2, %r1;373; CHECK-BF16-NEXT: st.param.b32 [func_retval0], %r2;374; CHECK-BF16-NEXT: ret;375entry:376 %exp2 = call <2 x bfloat> @llvm.exp2.v2bf16(<2 x bfloat> %in)377 ret <2 x bfloat> %exp2378}379 380declare float @llvm.exp2.f32(float %val)381 382declare <2 x float> @llvm.exp2.v2f32(<2 x float> %val)383 384declare half @llvm.exp2.f16(half %val)385 386declare <2 x half> @llvm.exp2.v2f16(<2 x half> %val)387 388declare bfloat @llvm.exp2.bf16(bfloat %val)389 390declare <2 x bfloat> @llvm.exp2.v2bf16(<2 x bfloat> %val)391 392attributes #0 = {"denormal-fp-math"="preserve-sign"}393