; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx-unknown-unknown | %ptxas-verify %}
;
; Check that parameters of a __device__ function with private or internal
; linkage called from a __global__ (kernel) function get increased alignment,
; and additional vectorization is performed on loads/stores with that
; parameters.
;
; Test IR is a minimized version of IR generated with the following command
; from the source code below:
; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu
;
; ----------------------------------------------------------------------------
; #include <stdint.h>
;
; struct St4x1 { uint32_t field[1]; };
; struct St4x2 { uint32_t field[2]; };
; struct St4x3 { uint32_t field[3]; };
; struct St4x4 { uint32_t field[4]; };
; struct St4x5 { uint32_t field[5]; };
; struct St4x6 { uint32_t field[6]; };
; struct St4x7 { uint32_t field[7]; };
; struct St4x8 { uint32_t field[8]; };
; struct St8x1 { uint64_t field[1]; };
; struct St8x2 { uint64_t field[2]; };
; struct St8x3 { uint64_t field[3]; };
; struct St8x4 { uint64_t field[4]; };
;
; #define DECLARE_CALLEE(StName) \
; static __device__ __attribute__((noinline)) \
; struct StName callee_##StName(struct StName in) { \
; struct StName ret; \
; const unsigned size = sizeof(ret.field) / sizeof(*ret.field); \
; for (unsigned i = 0; i != size; ++i) \
; ret.field[i] = in.field[i]; \
; return ret; \
; } \
; #define DECLARE_CALLER(StName) \
; __global__ \
; void caller_##StName(struct StName in, struct StName* ret) \
; { \
; *ret = callee_##StName(in); \
; } \
;
; #define DECLARE_CALL(StName) \
; DECLARE_CALLEE(StName) \
; DECLARE_CALLER(StName) \
;
; DECLARE_CALL(St4x1)
; DECLARE_CALL(St4x2)
; DECLARE_CALL(St4x3)
; DECLARE_CALL(St4x4)
; DECLARE_CALL(St4x5)
; DECLARE_CALL(St4x6)
; DECLARE_CALL(St4x7)
; DECLARE_CALL(St4x8)
; DECLARE_CALL(St8x1)
; DECLARE_CALL(St8x2)
; DECLARE_CALL(St8x3)
; DECLARE_CALL(St8x4)
; ----------------------------------------------------------------------------
%struct.St4x1 = type { [1 x i32] }
%struct.St4x2 = type { [2 x i32] }
%struct.St4x3 = type { [3 x i32] }
%struct.St4x4 = type { [4 x i32] }
%struct.St4x5 = type { [5 x i32] }
%struct.St4x6 = type { [6 x i32] }
%struct.St4x7 = type { [7 x i32] }
%struct.St4x8 = type { [8 x i32] }
%struct.St8x1 = type { [1 x i64] }
%struct.St8x2 = type { [2 x i64] }
%struct.St8x3 = type { [3 x i64] }
%struct.St8x4 = type { [4 x i64] }
; Section 1 - checking that:
; - function argument (including retval) vectorization is done with internal linkage;
; - caller and callee specify correct alignment for callee's params.
define dso_local void @caller_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x1(
; CHECK: .param .align 4 .b8 caller_St4x1_param_0[4],
; CHECK: .param .b32 caller_St4x1_param_1
; CHECK: )
; CHECK: .param .b32 param0;
; CHECK: st.param.b32 [param0+0], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[4];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x1,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+0];
%1 = load i32, ptr %in, align 4
%call = tail call fastcc [1 x i32] @callee_St4x1(i32 %1) #2
%.fca.0.extract = extractvalue [1 x i32] %call, 0
store i32 %.fca.0.extract, ptr %ret, align 4
ret void
}
define internal fastcc [1 x i32] @callee_St4x1(i32 %in.0.val) {
; CHECK: .func (.param .align 16 .b8 func_retval0[4])
; CHECK-LABEL: callee_St4x1(
; CHECK-NEXT: .param .b32 callee_St4x1_param_0
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [callee_St4x1_param_0];
; CHECK: st.param.b32 [func_retval0+0], [[R1]];
; CHECK-NEXT: ret;
%oldret = insertvalue [1 x i32] poison, i32 %in.0.val, 0
ret [1 x i32] %oldret
}
define dso_local void @caller_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x2(
; CHECK: .param .align 4 .b8 caller_St4x2_param_0[8],
; CHECK: .param .b32 caller_St4x2_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[8];
; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[8];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x2,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
%agg.tmp = alloca %struct.St4x2, align 8
%1 = load i64, ptr %in, align 4
store i64 %1, ptr %agg.tmp, align 8
%call = tail call fastcc [2 x i32] @callee_St4x2(ptr noundef nonnull byval(%struct.St4x2) align 4 %agg.tmp) #2
%.fca.0.extract = extractvalue [2 x i32] %call, 0
%.fca.1.extract = extractvalue [2 x i32] %call, 1
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
ret void
}
define internal fastcc [2 x i32] @callee_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[8])
; CHECK-LABEL: callee_St4x2(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x2_param_0[8]
; CHECK: ld.param.v2.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x2_param_0];
; CHECK: st.param.v2.b32 [func_retval0+0], {[[R1]], [[R2]]};
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [2 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%3 = insertvalue [2 x i32] poison, i32 %1, 0
%oldret = insertvalue [2 x i32] %3, i32 %2, 1
ret [2 x i32] %oldret
}
define dso_local void @caller_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x3(
; CHECK: .param .align 4 .b8 caller_St4x3_param_0[12],
; CHECK: .param .b32 caller_St4x3_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[12];
; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+8], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[12];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x3,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+8];
%call = tail call fastcc [3 x i32] @callee_St4x3(ptr noundef nonnull byval(%struct.St4x3) align 4 %in) #2
%.fca.0.extract = extractvalue [3 x i32] %call, 0
%.fca.1.extract = extractvalue [3 x i32] %call, 1
%.fca.2.extract = extractvalue [3 x i32] %call, 2
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
ret void
}
define internal fastcc [3 x i32] @callee_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[12])
; CHECK-LABEL: callee_St4x3(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x3_param_0[12]
; CHECK: ld.param.v2.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x3_param_0];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [callee_St4x3_param_0+8];
; CHECK: st.param.v2.b32 [func_retval0+0], {[[R1]], [[R2]]};
; CHECK: st.param.b32 [func_retval0+8], [[R3]];
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%4 = insertvalue [3 x i32] poison, i32 %1, 0
%5 = insertvalue [3 x i32] %4, i32 %2, 1
%oldret = insertvalue [3 x i32] %5, i32 %3, 2
ret [3 x i32] %oldret
}
define dso_local void @caller_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x4(
; CHECK: .param .align 4 .b8 caller_St4x4_param_0[16],
; CHECK: .param .b32 caller_St4x4_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[16];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[16];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x4,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
%call = tail call fastcc [4 x i32] @callee_St4x4(ptr noundef nonnull byval(%struct.St4x4) align 4 %in) #2
%.fca.0.extract = extractvalue [4 x i32] %call, 0
%.fca.1.extract = extractvalue [4 x i32] %call, 1
%.fca.2.extract = extractvalue [4 x i32] %call, 2
%.fca.3.extract = extractvalue [4 x i32] %call, 3
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
%ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12
store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4
ret void
}
define internal fastcc [4 x i32] @callee_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[16])
; CHECK-LABEL: callee_St4x4(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_param_0[16]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x4_param_0];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%5 = insertvalue [4 x i32] poison, i32 %1, 0
%6 = insertvalue [4 x i32] %5, i32 %2, 1
%7 = insertvalue [4 x i32] %6, i32 %3, 2
%oldret = insertvalue [4 x i32] %7, i32 %4, 3
ret [4 x i32] %oldret
}
define dso_local void @caller_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x5(
; CHECK: .param .align 4 .b8 caller_St4x5_param_0[20],
; CHECK: .param .b32 caller_St4x5_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[20];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+16], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[20];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x5,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+16];
%call = tail call fastcc [5 x i32] @callee_St4x5(ptr noundef nonnull byval(%struct.St4x5) align 4 %in) #2
%.fca.0.extract = extractvalue [5 x i32] %call, 0
%.fca.1.extract = extractvalue [5 x i32] %call, 1
%.fca.2.extract = extractvalue [5 x i32] %call, 2
%.fca.3.extract = extractvalue [5 x i32] %call, 3
%.fca.4.extract = extractvalue [5 x i32] %call, 4
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
%ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12
store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4
%ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4
ret void
}
define internal fastcc [5 x i32] @callee_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[20])
; CHECK-LABEL: callee_St4x5(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x5_param_0[20]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x5_param_0];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [callee_St4x5_param_0+16];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK: st.param.b32 [func_retval0+16], [[R5]];
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%arrayidx.4 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 4
%5 = load i32, ptr %arrayidx.4, align 4
%6 = insertvalue [5 x i32] poison, i32 %1, 0
%7 = insertvalue [5 x i32] %6, i32 %2, 1
%8 = insertvalue [5 x i32] %7, i32 %3, 2
%9 = insertvalue [5 x i32] %8, i32 %4, 3
%oldret = insertvalue [5 x i32] %9, i32 %5, 4
ret [5 x i32] %oldret
}
define dso_local void @caller_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x6(
; CHECK: .param .align 4 .b8 caller_St4x6_param_0[24],
; CHECK: .param .b32 caller_St4x6_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[24];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[24];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x6,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];
%call = tail call fastcc [6 x i32] @callee_St4x6(ptr noundef nonnull byval(%struct.St4x6) align 4 %in) #2
%.fca.0.extract = extractvalue [6 x i32] %call, 0
%.fca.1.extract = extractvalue [6 x i32] %call, 1
%.fca.2.extract = extractvalue [6 x i32] %call, 2
%.fca.3.extract = extractvalue [6 x i32] %call, 3
%.fca.4.extract = extractvalue [6 x i32] %call, 4
%.fca.5.extract = extractvalue [6 x i32] %call, 5
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
%ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12
store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4
%ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4
%ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20
store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4
ret void
}
define internal fastcc [6 x i32] @callee_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[24])
; CHECK-LABEL: callee_St4x6(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x6_param_0[24]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x6_param_0];
; CHECK: ld.param.v2.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x6_param_0+16];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]};
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%arrayidx.4 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 4
%5 = load i32, ptr %arrayidx.4, align 4
%arrayidx.5 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 5
%6 = load i32, ptr %arrayidx.5, align 4
%7 = insertvalue [6 x i32] poison, i32 %1, 0
%8 = insertvalue [6 x i32] %7, i32 %2, 1
%9 = insertvalue [6 x i32] %8, i32 %3, 2
%10 = insertvalue [6 x i32] %9, i32 %4, 3
%11 = insertvalue [6 x i32] %10, i32 %5, 4
%oldret = insertvalue [6 x i32] %11, i32 %6, 5
ret [6 x i32] %oldret
}
define dso_local void @caller_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x7(
; CHECK: .param .align 4 .b8 caller_St4x7_param_0[28],
; CHECK: .param .b32 caller_St4x7_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[28];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+24], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[28];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x7,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];
; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+24];
%call = tail call fastcc [7 x i32] @callee_St4x7(ptr noundef nonnull byval(%struct.St4x7) align 4 %in) #2
%.fca.0.extract = extractvalue [7 x i32] %call, 0
%.fca.1.extract = extractvalue [7 x i32] %call, 1
%.fca.2.extract = extractvalue [7 x i32] %call, 2
%.fca.3.extract = extractvalue [7 x i32] %call, 3
%.fca.4.extract = extractvalue [7 x i32] %call, 4
%.fca.5.extract = extractvalue [7 x i32] %call, 5
%.fca.6.extract = extractvalue [7 x i32] %call, 6
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
%ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12
store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4
%ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4
%ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20
store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4
%ref.tmp.sroa.9.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24
store i32 %.fca.6.extract, ptr %ref.tmp.sroa.9.0..sroa_idx, align 4
ret void
}
define internal fastcc [7 x i32] @callee_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[28])
; CHECK-LABEL: callee_St4x7(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x7_param_0[28]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x7_param_0];
; CHECK: ld.param.v2.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x7_param_0+16];
; CHECK: ld.param.u32 [[R7:%r[0-9]+]], [callee_St4x7_param_0+24];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]};
; CHECK: st.param.b32 [func_retval0+24], [[R7]];
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%arrayidx.4 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 4
%5 = load i32, ptr %arrayidx.4, align 4
%arrayidx.5 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 5
%6 = load i32, ptr %arrayidx.5, align 4
%arrayidx.6 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 6
%7 = load i32, ptr %arrayidx.6, align 4
%8 = insertvalue [7 x i32] poison, i32 %1, 0
%9 = insertvalue [7 x i32] %8, i32 %2, 1
%10 = insertvalue [7 x i32] %9, i32 %3, 2
%11 = insertvalue [7 x i32] %10, i32 %4, 3
%12 = insertvalue [7 x i32] %11, i32 %5, 4
%13 = insertvalue [7 x i32] %12, i32 %6, 5
%oldret = insertvalue [7 x i32] %13, i32 %7, 6
ret [7 x i32] %oldret
}
define dso_local void @caller_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x8(
; CHECK: .param .align 4 .b8 caller_St4x8_param_0[32],
; CHECK: .param .b32 caller_St4x8_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[32];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v4.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[32];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x8,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0];
; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];
%call = tail call fastcc [8 x i32] @callee_St4x8(ptr noundef nonnull byval(%struct.St4x8) align 4 %in) #2
%.fca.0.extract = extractvalue [8 x i32] %call, 0
%.fca.1.extract = extractvalue [8 x i32] %call, 1
%.fca.2.extract = extractvalue [8 x i32] %call, 2
%.fca.3.extract = extractvalue [8 x i32] %call, 3
%.fca.4.extract = extractvalue [8 x i32] %call, 4
%.fca.5.extract = extractvalue [8 x i32] %call, 5
%.fca.6.extract = extractvalue [8 x i32] %call, 6
%.fca.7.extract = extractvalue [8 x i32] %call, 7
store i32 %.fca.0.extract, ptr %ret, align 4
%ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4
store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4
%ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8
store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4
%ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12
store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4
%ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16
store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4
%ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20
store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4
%ref.tmp.sroa.9.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24
store i32 %.fca.6.extract, ptr %ref.tmp.sroa.9.0..sroa_idx, align 4
%ref.tmp.sroa.10.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 28
store i32 %.fca.7.extract, ptr %ref.tmp.sroa.10.0..sroa_idx, align 4
ret void
}
define internal fastcc [8 x i32] @callee_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[32])
; CHECK-LABEL: callee_St4x8(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x8_param_0[32]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x8_param_0];
; CHECK: ld.param.v4.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]], [[R7:%r[0-9]+]], [[R8:%r[0-9]+]]}, [callee_St4x8_param_0+16];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK: st.param.v4.b32 [func_retval0+16], {[[R5]], [[R6]], [[R7]], [[R8]]};
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%arrayidx.4 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 4
%5 = load i32, ptr %arrayidx.4, align 4
%arrayidx.5 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 5
%6 = load i32, ptr %arrayidx.5, align 4
%arrayidx.6 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 6
%7 = load i32, ptr %arrayidx.6, align 4
%arrayidx.7 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 7
%8 = load i32, ptr %arrayidx.7, align 4
%9 = insertvalue [8 x i32] poison, i32 %1, 0
%10 = insertvalue [8 x i32] %9, i32 %2, 1
%11 = insertvalue [8 x i32] %10, i32 %3, 2
%12 = insertvalue [8 x i32] %11, i32 %4, 3
%13 = insertvalue [8 x i32] %12, i32 %5, 4
%14 = insertvalue [8 x i32] %13, i32 %6, 5
%15 = insertvalue [8 x i32] %14, i32 %7, 6
%oldret = insertvalue [8 x i32] %15, i32 %8, 7
ret [8 x i32] %oldret
}
define dso_local void @caller_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x1(
; CHECK: .param .align 8 .b8 caller_St8x1_param_0[8],
; CHECK: .param .b32 caller_St8x1_param_1
; CHECK: )
; CHECK: .param .b64 param0;
; CHECK: st.param.b64 [param0+0], {{%rd[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[8];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x1,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0+0];
%1 = load i64, ptr %in, align 8
%call = tail call fastcc [1 x i64] @callee_St8x1(i64 %1) #2
%.fca.0.extract = extractvalue [1 x i64] %call, 0
store i64 %.fca.0.extract, ptr %ret, align 8
ret void
}
define internal fastcc [1 x i64] @callee_St8x1(i64 %in.0.val) {
; CHECK: .func (.param .align 16 .b8 func_retval0[8])
; CHECK-LABEL: callee_St8x1(
; CHECK-NEXT: .param .b64 callee_St8x1_param_0
; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [callee_St8x1_param_0];
; CHECK: st.param.b64 [func_retval0+0], [[RD1]];
; CHECK-NEXT: ret;
%oldret = insertvalue [1 x i64] poison, i64 %in.0.val, 0
ret [1 x i64] %oldret
}
define dso_local void @caller_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x2(
; CHECK: .param .align 8 .b8 caller_St8x2_param_0[16],
; CHECK: .param .b32 caller_St8x2_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[16];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[16];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x2,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0];
%call = tail call fastcc [2 x i64] @callee_St8x2(ptr noundef nonnull byval(%struct.St8x2) align 8 %in) #2
%.fca.0.extract = extractvalue [2 x i64] %call, 0
%.fca.1.extract = extractvalue [2 x i64] %call, 1
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
ret void
}
define internal fastcc [2 x i64] @callee_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[16])
; CHECK-LABEL: callee_St8x2(
; CHECK-NEXT: .param .align 16 .b8 callee_St8x2_param_0[16]
; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x2_param_0];
; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]};
; CHECK-NEXT: ret;
%1 = load i64, ptr %in, align 8
%arrayidx.1 = getelementptr inbounds [2 x i64], ptr %in, i64 0, i64 1
%2 = load i64, ptr %arrayidx.1, align 8
%3 = insertvalue [2 x i64] poison, i64 %1, 0
%oldret = insertvalue [2 x i64] %3, i64 %2, 1
ret [2 x i64] %oldret
}
define dso_local void @caller_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x3(
; CHECK: .param .align 8 .b8 caller_St8x3_param_0[24],
; CHECK: .param .b32 caller_St8x3_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[24];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: st.param.b64 [param0+16], {{%rd[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[24];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x3,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0];
; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0+16];
%call = tail call fastcc [3 x i64] @callee_St8x3(ptr noundef nonnull byval(%struct.St8x3) align 8 %in) #2
%.fca.0.extract = extractvalue [3 x i64] %call, 0
%.fca.1.extract = extractvalue [3 x i64] %call, 1
%.fca.2.extract = extractvalue [3 x i64] %call, 2
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
ret void
}
define internal fastcc [3 x i64] @callee_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[24])
; CHECK-LABEL: callee_St8x3(
; CHECK-NEXT: .param .align 16 .b8 callee_St8x3_param_0[24]
; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x3_param_0];
; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [callee_St8x3_param_0+16];
; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]};
; CHECK: st.param.b64 [func_retval0+16], [[RD3]];
; CHECK-NEXT: ret;
%1 = load i64, ptr %in, align 8
%arrayidx.1 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 1
%2 = load i64, ptr %arrayidx.1, align 8
%arrayidx.2 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 2
%3 = load i64, ptr %arrayidx.2, align 8
%4 = insertvalue [3 x i64] poison, i64 %1, 0
%5 = insertvalue [3 x i64] %4, i64 %2, 1
%oldret = insertvalue [3 x i64] %5, i64 %3, 2
ret [3 x i64] %oldret
}
define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x4(
; CHECK: .param .align 8 .b8 caller_St8x4_param_0[32],
; CHECK: .param .b32 caller_St8x4_param_1
; CHECK: )
; CHECK: .param .align 16 .b8 param0[32];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: st.param.v2.b64 [param0+16], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[32];
; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x4,
; CHECK-NEXT: (
; CHECK-NEXT: param0
; CHECK-NEXT: );
; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0];
; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+16];
%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: .func (.param .align 16 .b8 func_retval0[32])
; CHECK-LABEL: callee_St8x4(
; CHECK-NEXT: .param .align 16 .b8 callee_St8x4_param_0[32]
; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x4_param_0];
; CHECK: ld.param.v2.u64 {[[RD3:%rd[0-9]+]], [[RD4:%rd[0-9]+]]}, [callee_St8x4_param_0+16];
; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]};
; CHECK: 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
}
; Section 2 - checking that function argument (including retval) vectorization is done with private linkage.
define private fastcc [4 x i32] @callee_St4x4_private(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {
; CHECK: .func (.param .align 16 .b8 func_retval0[16])
; CHECK-LABEL: callee_St4x4_private(
; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_private_param_0[16]
; CHECK: ld.param.v4.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x4_private_param_0];
; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]};
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%5 = insertvalue [4 x i32] poison, i32 %1, 0
%6 = insertvalue [4 x i32] %5, i32 %2, 1
%7 = insertvalue [4 x i32] %6, i32 %3, 2
%oldret = insertvalue [4 x i32] %7, i32 %4, 3
ret [4 x i32] %oldret
}
; Section 3 - checking that function argument (including retval) vectorization
; is NOT done with linkage types other than internal and private.
define external fastcc [4 x i32] @callee_St4x4_external(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {
; CHECK: .func (.param .align 4 .b8 func_retval0[16])
; CHECK-LABEL: callee_St4x4_external(
; CHECK-NEXT: .param .align 4 .b8 callee_St4x4_external_param_0[16]
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [callee_St4x4_external_param_0];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [callee_St4x4_external_param_0+4];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [callee_St4x4_external_param_0+8];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [callee_St4x4_external_param_0+12];
; CHECK: st.param.b32 [func_retval0+0], [[R1]];
; CHECK: st.param.b32 [func_retval0+4], [[R2]];
; CHECK: st.param.b32 [func_retval0+8], [[R3]];
; CHECK: st.param.b32 [func_retval0+12], [[R4]];
; CHECK-NEXT: ret;
%1 = load i32, ptr %in, align 4
%arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4
%arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4
%arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3
%4 = load i32, ptr %arrayidx.3, align 4
%5 = insertvalue [4 x i32] poison, i32 %1, 0
%6 = insertvalue [4 x i32] %5, i32 %2, 1
%7 = insertvalue [4 x i32] %6, i32 %3, 2
%oldret = insertvalue [4 x i32] %7, i32 %4, 3
ret [4 x i32] %oldret
}