1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s | FileCheck %s 3; RUN: %if ptxas %{ llc < %s | %ptxas-verify %} 4 5target triple = "nvptx64-nvidia-cuda" 6 7@out = addrspace(1) global i32 0, align 4 8 9define void @foo(i32 %i) { 10; CHECK-LABEL: foo( 11; CHECK: { 12; CHECK-NEXT: .reg .pred %p<2>; 13; CHECK-NEXT: .reg .b32 %r<7>; 14; CHECK-EMPTY: 15; CHECK-NEXT: // %bb.0: // %entry 16; CHECK-NEXT: ld.param.u32 %r2, [foo_param_0]; 17; CHECK-NEXT: setp.gt.u32 %p1, %r2, 3; 18; CHECK-NEXT: @%p1 bra $L__BB0_6; 19; CHECK-NEXT: // %bb.1: // %entry 20; CHECK-NEXT: $L_brx_0: .branchtargets 21; 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 %r2, $L_brx_0; 26; CHECK-NEXT: $L__BB0_2: // %case0 27; CHECK-NEXT: mov.b32 %r6, 0; 28; CHECK-NEXT: st.global.u32 [out], %r6; 29; CHECK-NEXT: bra.uni $L__BB0_6; 30; CHECK-NEXT: $L__BB0_4: // %case2 31; CHECK-NEXT: mov.b32 %r4, 2; 32; CHECK-NEXT: st.global.u32 [out], %r4; 33; CHECK-NEXT: bra.uni $L__BB0_6; 34; CHECK-NEXT: $L__BB0_5: // %case3 35; CHECK-NEXT: mov.b32 %r3, 3; 36; CHECK-NEXT: st.global.u32 [out], %r3; 37; CHECK-NEXT: bra.uni $L__BB0_6; 38; CHECK-NEXT: $L__BB0_3: // %case1 39; CHECK-NEXT: mov.b32 %r5, 1; 40; CHECK-NEXT: st.global.u32 [out], %r5; 41; CHECK-NEXT: $L__BB0_6: // %end 42; CHECK-NEXT: ret; 43entry: 44 switch i32 %i, label %end [ 45 i32 0, label %case0 46 i32 1, label %case1 47 i32 2, label %case2 48 i32 3, label %case3 49 ] 50 51case0: 52 store i32 0, ptr addrspace(1) @out, align 4 53 br label %end 54 55case1: 56 store i32 1, ptr addrspace(1) @out, align 4 57 br label %end 58 59case2: 60 store i32 2, ptr addrspace(1) @out, align 4 61 br label %end 62 63case3: 64 store i32 3, ptr addrspace(1) @out, align 4 65 br label %end 66 67end: 68 ret void 69} 70 71 72define i32 @test2(i32 %tmp158) { 73; CHECK-LABEL: test2( 74; CHECK: { 75; CHECK-NEXT: .reg .pred %p<6>; 76; CHECK-NEXT: .reg .b32 %r<10>; 77; CHECK-EMPTY: 78; CHECK-NEXT: // %bb.0: // %entry 79; CHECK-NEXT: ld.param.u32 %r1, [test2_param_0]; 80; CHECK-NEXT: setp.gt.s32 %p1, %r1, 119; 81; CHECK-NEXT: @%p1 bra $L__BB1_4; 82; CHECK-NEXT: // %bb.1: // %entry 83; CHECK-NEXT: setp.lt.u32 %p4, %r1, 6; 84; CHECK-NEXT: @%p4 bra $L__BB1_3; 85; CHECK-NEXT: // %bb.2: // %entry 86; CHECK-NEXT: setp.lt.s32 %p5, %r1, -2147483645; 87; CHECK-NEXT: @%p5 bra $L__BB1_3; 88; CHECK-NEXT: bra.uni $L__BB1_6; 89; CHECK-NEXT: $L__BB1_4: // %entry 90; CHECK-NEXT: add.s32 %r2, %r1, -120; 91; CHECK-NEXT: setp.gt.u32 %p2, %r2, 5; 92; CHECK-NEXT: @%p2 bra $L__BB1_5; 93; CHECK-NEXT: // %bb.12: // %entry 94; CHECK-NEXT: $L_brx_0: .branchtargets 95; CHECK-NEXT: $L__BB1_3, 96; CHECK-NEXT: $L__BB1_7, 97; CHECK-NEXT: $L__BB1_8, 98; CHECK-NEXT: $L__BB1_9, 99; CHECK-NEXT: $L__BB1_10, 100; CHECK-NEXT: $L__BB1_11; 101; CHECK-NEXT: brx.idx %r2, $L_brx_0; 102; CHECK-NEXT: $L__BB1_7: // %bb339 103; CHECK-NEXT: mov.b32 %r7, 12; 104; CHECK-NEXT: st.param.b32 [func_retval0], %r7; 105; CHECK-NEXT: ret; 106; CHECK-NEXT: $L__BB1_5: // %entry 107; CHECK-NEXT: setp.eq.s32 %p3, %r1, 1024; 108; CHECK-NEXT: @%p3 bra $L__BB1_3; 109; CHECK-NEXT: bra.uni $L__BB1_6; 110; CHECK-NEXT: $L__BB1_3: // %bb338 111; CHECK-NEXT: mov.b32 %r8, 11; 112; CHECK-NEXT: st.param.b32 [func_retval0], %r8; 113; CHECK-NEXT: ret; 114; CHECK-NEXT: $L__BB1_10: // %bb342 115; CHECK-NEXT: mov.b32 %r4, 15; 116; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 117; CHECK-NEXT: ret; 118; CHECK-NEXT: $L__BB1_6: // %bb336 119; CHECK-NEXT: mov.b32 %r9, 10; 120; CHECK-NEXT: st.param.b32 [func_retval0], %r9; 121; CHECK-NEXT: ret; 122; CHECK-NEXT: $L__BB1_8: // %bb340 123; CHECK-NEXT: mov.b32 %r6, 13; 124; CHECK-NEXT: st.param.b32 [func_retval0], %r6; 125; CHECK-NEXT: ret; 126; CHECK-NEXT: $L__BB1_9: // %bb341 127; CHECK-NEXT: mov.b32 %r5, 14; 128; CHECK-NEXT: st.param.b32 [func_retval0], %r5; 129; CHECK-NEXT: ret; 130; CHECK-NEXT: $L__BB1_11: // %bb343 131; CHECK-NEXT: mov.b32 %r3, 18; 132; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 133; CHECK-NEXT: ret; 134entry: 135 switch i32 %tmp158, label %bb336 [ 136 i32 -2147483648, label %bb338 137 i32 -2147483647, label %bb338 138 i32 -2147483646, label %bb338 139 i32 120, label %bb338 140 i32 121, label %bb339 141 i32 122, label %bb340 142 i32 123, label %bb341 143 i32 124, label %bb342 144 i32 125, label %bb343 145 i32 126, label %bb336 146 i32 1024, label %bb338 147 i32 0, label %bb338 148 i32 1, label %bb338 149 i32 2, label %bb338 150 i32 3, label %bb338 151 i32 4, label %bb338 152 i32 5, label %bb338 153 ] 154 155bb336: 156 ret i32 10 157bb338: 158 ret i32 11 159bb339: 160 ret i32 12 161bb340: 162 ret i32 13 163bb341: 164 ret i32 14 165bb342: 166 ret i32 15 167bb343: 168 ret i32 18 169 170} 171