219 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 -nvptx-approx-log2f32 | FileCheck --check-prefixes=CHECK %s3; RUN: %if ptxas-sm_50 && ptxas-isa-3.2 %{ llc < %s -mcpu=sm_50 -mattr=+ptx32 -nvptx-approx-log2f32 | %ptxas-verify -arch=sm_50 %}4target triple = "nvptx64-nvidia-cuda"5 6; CHECK-LABEL: log2_test7define float @log2_test(float %in) {8; CHECK-LABEL: log2_test(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0: // %entry13; CHECK-NEXT: ld.param.b32 %r1, [log2_test_param_0];14; CHECK-NEXT: lg2.approx.f32 %r2, %r1;15; CHECK-NEXT: st.param.b32 [func_retval0], %r2;16; CHECK-NEXT: ret;17entry:18 %log2 = call float @llvm.log2.f32(float %in)19 ret float %log220}21 22; CHECK-LABEL: log2_ftz_test23define float @log2_ftz_test(float %in) #0 {24; CHECK-LABEL: log2_ftz_test(25; CHECK: {26; CHECK-NEXT: .reg .b32 %r<3>;27; CHECK-EMPTY:28; CHECK-NEXT: // %bb.0: // %entry29; CHECK-NEXT: ld.param.b32 %r1, [log2_ftz_test_param_0];30; CHECK-NEXT: lg2.approx.ftz.f32 %r2, %r1;31; CHECK-NEXT: st.param.b32 [func_retval0], %r2;32; CHECK-NEXT: ret;33entry:34 %log2 = call float @llvm.log2.f32(float %in)35 ret float %log236}37 38; CHECK-LABEL: log2_test_v39define <2 x float> @log2_test_v(<2 x float> %in) {40; CHECK-LABEL: log2_test_v(41; CHECK: {42; CHECK-NEXT: .reg .b32 %r<5>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0: // %entry45; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [log2_test_v_param_0];46; CHECK-NEXT: lg2.approx.f32 %r3, %r2;47; CHECK-NEXT: lg2.approx.f32 %r4, %r1;48; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};49; CHECK-NEXT: ret;50entry:51 %log2 = call <2 x float> @llvm.log2.v2f32(<2 x float> %in)52 ret <2 x float> %log253}54 55; --- f16 ---56 57; CHECK-LABEL: log2_f16_test58define half @log2_f16_test(half %in) {59; CHECK-LABEL: log2_f16_test(60; CHECK: {61; CHECK-NEXT: .reg .b16 %rs<3>;62; CHECK-NEXT: .reg .b32 %r<3>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0: // %entry65; CHECK-NEXT: ld.param.b16 %rs1, [log2_f16_test_param_0];66; CHECK-NEXT: cvt.f32.f16 %r1, %rs1;67; CHECK-NEXT: lg2.approx.f32 %r2, %r1;68; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;69; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;70; CHECK-NEXT: ret;71entry:72 %log2 = call half @llvm.log2.f16(half %in)73 ret half %log274}75 76; CHECK-LABEL: log2_f16_ftz_test77define half @log2_f16_ftz_test(half %in) #0 {78; CHECK-LABEL: log2_f16_ftz_test(79; CHECK: {80; CHECK-NEXT: .reg .b16 %rs<3>;81; CHECK-NEXT: .reg .b32 %r<3>;82; CHECK-EMPTY:83; CHECK-NEXT: // %bb.0: // %entry84; CHECK-NEXT: ld.param.b16 %rs1, [log2_f16_ftz_test_param_0];85; CHECK-NEXT: cvt.ftz.f32.f16 %r1, %rs1;86; CHECK-NEXT: lg2.approx.ftz.f32 %r2, %r1;87; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;88; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;89; CHECK-NEXT: ret;90entry:91 %log2 = call half @llvm.log2.f16(half %in)92 ret half %log293}94 95; CHECK-LABEL: log2_f16_test_v96define <2 x half> @log2_f16_test_v(<2 x half> %in) {97; CHECK-LABEL: log2_f16_test_v(98; CHECK: {99; CHECK-NEXT: .reg .b16 %rs<5>;100; CHECK-NEXT: .reg .b32 %r<6>;101; CHECK-EMPTY:102; CHECK-NEXT: // %bb.0: // %entry103; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [log2_f16_test_v_param_0];104; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;105; CHECK-NEXT: lg2.approx.f32 %r2, %r1;106; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r2;107; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;108; CHECK-NEXT: lg2.approx.f32 %r4, %r3;109; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r4;110; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3};111; CHECK-NEXT: st.param.b32 [func_retval0], %r5;112; CHECK-NEXT: ret;113entry:114 %log2 = call <2 x half> @llvm.log2.v2f16(<2 x half> %in)115 ret <2 x half> %log2116}117 118; --- bf16 ---119 120; CHECK-LABEL: log2_bf16_test121define bfloat @log2_bf16_test(bfloat %in) {122; CHECK-LABEL: log2_bf16_test(123; CHECK: {124; CHECK-NEXT: .reg .pred %p<2>;125; CHECK-NEXT: .reg .b32 %r<10>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0: // %entry128; CHECK-NEXT: ld.param.b16 %r1, [log2_bf16_test_param_0];129; CHECK-NEXT: shl.b32 %r2, %r1, 16;130; CHECK-NEXT: lg2.approx.f32 %r3, %r2;131; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;132; CHECK-NEXT: add.s32 %r5, %r4, %r3;133; CHECK-NEXT: add.s32 %r6, %r5, 32767;134; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;135; CHECK-NEXT: or.b32 %r7, %r3, 4194304;136; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;137; CHECK-NEXT: shr.u32 %r9, %r8, 16;138; CHECK-NEXT: st.param.b16 [func_retval0], %r9;139; CHECK-NEXT: ret;140entry:141 %log2 = call bfloat @llvm.log2.bf16(bfloat %in)142 ret bfloat %log2143}144 145; CHECK-LABEL: log2_bf16_ftz_test146define bfloat @log2_bf16_ftz_test(bfloat %in) #0 {147; CHECK-LABEL: log2_bf16_ftz_test(148; CHECK: {149; CHECK-NEXT: .reg .pred %p<2>;150; CHECK-NEXT: .reg .b32 %r<10>;151; CHECK-EMPTY:152; CHECK-NEXT: // %bb.0: // %entry153; CHECK-NEXT: ld.param.b16 %r1, [log2_bf16_ftz_test_param_0];154; CHECK-NEXT: shl.b32 %r2, %r1, 16;155; CHECK-NEXT: lg2.approx.ftz.f32 %r3, %r2;156; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;157; CHECK-NEXT: add.s32 %r5, %r4, %r3;158; CHECK-NEXT: add.s32 %r6, %r5, 32767;159; CHECK-NEXT: setp.nan.ftz.f32 %p1, %r3, %r3;160; CHECK-NEXT: or.b32 %r7, %r3, 4194304;161; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;162; CHECK-NEXT: shr.u32 %r9, %r8, 16;163; CHECK-NEXT: st.param.b16 [func_retval0], %r9;164; CHECK-NEXT: ret;165entry:166 %log2 = call bfloat @llvm.log2.bf16(bfloat %in)167 ret bfloat %log2168}169 170; CHECK-LABEL: log2_bf16_test_v171define <2 x bfloat> @log2_bf16_test_v(<2 x bfloat> %in) {172; CHECK-LABEL: log2_bf16_test_v(173; CHECK: {174; CHECK-NEXT: .reg .pred %p<3>;175; CHECK-NEXT: .reg .b16 %rs<3>;176; CHECK-NEXT: .reg .b32 %r<18>;177; CHECK-EMPTY:178; CHECK-NEXT: // %bb.0: // %entry179; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [log2_bf16_test_v_param_0];180; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;181; CHECK-NEXT: shl.b32 %r2, %r1, 16;182; CHECK-NEXT: lg2.approx.f32 %r3, %r2;183; CHECK-NEXT: bfe.u32 %r4, %r3, 16, 1;184; CHECK-NEXT: add.s32 %r5, %r4, %r3;185; CHECK-NEXT: add.s32 %r6, %r5, 32767;186; CHECK-NEXT: setp.nan.f32 %p1, %r3, %r3;187; CHECK-NEXT: or.b32 %r7, %r3, 4194304;188; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1;189; CHECK-NEXT: cvt.u32.u16 %r9, %rs1;190; CHECK-NEXT: shl.b32 %r10, %r9, 16;191; CHECK-NEXT: lg2.approx.f32 %r11, %r10;192; CHECK-NEXT: bfe.u32 %r12, %r11, 16, 1;193; CHECK-NEXT: add.s32 %r13, %r12, %r11;194; CHECK-NEXT: add.s32 %r14, %r13, 32767;195; CHECK-NEXT: setp.nan.f32 %p2, %r11, %r11;196; CHECK-NEXT: or.b32 %r15, %r11, 4194304;197; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p2;198; CHECK-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;199; CHECK-NEXT: st.param.b32 [func_retval0], %r17;200; CHECK-NEXT: ret;201entry:202 %log2 = call <2 x bfloat> @llvm.log2.v2bf16(<2 x bfloat> %in)203 ret <2 x bfloat> %log2204}205 206declare float @llvm.log2.f32(float %val)207 208declare <2 x float> @llvm.log2.v2f32(<2 x float> %val)209 210declare half @llvm.log2.f16(half %val)211 212declare <2 x half> @llvm.log2.v2f16(<2 x half> %val)213 214declare bfloat @llvm.log2.bf16(bfloat %val)215 216declare <2 x bfloat> @llvm.log2.v2bf16(<2 x bfloat> %val)217 218attributes #0 = {"denormal-fp-math"="preserve-sign"}219