160 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s | FileCheck %s3; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7@out = addrspace(1) global i32 0, align 48 9define void @foo(i32 %i) {10; CHECK-LABEL: foo(11; CHECK: {12; CHECK-NEXT: .reg .pred %p<2>;13; CHECK-NEXT: .reg .b32 %r<2>;14; CHECK-EMPTY:15; CHECK-NEXT: // %bb.0: // %entry16; CHECK-NEXT: ld.param.b32 %r1, [foo_param_0];17; CHECK-NEXT: setp.gt.u32 %p1, %r1, 3;18; CHECK-NEXT: @%p1 bra $L__BB0_6;19; CHECK-NEXT: // %bb.1: // %entry20; CHECK-NEXT: $L_brx_0: .branchtargets21; CHECK-NEXT: $L__BB0_2,22; CHECK-NEXT: $L__BB0_3,23; CHECK-NEXT: $L__BB0_4,24; CHECK-NEXT: $L__BB0_5;25; CHECK-NEXT: brx.idx %r1, $L_brx_0;26; CHECK-NEXT: $L__BB0_2: // %case027; CHECK-NEXT: st.global.b32 [out], 0;28; CHECK-NEXT: bra.uni $L__BB0_6;29; CHECK-NEXT: $L__BB0_4: // %case230; CHECK-NEXT: st.global.b32 [out], 2;31; CHECK-NEXT: bra.uni $L__BB0_6;32; CHECK-NEXT: $L__BB0_5: // %case333; CHECK-NEXT: st.global.b32 [out], 3;34; CHECK-NEXT: bra.uni $L__BB0_6;35; CHECK-NEXT: $L__BB0_3: // %case136; CHECK-NEXT: st.global.b32 [out], 1;37; CHECK-NEXT: $L__BB0_6: // %end38; CHECK-NEXT: ret;39entry:40 switch i32 %i, label %end [41 i32 0, label %case042 i32 1, label %case143 i32 2, label %case244 i32 3, label %case345 ]46 47case0:48 store i32 0, ptr addrspace(1) @out, align 449 br label %end50 51case1:52 store i32 1, ptr addrspace(1) @out, align 453 br label %end54 55case2:56 store i32 2, ptr addrspace(1) @out, align 457 br label %end58 59case3:60 store i32 3, ptr addrspace(1) @out, align 461 br label %end62 63end:64 ret void65}66 67 68define i32 @test2(i32 %tmp158) {69; CHECK-LABEL: test2(70; CHECK: {71; CHECK-NEXT: .reg .pred %p<6>;72; CHECK-NEXT: .reg .b32 %r<3>;73; CHECK-EMPTY:74; CHECK-NEXT: // %bb.0: // %entry75; CHECK-NEXT: ld.param.b32 %r1, [test2_param_0];76; CHECK-NEXT: setp.gt.s32 %p1, %r1, 119;77; CHECK-NEXT: @%p1 bra $L__BB1_4;78; CHECK-NEXT: // %bb.1: // %entry79; CHECK-NEXT: setp.lt.u32 %p4, %r1, 6;80; CHECK-NEXT: @%p4 bra $L__BB1_3;81; CHECK-NEXT: // %bb.2: // %entry82; CHECK-NEXT: setp.lt.s32 %p5, %r1, -2147483645;83; CHECK-NEXT: @%p5 bra $L__BB1_3;84; CHECK-NEXT: bra.uni $L__BB1_6;85; CHECK-NEXT: $L__BB1_4: // %entry86; CHECK-NEXT: add.s32 %r2, %r1, -120;87; CHECK-NEXT: setp.gt.u32 %p2, %r2, 5;88; CHECK-NEXT: @%p2 bra $L__BB1_5;89; CHECK-NEXT: // %bb.12: // %entry90; CHECK-NEXT: $L_brx_0: .branchtargets91; CHECK-NEXT: $L__BB1_3,92; CHECK-NEXT: $L__BB1_7,93; CHECK-NEXT: $L__BB1_8,94; CHECK-NEXT: $L__BB1_9,95; CHECK-NEXT: $L__BB1_10,96; CHECK-NEXT: $L__BB1_11;97; CHECK-NEXT: brx.idx %r2, $L_brx_0;98; CHECK-NEXT: $L__BB1_7: // %bb33999; CHECK-NEXT: st.param.b32 [func_retval0], 12;100; CHECK-NEXT: ret;101; CHECK-NEXT: $L__BB1_5: // %entry102; CHECK-NEXT: setp.eq.b32 %p3, %r1, 1024;103; CHECK-NEXT: @%p3 bra $L__BB1_3;104; CHECK-NEXT: bra.uni $L__BB1_6;105; CHECK-NEXT: $L__BB1_3: // %bb338106; CHECK-NEXT: st.param.b32 [func_retval0], 11;107; CHECK-NEXT: ret;108; CHECK-NEXT: $L__BB1_10: // %bb342109; CHECK-NEXT: st.param.b32 [func_retval0], 15;110; CHECK-NEXT: ret;111; CHECK-NEXT: $L__BB1_6: // %bb336112; CHECK-NEXT: st.param.b32 [func_retval0], 10;113; CHECK-NEXT: ret;114; CHECK-NEXT: $L__BB1_8: // %bb340115; CHECK-NEXT: st.param.b32 [func_retval0], 13;116; CHECK-NEXT: ret;117; CHECK-NEXT: $L__BB1_9: // %bb341118; CHECK-NEXT: st.param.b32 [func_retval0], 14;119; CHECK-NEXT: ret;120; CHECK-NEXT: $L__BB1_11: // %bb343121; CHECK-NEXT: st.param.b32 [func_retval0], 18;122; CHECK-NEXT: ret;123entry:124 switch i32 %tmp158, label %bb336 [125 i32 -2147483648, label %bb338126 i32 -2147483647, label %bb338127 i32 -2147483646, label %bb338128 i32 120, label %bb338129 i32 121, label %bb339130 i32 122, label %bb340131 i32 123, label %bb341132 i32 124, label %bb342133 i32 125, label %bb343134 i32 126, label %bb336135 i32 1024, label %bb338136 i32 0, label %bb338137 i32 1, label %bb338138 i32 2, label %bb338139 i32 3, label %bb338140 i32 4, label %bb338141 i32 5, label %bb338142 ]143 144bb336:145 ret i32 10146bb338:147 ret i32 11148bb339:149 ret i32 12150bb340:151 ret i32 13152bb341:153 ret i32 14154bb342:155 ret i32 15156bb343:157 ret i32 18158 159}160