40 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 --debug-counter=dagcombine=0 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -march=nvptx64 --debug-counter=dagcombine=0 | %ptxas-verify %}4 5; REQUIRES: asserts6; asserts are required for --debug-counter=dagcombine=0 to have the intended7; effect of disabling DAG combines, which exposes the bug. When combines are8; enabled the bug does not occur.9 10%struct.1float = type <{ [1 x float] }>11 12declare i32 @callee(%struct.1float %a)13 14define i32 @test(%struct.1float alignstack(32) %data) {15; CHECK-LABEL: test(16; CHECK: {17; CHECK-NEXT: .reg .b32 %r<6>;18; CHECK-EMPTY:19; CHECK-NEXT: // %bb.0:20; CHECK-NEXT: ld.param.b32 %r1, [test_param_0];21; CHECK-NEXT: { // callseq 0, 022; CHECK-NEXT: .param .align 1 .b8 param0[4];23; CHECK-NEXT: .param .b32 retval0;24; CHECK-NEXT: st.param.b8 [param0], %r1;25; CHECK-NEXT: shr.u32 %r2, %r1, 8;26; CHECK-NEXT: st.param.b8 [param0+1], %r2;27; CHECK-NEXT: shr.u32 %r3, %r1, 16;28; CHECK-NEXT: st.param.b8 [param0+2], %r3;29; CHECK-NEXT: shr.u32 %r4, %r3, 8;30; CHECK-NEXT: st.param.b8 [param0+3], %r4;31; CHECK-NEXT: call.uni (retval0), callee, (param0);32; CHECK-NEXT: ld.param.b32 %r5, [retval0];33; CHECK-NEXT: } // callseq 034; CHECK-NEXT: st.param.b32 [func_retval0], %r5;35; CHECK-NEXT: ret;36 37 %1 = call i32 @callee(%struct.1float %data)38 ret i32 %139}40