82 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py2; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s3 4%struct.St8x4 = type { [4 x i64] }5 6define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {7; CHECK-LABEL: caller_St8x4(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<2>;10; CHECK-NEXT: .reg .b64 %rd<9>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: { // callseq 0, 014; CHECK-NEXT: .param .align 16 .b8 param0[32];15; CHECK-NEXT: .param .align 16 .b8 retval0[32];16; CHECK-NEXT: ld.param.b64 %rd1, [caller_St8x4_param_0+24];17; CHECK-NEXT: ld.param.b64 %rd2, [caller_St8x4_param_0+16];18; CHECK-NEXT: st.param.v2.b64 [param0+16], {%rd2, %rd1};19; CHECK-NEXT: ld.param.b64 %rd3, [caller_St8x4_param_0+8];20; CHECK-NEXT: ld.param.b64 %rd4, [caller_St8x4_param_0];21; CHECK-NEXT: st.param.v2.b64 [param0], {%rd4, %rd3};22; CHECK-NEXT: call.uni (retval0), callee_St8x4, (param0);23; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [retval0];24; CHECK-NEXT: ld.param.v2.b64 {%rd7, %rd8}, [retval0+16];25; CHECK-NEXT: } // callseq 026; CHECK-NEXT: ld.param.b32 %r1, [caller_St8x4_param_1];27; CHECK-NEXT: st.b64 [%r1], %rd5;28; CHECK-NEXT: st.b64 [%r1+8], %rd6;29; CHECK-NEXT: st.b64 [%r1+16], %rd7;30; CHECK-NEXT: st.b64 [%r1+24], %rd8;31; CHECK-NEXT: ret;32 %call = tail call fastcc [4 x i64] @callee_St8x4(ptr noundef nonnull byval(%struct.St8x4) align 8 %in) #233 %.fca.0.extract = extractvalue [4 x i64] %call, 034 %.fca.1.extract = extractvalue [4 x i64] %call, 135 %.fca.2.extract = extractvalue [4 x i64] %call, 236 %.fca.3.extract = extractvalue [4 x i64] %call, 337 store i64 %.fca.0.extract, ptr %ret, align 838 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 839 store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 840 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 1641 store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 842 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 2443 store i64 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 844 ret void45}46 47define internal fastcc [4 x i64] @callee_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in) {48; CHECK-LABEL: callee_St8x4(49; CHECK: // @callee_St8x450; CHECK-NEXT: {51; CHECK-NEXT: .reg .b64 %rd<5>;52; CHECK-EMPTY:53; CHECK-NEXT: // %bb.0:54; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [callee_St8x4_param_0];55; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [callee_St8x4_param_0+16];56; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd3, %rd4};57; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};58; CHECK-NEXT: ret;59 %1 = load i64, ptr %in, align 860 %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 161 %2 = load i64, ptr %arrayidx.1, align 862 %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 263 %3 = load i64, ptr %arrayidx.2, align 864 %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 365 %4 = load i64, ptr %arrayidx.3, align 866 %5 = insertvalue [4 x i64] poison, i64 %1, 067 %6 = insertvalue [4 x i64] %5, i64 %2, 168 %7 = insertvalue [4 x i64] %6, i64 %3, 269 %oldret = insertvalue [4 x i64] %7, i64 %4, 370 ret [4 x i64] %oldret71}72 73define void @call_void() {74; CHECK-LABEL: call_void(75; CHECK: {76; CHECK-EMPTY:77; CHECK-EMPTY:78; CHECK-NEXT: // %bb.0:79; CHECK-NEXT: ret;80 ret void81}82