222 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=CHECK,CHECK-O33; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefixes=CHECK,CHECK-O04; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}6 7target triple = "nvptx64-nvidia-cuda"8 9define i32 @bfe0(i32 %a) {10; CHECK-LABEL: bfe0(11; CHECK: {12; CHECK-NEXT: .reg .b32 %r<3>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b32 %r1, [bfe0_param_0];16; CHECK-NEXT: bfe.u32 %r2, %r1, 4, 4;17; CHECK-NEXT: st.param.b32 [func_retval0], %r2;18; CHECK-NEXT: ret;19 %val0 = ashr i32 %a, 420 %val1 = and i32 %val0, 1521 ret i32 %val122}23 24define i32 @bfe1(i32 %a) {25; CHECK-LABEL: bfe1(26; CHECK: {27; CHECK-NEXT: .reg .b32 %r<3>;28; CHECK-EMPTY:29; CHECK-NEXT: // %bb.0:30; CHECK-NEXT: ld.param.b32 %r1, [bfe1_param_0];31; CHECK-NEXT: bfe.u32 %r2, %r1, 3, 3;32; CHECK-NEXT: st.param.b32 [func_retval0], %r2;33; CHECK-NEXT: ret;34 %val0 = ashr i32 %a, 335 %val1 = and i32 %val0, 736 ret i32 %val137}38 39define i32 @bfe2(i32 %a) {40; CHECK-LABEL: bfe2(41; CHECK: {42; CHECK-NEXT: .reg .b32 %r<3>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b32 %r1, [bfe2_param_0];46; CHECK-NEXT: bfe.u32 %r2, %r1, 5, 3;47; CHECK-NEXT: st.param.b32 [func_retval0], %r2;48; CHECK-NEXT: ret;49 %val0 = ashr i32 %a, 550 %val1 = and i32 %val0, 751 ret i32 %val152}53 54define i32 @no_bfe_on_32bit_overflow(i32 %a) {55; CHECK-LABEL: no_bfe_on_32bit_overflow(56; CHECK: {57; CHECK-NEXT: .reg .b32 %r<4>;58; CHECK-EMPTY:59; CHECK-NEXT: // %bb.0:60; CHECK-NEXT: ld.param.b32 %r1, [no_bfe_on_32bit_overflow_param_0];61; CHECK-NEXT: shr.s32 %r2, %r1, 31;62; CHECK-NEXT: and.b32 %r3, %r2, 15;63; CHECK-NEXT: st.param.b32 [func_retval0], %r3;64; CHECK-NEXT: ret;65 %val0 = ashr i32 %a, 3166 %val1 = and i32 %val0, 1567 ret i32 %val168}69 70define i32 @no_bfe_on_32bit_overflow_shr_and_pair(i32 %a) {71; CHECK-LABEL: no_bfe_on_32bit_overflow_shr_and_pair(72; CHECK: {73; CHECK-NEXT: .reg .b32 %r<4>;74; CHECK-EMPTY:75; CHECK-NEXT: // %bb.0:76; CHECK-NEXT: ld.param.b32 %r1, [no_bfe_on_32bit_overflow_shr_and_pair_param_0];77; CHECK-NEXT: shr.s32 %r2, %r1, 31;78; CHECK-NEXT: and.b32 %r3, %r2, 15;79; CHECK-NEXT: st.param.b32 [func_retval0], %r3;80; CHECK-NEXT: ret;81 %val0 = ashr i32 %a, 3182 %val1 = and i32 %val0, 1583 ret i32 %val184}85 86define i64 @no_bfe_on_64bit_overflow(i64 %a) {87; CHECK-LABEL: no_bfe_on_64bit_overflow(88; CHECK: {89; CHECK-NEXT: .reg .b64 %rd<4>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b64 %rd1, [no_bfe_on_64bit_overflow_param_0];93; CHECK-NEXT: shr.s64 %rd2, %rd1, 63;94; CHECK-NEXT: and.b64 %rd3, %rd2, 7;95; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;96; CHECK-NEXT: ret;97 %val0 = ashr i64 %a, 6398 %val1 = and i64 %val0, 799 ret i64 %val1100}101 102define i64 @no_bfe_on_64bit_overflow_shr_and_pair(i64 %a) {103; CHECK-LABEL: no_bfe_on_64bit_overflow_shr_and_pair(104; CHECK: {105; CHECK-NEXT: .reg .b64 %rd<4>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b64 %rd1, [no_bfe_on_64bit_overflow_shr_and_pair_param_0];109; CHECK-NEXT: shr.s64 %rd2, %rd1, 63;110; CHECK-NEXT: and.b64 %rd3, %rd2, 7;111; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;112; CHECK-NEXT: ret;113 %val0 = ashr i64 %a, 63114 %val1 = and i64 %val0, 7115 ret i64 %val1116}117 118define i32 @bfe_ashr_signed_32(i32 %x) {119; CHECK-O3-LABEL: bfe_ashr_signed_32(120; CHECK-O3: {121; CHECK-O3-NEXT: .reg .b32 %r<3>;122; CHECK-O3-EMPTY:123; CHECK-O3-NEXT: // %bb.0:124; CHECK-O3-NEXT: ld.param.b16 %r1, [bfe_ashr_signed_32_param_0+2];125; CHECK-O3-NEXT: bfe.s32 %r2, %r1, 4, 12;126; CHECK-O3-NEXT: st.param.b32 [func_retval0], %r2;127; CHECK-O3-NEXT: ret;128;129; CHECK-O0-LABEL: bfe_ashr_signed_32(130; CHECK-O0: {131; CHECK-O0-NEXT: .reg .b32 %r<3>;132; CHECK-O0-EMPTY:133; CHECK-O0-NEXT: // %bb.0:134; CHECK-O0-NEXT: ld.param.b32 %r1, [bfe_ashr_signed_32_param_0];135; CHECK-O0-NEXT: bfe.s32 %r2, %r1, 20, 12;136; CHECK-O0-NEXT: st.param.b32 [func_retval0], %r2;137; CHECK-O0-NEXT: ret;138 %and = and i32 %x, -65536139 %shr = ashr exact i32 %and, 20140 ret i32 %shr141}142 143define i32 @bfe_ashr_unsigned_32(i32 %x) {144; CHECK-LABEL: bfe_ashr_unsigned_32(145; CHECK: {146; CHECK-NEXT: .reg .b32 %r<3>;147; CHECK-EMPTY:148; CHECK-NEXT: // %bb.0:149; CHECK-NEXT: ld.param.b32 %r1, [bfe_ashr_unsigned_32_param_0];150; CHECK-NEXT: bfe.u32 %r2, %r1, 5, 6;151; CHECK-NEXT: st.param.b32 [func_retval0], %r2;152; CHECK-NEXT: ret;153 %and = and i32 %x, 2047154 %shr = ashr exact i32 %and, 5155 ret i32 %shr156}157 158define i64 @bfe_ashr_signed_64(i64 %x) {159; CHECK-LABEL: bfe_ashr_signed_64(160; CHECK: {161; CHECK-NEXT: .reg .b64 %rd<3>;162; CHECK-EMPTY:163; CHECK-NEXT: // %bb.0:164; CHECK-NEXT: ld.param.b64 %rd1, [bfe_ashr_signed_64_param_0];165; CHECK-NEXT: bfe.s64 %rd2, %rd1, 16, 48;166; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;167; CHECK-NEXT: ret;168 %and = and i64 %x, -65536169 %shr = ashr exact i64 %and, 16170 ret i64 %shr171}172 173define i64 @bfe_ashr_unsigned_64(i64 %x) {174; CHECK-LABEL: bfe_ashr_unsigned_64(175; CHECK: {176; CHECK-NEXT: .reg .b64 %rd<3>;177; CHECK-EMPTY:178; CHECK-NEXT: // %bb.0:179; CHECK-NEXT: ld.param.b64 %rd1, [bfe_ashr_unsigned_64_param_0];180; CHECK-NEXT: bfe.u64 %rd2, %rd1, 5, 6;181; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;182; CHECK-NEXT: ret;183 %and = and i64 %x, 2047184 %shr = ashr exact i64 %and, 5185 ret i64 %shr186}187 188define i32 @bfe3(i128 %a) {189; CHECK-LABEL: bfe3(190; CHECK: {191; CHECK-NEXT: .reg .b32 %r<3>;192; CHECK-NEXT: .reg .b64 %rd<3>;193; CHECK-EMPTY:194; CHECK-NEXT: // %bb.0:195; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [bfe3_param_0];196; CHECK-NEXT: cvt.u32.u64 %r1, %rd1;197; CHECK-NEXT: bfe.s32 %r2, %r1, 15, 17;198; CHECK-NEXT: st.param.b32 [func_retval0], %r2;199; CHECK-NEXT: ret;200 %trunc = trunc i128 %a to i32201 %and = and i32 %trunc, -32768202 %shr = ashr exact i32 %and, 15203 ret i32 %shr204}205 206define i64 @bfe4(i128 %a) {207; CHECK-LABEL: bfe4(208; CHECK: {209; CHECK-NEXT: .reg .b64 %rd<4>;210; CHECK-EMPTY:211; CHECK-NEXT: // %bb.0:212; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [bfe4_param_0];213; CHECK-NEXT: bfe.s64 %rd3, %rd1, 17, 47;214; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;215; CHECK-NEXT: ret;216 %trunc = trunc i128 %a to i64217 %and = and i64 %trunc, -131072218 %shr = ashr exact i64 %and, 17219 ret i64 %shr220}221 222