1*ec3525f7SKevin McAfee; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2*ec3525f7SKevin McAfee; RUN: llc < %s -march=nvptx64 --debug-counter=dagcombine=0 | FileCheck %s 3*ec3525f7SKevin McAfee; RUN: %if ptxas %{ llc < %s -march=nvptx64 --debug-counter=dagcombine=0 | %ptxas-verify %} 4*ec3525f7SKevin McAfee 5*ec3525f7SKevin McAfee; REQUIRES: asserts 6*ec3525f7SKevin McAfee; asserts are required for --debug-counter=dagcombine=0 to have the intended 7*ec3525f7SKevin McAfee; effect of disabling DAG combines, which exposes the bug. When combines are 8*ec3525f7SKevin McAfee; enabled the bug does not occur. 9*ec3525f7SKevin McAfee 10*ec3525f7SKevin McAfee%struct.1float = type <{ [1 x float] }> 11*ec3525f7SKevin McAfee 12*ec3525f7SKevin McAfeedeclare i32 @callee(%struct.1float %a) 13*ec3525f7SKevin McAfee 14*ec3525f7SKevin McAfeedefine i32 @test(%struct.1float alignstack(32) %data) { 15*ec3525f7SKevin McAfee; CHECK-LABEL: test( 16*ec3525f7SKevin McAfee; CHECK: { 17*ec3525f7SKevin McAfee; CHECK-NEXT: .reg .b32 %r<18>; 18*ec3525f7SKevin McAfee; CHECK-NEXT: .reg .f32 %f<2>; 19*ec3525f7SKevin McAfee; CHECK-EMPTY: 20*ec3525f7SKevin McAfee; CHECK-NEXT: // %bb.0: 21*ec3525f7SKevin McAfee; CHECK-NEXT: ld.param.u8 %r1, [test_param_0+1]; 22*ec3525f7SKevin McAfee; CHECK-NEXT: shl.b32 %r2, %r1, 8; 23*ec3525f7SKevin McAfee; CHECK-NEXT: ld.param.u8 %r3, [test_param_0]; 24*ec3525f7SKevin McAfee; CHECK-NEXT: or.b32 %r4, %r2, %r3; 25*ec3525f7SKevin McAfee; CHECK-NEXT: ld.param.u8 %r5, [test_param_0+3]; 26*ec3525f7SKevin McAfee; CHECK-NEXT: shl.b32 %r6, %r5, 8; 27*ec3525f7SKevin McAfee; CHECK-NEXT: ld.param.u8 %r7, [test_param_0+2]; 28*ec3525f7SKevin McAfee; CHECK-NEXT: or.b32 %r8, %r6, %r7; 29*ec3525f7SKevin McAfee; CHECK-NEXT: shl.b32 %r9, %r8, 16; 30*ec3525f7SKevin McAfee; CHECK-NEXT: or.b32 %r17, %r9, %r4; 31*ec3525f7SKevin McAfee; CHECK-NEXT: mov.b32 %f1, %r17; 32*ec3525f7SKevin McAfee; CHECK-NEXT: shr.u32 %r12, %r17, 8; 33*ec3525f7SKevin McAfee; CHECK-NEXT: shr.u32 %r13, %r17, 16; 34*ec3525f7SKevin McAfee; CHECK-NEXT: shr.u32 %r14, %r17, 24; 35*ec3525f7SKevin McAfee; CHECK-NEXT: { // callseq 0, 0 36*ec3525f7SKevin McAfee; CHECK-NEXT: .param .align 1 .b8 param0[4]; 37*ec3525f7SKevin McAfee; CHECK-NEXT: st.param.b8 [param0], %r17; 38*ec3525f7SKevin McAfee; CHECK-NEXT: st.param.b8 [param0+1], %r12; 39*ec3525f7SKevin McAfee; CHECK-NEXT: st.param.b8 [param0+2], %r13; 40*ec3525f7SKevin McAfee; CHECK-NEXT: st.param.b8 [param0+3], %r14; 41*ec3525f7SKevin McAfee; CHECK-NEXT: .param .b32 retval0; 42*ec3525f7SKevin McAfee; CHECK-NEXT: call.uni (retval0), 43*ec3525f7SKevin McAfee; CHECK-NEXT: callee, 44*ec3525f7SKevin McAfee; CHECK-NEXT: ( 45*ec3525f7SKevin McAfee; CHECK-NEXT: param0 46*ec3525f7SKevin McAfee; CHECK-NEXT: ); 47*ec3525f7SKevin McAfee; CHECK-NEXT: ld.param.b32 %r15, [retval0]; 48*ec3525f7SKevin McAfee; CHECK-NEXT: } // callseq 0 49*ec3525f7SKevin McAfee; CHECK-NEXT: st.param.b32 [func_retval0], %r15; 50*ec3525f7SKevin McAfee; CHECK-NEXT: ret; 51*ec3525f7SKevin McAfee 52*ec3525f7SKevin McAfee %1 = call i32 @callee(%struct.1float %data) 53*ec3525f7SKevin McAfee ret i32 %1 54*ec3525f7SKevin McAfee} 55