| ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py |
| ; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s |
| |
| %struct.St8x4 = type { [4 x i64] } |
| |
| define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) { |
| ; CHECK-LABEL: caller_St8x4( |
| ; CHECK: { |
| ; CHECK-NEXT: .local .align 8 .b8 __local_depot0[32]; |
| ; CHECK-NEXT: .reg .b32 %SP; |
| ; CHECK-NEXT: .reg .b32 %SPL; |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-NEXT: .reg .b64 %rd<17>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: mov.u32 %SPL, __local_depot0; |
| ; CHECK-NEXT: cvta.local.u32 %SP, %SPL; |
| ; CHECK-NEXT: ld.param.u32 %r1, [caller_St8x4_param_1]; |
| ; CHECK-NEXT: add.u32 %r3, %SPL, 0; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [caller_St8x4_param_0+24]; |
| ; CHECK-NEXT: ld.param.u64 %rd2, [caller_St8x4_param_0+16]; |
| ; CHECK-NEXT: ld.param.u64 %rd3, [caller_St8x4_param_0+8]; |
| ; CHECK-NEXT: ld.param.u64 %rd4, [caller_St8x4_param_0]; |
| ; CHECK-NEXT: st.local.u64 [%r3], %rd4; |
| ; CHECK-NEXT: st.local.u64 [%r3+8], %rd3; |
| ; CHECK-NEXT: st.local.u64 [%r3+16], %rd2; |
| ; CHECK-NEXT: st.local.u64 [%r3+24], %rd1; |
| ; CHECK-NEXT: ld.u64 %rd5, [%SP+8]; |
| ; CHECK-NEXT: ld.u64 %rd6, [%SP+0]; |
| ; CHECK-NEXT: ld.u64 %rd7, [%SP+24]; |
| ; CHECK-NEXT: ld.u64 %rd8, [%SP+16]; |
| ; CHECK-NEXT: { // callseq 0, 0 |
| ; CHECK-NEXT: .param .align 16 .b8 param0[32]; |
| ; CHECK-NEXT: st.param.v2.b64 [param0+0], {%rd6, %rd5}; |
| ; CHECK-NEXT: st.param.v2.b64 [param0+16], {%rd8, %rd7}; |
| ; CHECK-NEXT: .param .align 16 .b8 retval0[32]; |
| ; CHECK-NEXT: call.uni (retval0), |
| ; CHECK-NEXT: callee_St8x4, |
| ; CHECK-NEXT: ( |
| ; CHECK-NEXT: param0 |
| ; CHECK-NEXT: ); |
| ; CHECK-NEXT: ld.param.v2.b64 {%rd9, %rd10}, [retval0+0]; |
| ; CHECK-NEXT: ld.param.v2.b64 {%rd11, %rd12}, [retval0+16]; |
| ; CHECK-NEXT: } // callseq 0 |
| ; CHECK-NEXT: st.u64 [%r1], %rd9; |
| ; CHECK-NEXT: st.u64 [%r1+8], %rd10; |
| ; CHECK-NEXT: st.u64 [%r1+16], %rd11; |
| ; CHECK-NEXT: st.u64 [%r1+24], %rd12; |
| ; CHECK-NEXT: ret; |
| %call = tail call fastcc [4 x i64] @callee_St8x4(ptr noundef nonnull byval(%struct.St8x4) align 8 %in) #2 |
| %.fca.0.extract = extractvalue [4 x i64] %call, 0 |
| %.fca.1.extract = extractvalue [4 x i64] %call, 1 |
| %.fca.2.extract = extractvalue [4 x i64] %call, 2 |
| %.fca.3.extract = extractvalue [4 x i64] %call, 3 |
| store i64 %.fca.0.extract, ptr %ret, align 8 |
| %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8 |
| store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8 |
| %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16 |
| store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 8 |
| %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24 |
| store i64 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 8 |
| ret void |
| } |
| |
| define internal fastcc [4 x i64] @callee_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in) { |
| ; CHECK-LABEL: callee_St8x4( |
| ; CHECK: // @callee_St8x4 |
| ; CHECK-NEXT: { |
| ; CHECK-NEXT: .reg .b64 %rd<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [callee_St8x4_param_0]; |
| ; CHECK-NEXT: ld.param.v2.u64 {%rd3, %rd4}, [callee_St8x4_param_0+16]; |
| ; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd1, %rd2}; |
| ; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4}; |
| ; CHECK-NEXT: ret; |
| %1 = load i64, ptr %in, align 8 |
| %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 1 |
| %2 = load i64, ptr %arrayidx.1, align 8 |
| %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 2 |
| %3 = load i64, ptr %arrayidx.2, align 8 |
| %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 3 |
| %4 = load i64, ptr %arrayidx.3, align 8 |
| %5 = insertvalue [4 x i64] poison, i64 %1, 0 |
| %6 = insertvalue [4 x i64] %5, i64 %2, 1 |
| %7 = insertvalue [4 x i64] %6, i64 %3, 2 |
| %oldret = insertvalue [4 x i64] %7, i64 %4, 3 |
| ret [4 x i64] %oldret |
| } |
| |
| define void @call_void() { |
| ; CHECK-LABEL: call_void( |
| ; CHECK: { |
| ; CHECK-EMPTY: |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ret; |
| ret void |
| } |