d494eb0fa3
When assigning numbers to registers, skip any with neither uses nor defs. This is will not have any impact at all on the final SASS but it makes for slightly more readable PTX. This change should also ensure that future minor changes are less likely to cause noisy diffs in register numbering.
160 lines
4.2 KiB
LLVM
160 lines
4.2 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s | FileCheck %s
|
|
; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}
|
|
|
|
target triple = "nvptx64-nvidia-cuda"
|
|
|
|
@out = addrspace(1) global i32 0, align 4
|
|
|
|
define void @foo(i32 %i) {
|
|
; CHECK-LABEL: foo(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<2>;
|
|
; CHECK-NEXT: .reg .b32 %r<2>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0: // %entry
|
|
; CHECK-NEXT: ld.param.b32 %r1, [foo_param_0];
|
|
; CHECK-NEXT: setp.gt.u32 %p1, %r1, 3;
|
|
; CHECK-NEXT: @%p1 bra $L__BB0_6;
|
|
; CHECK-NEXT: // %bb.1: // %entry
|
|
; CHECK-NEXT: $L_brx_0: .branchtargets
|
|
; CHECK-NEXT: $L__BB0_2,
|
|
; CHECK-NEXT: $L__BB0_3,
|
|
; CHECK-NEXT: $L__BB0_4,
|
|
; CHECK-NEXT: $L__BB0_5;
|
|
; CHECK-NEXT: brx.idx %r1, $L_brx_0;
|
|
; CHECK-NEXT: $L__BB0_2: // %case0
|
|
; CHECK-NEXT: st.global.b32 [out], 0;
|
|
; CHECK-NEXT: bra.uni $L__BB0_6;
|
|
; CHECK-NEXT: $L__BB0_4: // %case2
|
|
; CHECK-NEXT: st.global.b32 [out], 2;
|
|
; CHECK-NEXT: bra.uni $L__BB0_6;
|
|
; CHECK-NEXT: $L__BB0_5: // %case3
|
|
; CHECK-NEXT: st.global.b32 [out], 3;
|
|
; CHECK-NEXT: bra.uni $L__BB0_6;
|
|
; CHECK-NEXT: $L__BB0_3: // %case1
|
|
; CHECK-NEXT: st.global.b32 [out], 1;
|
|
; CHECK-NEXT: $L__BB0_6: // %end
|
|
; CHECK-NEXT: ret;
|
|
entry:
|
|
switch i32 %i, label %end [
|
|
i32 0, label %case0
|
|
i32 1, label %case1
|
|
i32 2, label %case2
|
|
i32 3, label %case3
|
|
]
|
|
|
|
case0:
|
|
store i32 0, ptr addrspace(1) @out, align 4
|
|
br label %end
|
|
|
|
case1:
|
|
store i32 1, ptr addrspace(1) @out, align 4
|
|
br label %end
|
|
|
|
case2:
|
|
store i32 2, ptr addrspace(1) @out, align 4
|
|
br label %end
|
|
|
|
case3:
|
|
store i32 3, ptr addrspace(1) @out, align 4
|
|
br label %end
|
|
|
|
end:
|
|
ret void
|
|
}
|
|
|
|
|
|
define i32 @test2(i32 %tmp158) {
|
|
; CHECK-LABEL: test2(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .pred %p<6>;
|
|
; CHECK-NEXT: .reg .b32 %r<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0: // %entry
|
|
; CHECK-NEXT: ld.param.b32 %r1, [test2_param_0];
|
|
; CHECK-NEXT: setp.gt.s32 %p1, %r1, 119;
|
|
; CHECK-NEXT: @%p1 bra $L__BB1_4;
|
|
; CHECK-NEXT: // %bb.1: // %entry
|
|
; CHECK-NEXT: setp.lt.u32 %p4, %r1, 6;
|
|
; CHECK-NEXT: @%p4 bra $L__BB1_3;
|
|
; CHECK-NEXT: // %bb.2: // %entry
|
|
; CHECK-NEXT: setp.lt.s32 %p5, %r1, -2147483645;
|
|
; CHECK-NEXT: @%p5 bra $L__BB1_3;
|
|
; CHECK-NEXT: bra.uni $L__BB1_6;
|
|
; CHECK-NEXT: $L__BB1_4: // %entry
|
|
; CHECK-NEXT: add.s32 %r2, %r1, -120;
|
|
; CHECK-NEXT: setp.gt.u32 %p2, %r2, 5;
|
|
; CHECK-NEXT: @%p2 bra $L__BB1_5;
|
|
; CHECK-NEXT: // %bb.12: // %entry
|
|
; CHECK-NEXT: $L_brx_0: .branchtargets
|
|
; CHECK-NEXT: $L__BB1_3,
|
|
; CHECK-NEXT: $L__BB1_7,
|
|
; CHECK-NEXT: $L__BB1_8,
|
|
; CHECK-NEXT: $L__BB1_9,
|
|
; CHECK-NEXT: $L__BB1_10,
|
|
; CHECK-NEXT: $L__BB1_11;
|
|
; CHECK-NEXT: brx.idx %r2, $L_brx_0;
|
|
; CHECK-NEXT: $L__BB1_7: // %bb339
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 12;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_5: // %entry
|
|
; CHECK-NEXT: setp.eq.b32 %p3, %r1, 1024;
|
|
; CHECK-NEXT: @%p3 bra $L__BB1_3;
|
|
; CHECK-NEXT: bra.uni $L__BB1_6;
|
|
; CHECK-NEXT: $L__BB1_3: // %bb338
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 11;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_10: // %bb342
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 15;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_6: // %bb336
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 10;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_8: // %bb340
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 13;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_9: // %bb341
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 14;
|
|
; CHECK-NEXT: ret;
|
|
; CHECK-NEXT: $L__BB1_11: // %bb343
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], 18;
|
|
; CHECK-NEXT: ret;
|
|
entry:
|
|
switch i32 %tmp158, label %bb336 [
|
|
i32 -2147483648, label %bb338
|
|
i32 -2147483647, label %bb338
|
|
i32 -2147483646, label %bb338
|
|
i32 120, label %bb338
|
|
i32 121, label %bb339
|
|
i32 122, label %bb340
|
|
i32 123, label %bb341
|
|
i32 124, label %bb342
|
|
i32 125, label %bb343
|
|
i32 126, label %bb336
|
|
i32 1024, label %bb338
|
|
i32 0, label %bb338
|
|
i32 1, label %bb338
|
|
i32 2, label %bb338
|
|
i32 3, label %bb338
|
|
i32 4, label %bb338
|
|
i32 5, label %bb338
|
|
]
|
|
|
|
bb336:
|
|
ret i32 10
|
|
bb338:
|
|
ret i32 11
|
|
bb339:
|
|
ret i32 12
|
|
bb340:
|
|
ret i32 13
|
|
bb341:
|
|
ret i32 14
|
|
bb342:
|
|
ret i32 15
|
|
bb343:
|
|
ret i32 18
|
|
|
|
}
|