709 lines · plain
1; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s2; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-unknown-unknown | %ptxas-verify %}3;4; Check that parameters of a __device__ function with private or internal5; linkage called from a __global__ (kernel) function get increased alignment,6; and additional vectorization is performed on loads/stores with that7; parameters.8;9; Test IR is a minimized version of IR generated with the following command10; from the source code below:11; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu12;13; ----------------------------------------------------------------------------14; #include <stdint.h>15;16; struct St4x1 { uint32_t field[1]; };17; struct St4x2 { uint32_t field[2]; };18; struct St4x3 { uint32_t field[3]; };19; struct St4x4 { uint32_t field[4]; };20; struct St4x5 { uint32_t field[5]; };21; struct St4x6 { uint32_t field[6]; };22; struct St4x7 { uint32_t field[7]; };23; struct St4x8 { uint32_t field[8]; };24; struct St8x1 { uint64_t field[1]; };25; struct St8x2 { uint64_t field[2]; };26; struct St8x3 { uint64_t field[3]; };27; struct St8x4 { uint64_t field[4]; };28;29; #define DECLARE_CALLEE(StName) \30; static __device__ __attribute__((noinline)) \31; struct StName callee_##StName(struct StName in) { \32; struct StName ret; \33; const unsigned size = sizeof(ret.field) / sizeof(*ret.field); \34; for (unsigned i = 0; i != size; ++i) \35; ret.field[i] = in.field[i]; \36; return ret; \37; } \38 39; #define DECLARE_CALLER(StName) \40; __global__ \41; void caller_##StName(struct StName in, struct StName* ret) \42; { \43; *ret = callee_##StName(in); \44; } \45;46; #define DECLARE_CALL(StName) \47; DECLARE_CALLEE(StName) \48; DECLARE_CALLER(StName) \49;50; DECLARE_CALL(St4x1)51; DECLARE_CALL(St4x2)52; DECLARE_CALL(St4x3)53; DECLARE_CALL(St4x4)54; DECLARE_CALL(St4x5)55; DECLARE_CALL(St4x6)56; DECLARE_CALL(St4x7)57; DECLARE_CALL(St4x8)58; DECLARE_CALL(St8x1)59; DECLARE_CALL(St8x2)60; DECLARE_CALL(St8x3)61; DECLARE_CALL(St8x4)62; ----------------------------------------------------------------------------63 64%struct.St4x1 = type { [1 x i32] }65%struct.St4x2 = type { [2 x i32] }66%struct.St4x3 = type { [3 x i32] }67%struct.St4x4 = type { [4 x i32] }68%struct.St4x5 = type { [5 x i32] }69%struct.St4x6 = type { [6 x i32] }70%struct.St4x7 = type { [7 x i32] }71%struct.St4x8 = type { [8 x i32] }72%struct.St8x1 = type { [1 x i64] }73%struct.St8x2 = type { [2 x i64] }74%struct.St8x3 = type { [3 x i64] }75%struct.St8x4 = type { [4 x i64] }76 77; Section 1 - checking that:78; - function argument (including retval) vectorization is done with internal linkage;79; - caller and callee specify correct alignment for callee's params.80 81define dso_local void @caller_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) {82 ; CHECK-LABEL: .visible .func caller_St4x1(83 ; CHECK: .param .align 4 .b8 caller_St4x1_param_0[4],84 ; CHECK: .param .b64 caller_St4x1_param_185 ; CHECK: )86 ; CHECK: .param .b32 param0;87 ; CHECK: .param .align 16 .b8 retval0[4];88 ; CHECK: st.param.b32 [param0], {{%r[0-9]+}};89 ; CHECK: call.uni (retval0), callee_St4x1, (param0);90 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0];91 %1 = load i32, ptr %in, align 492 %call = tail call fastcc [1 x i32] @callee_St4x1(i32 %1) #293 %.fca.0.extract = extractvalue [1 x i32] %call, 094 store i32 %.fca.0.extract, ptr %ret, align 495 ret void96}97 98define internal fastcc [1 x i32] @callee_St4x1(i32 %in.0.val) {99 ; CHECK: .func (.param .align 16 .b8 func_retval0[4])100 ; CHECK-LABEL: callee_St4x1(101 ; CHECK-NEXT: .param .b32 callee_St4x1_param_0102 ; CHECK: ld.param.b32 [[R1:%r[0-9]+]], [callee_St4x1_param_0];103 ; CHECK: st.param.b32 [func_retval0], [[R1]];104 ; CHECK-NEXT: ret;105 %oldret = insertvalue [1 x i32] poison, i32 %in.0.val, 0106 ret [1 x i32] %oldret107}108 109define dso_local void @caller_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) {110 ; CHECK-LABEL: .visible .func caller_St4x2(111 ; CHECK: .param .align 4 .b8 caller_St4x2_param_0[8],112 ; CHECK: .param .b64 caller_St4x2_param_1113 ; CHECK: )114 ; CHECK: .param .align 16 .b8 param0[8];115 ; CHECK: .param .align 16 .b8 retval0[8];116 ; CHECK: st.param.v2.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}};117 ; CHECK: call.uni (retval0), callee_St4x2, (param0);118 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];119 %agg.tmp = alloca %struct.St4x2, align 8120 %1 = load i64, ptr %in, align 4121 store i64 %1, ptr %agg.tmp, align 8122 %call = tail call fastcc [2 x i32] @callee_St4x2(ptr noundef nonnull byval(%struct.St4x2) align 4 %agg.tmp) #2123 %.fca.0.extract = extractvalue [2 x i32] %call, 0124 %.fca.1.extract = extractvalue [2 x i32] %call, 1125 store i32 %.fca.0.extract, ptr %ret, align 4126 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4127 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4128 ret void129}130 131define internal fastcc [2 x i32] @callee_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in) {132 ; CHECK: .func (.param .align 16 .b8 func_retval0[8])133 ; CHECK-LABEL: callee_St4x2(134 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x2_param_0[8]135 ; CHECK: ld.param.v2.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x2_param_0];136 ; CHECK: st.param.v2.b32 [func_retval0], {[[R1]], [[R2]]};137 ; CHECK-NEXT: ret;138 %1 = load i32, ptr %in, align 4139 %arrayidx.1 = getelementptr inbounds [2 x i32], ptr %in, i64 0, i64 1140 %2 = load i32, ptr %arrayidx.1, align 4141 %3 = insertvalue [2 x i32] poison, i32 %1, 0142 %oldret = insertvalue [2 x i32] %3, i32 %2, 1143 ret [2 x i32] %oldret144}145 146define dso_local void @caller_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) {147 ; CHECK-LABEL: .visible .func caller_St4x3(148 ; CHECK: .param .align 4 .b8 caller_St4x3_param_0[12],149 ; CHECK: .param .b64 caller_St4x3_param_1150 ; CHECK: )151 ; CHECK: .param .align 16 .b8 param0[12];152 ; CHECK: .param .align 16 .b8 retval0[12];153 ; CHECK-DAG: st.param.v2.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}};154 ; CHECK-DAG: st.param.b32 [param0+8], {{%r[0-9]+}};155 ; CHECK: call.uni (retval0), callee_St4x3, (param0);156 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];157 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+8];158 %call = tail call fastcc [3 x i32] @callee_St4x3(ptr noundef nonnull byval(%struct.St4x3) align 4 %in) #2159 %.fca.0.extract = extractvalue [3 x i32] %call, 0160 %.fca.1.extract = extractvalue [3 x i32] %call, 1161 %.fca.2.extract = extractvalue [3 x i32] %call, 2162 store i32 %.fca.0.extract, ptr %ret, align 4163 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4164 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4165 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8166 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4167 ret void168}169 170define internal fastcc [3 x i32] @callee_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in) {171 ; CHECK: .func (.param .align 16 .b8 func_retval0[12])172 ; CHECK-LABEL: callee_St4x3(173 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x3_param_0[12]174 ; CHECK: ld.param.v2.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x3_param_0];175 ; CHECK: ld.param.b32 [[R3:%r[0-9]+]], [callee_St4x3_param_0+8];176 ; CHECK-DAG: st.param.v2.b32 [func_retval0], {[[R1]], [[R2]]};177 ; CHECK-DAG: st.param.b32 [func_retval0+8], [[R3]];178 ; CHECK-NEXT: ret;179 %1 = load i32, ptr %in, align 4180 %arrayidx.1 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 1181 %2 = load i32, ptr %arrayidx.1, align 4182 %arrayidx.2 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 2183 %3 = load i32, ptr %arrayidx.2, align 4184 %4 = insertvalue [3 x i32] poison, i32 %1, 0185 %5 = insertvalue [3 x i32] %4, i32 %2, 1186 %oldret = insertvalue [3 x i32] %5, i32 %3, 2187 ret [3 x i32] %oldret188}189 190define dso_local void @caller_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {191 ; CHECK-LABEL: .visible .func caller_St4x4(192 ; CHECK: .param .align 4 .b8 caller_St4x4_param_0[16],193 ; CHECK: .param .b64 caller_St4x4_param_1194 ; CHECK: )195 ; CHECK: .param .align 16 .b8 param0[16];196 ; CHECK: .param .align 16 .b8 retval0[16];197 ; CHECK: st.param.v4.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};198 ; CHECK: call.uni (retval0), callee_St4x4, (param0);199 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];200 %call = tail call fastcc [4 x i32] @callee_St4x4(ptr noundef nonnull byval(%struct.St4x4) align 4 %in) #2201 %.fca.0.extract = extractvalue [4 x i32] %call, 0202 %.fca.1.extract = extractvalue [4 x i32] %call, 1203 %.fca.2.extract = extractvalue [4 x i32] %call, 2204 %.fca.3.extract = extractvalue [4 x i32] %call, 3205 store i32 %.fca.0.extract, ptr %ret, align 4206 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4207 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4208 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8209 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4210 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12211 store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4212 ret void213}214 215define internal fastcc [4 x i32] @callee_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {216 ; CHECK: .func (.param .align 16 .b8 func_retval0[16])217 ; CHECK-LABEL: callee_St4x4(218 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_param_0[16]219 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x4_param_0];220 ; CHECK: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};221 ; CHECK-NEXT: ret;222 %1 = load i32, ptr %in, align 4223 %arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1224 %2 = load i32, ptr %arrayidx.1, align 4225 %arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2226 %3 = load i32, ptr %arrayidx.2, align 4227 %arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3228 %4 = load i32, ptr %arrayidx.3, align 4229 %5 = insertvalue [4 x i32] poison, i32 %1, 0230 %6 = insertvalue [4 x i32] %5, i32 %2, 1231 %7 = insertvalue [4 x i32] %6, i32 %3, 2232 %oldret = insertvalue [4 x i32] %7, i32 %4, 3233 ret [4 x i32] %oldret234}235 236define dso_local void @caller_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {237 ; CHECK-LABEL: .visible .func caller_St4x5(238 ; CHECK: .param .align 4 .b8 caller_St4x5_param_0[20],239 ; CHECK: .param .b64 caller_St4x5_param_1240 ; CHECK: )241 ; CHECK: .param .align 16 .b8 param0[20];242 ; CHECK: .param .align 16 .b8 retval0[20];243 ; CHECK-DAG: st.param.v4.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};244 ; CHECK-DAG: st.param.b32 [param0+16], {{%r[0-9]+}};245 ; CHECK: call.uni (retval0), callee_St4x5, (param0);246 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];247 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+16];248 %call = tail call fastcc [5 x i32] @callee_St4x5(ptr noundef nonnull byval(%struct.St4x5) align 4 %in) #2249 %.fca.0.extract = extractvalue [5 x i32] %call, 0250 %.fca.1.extract = extractvalue [5 x i32] %call, 1251 %.fca.2.extract = extractvalue [5 x i32] %call, 2252 %.fca.3.extract = extractvalue [5 x i32] %call, 3253 %.fca.4.extract = extractvalue [5 x i32] %call, 4254 store i32 %.fca.0.extract, ptr %ret, align 4255 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4256 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4257 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8258 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4259 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12260 store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4261 %ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16262 store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4263 ret void264}265 266define internal fastcc [5 x i32] @callee_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in) {267 ; CHECK: .func (.param .align 16 .b8 func_retval0[20])268 ; CHECK-LABEL: callee_St4x5(269 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x5_param_0[20]270 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x5_param_0];271 ; CHECK: ld.param.b32 [[R5:%r[0-9]+]], [callee_St4x5_param_0+16];272 ; CHECK-DAG: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};273 ; CHECK-DAG: st.param.b32 [func_retval0+16], [[R5]];274 ; CHECK-NEXT: ret;275 %1 = load i32, ptr %in, align 4276 %arrayidx.1 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 1277 %2 = load i32, ptr %arrayidx.1, align 4278 %arrayidx.2 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 2279 %3 = load i32, ptr %arrayidx.2, align 4280 %arrayidx.3 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 3281 %4 = load i32, ptr %arrayidx.3, align 4282 %arrayidx.4 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 4283 %5 = load i32, ptr %arrayidx.4, align 4284 %6 = insertvalue [5 x i32] poison, i32 %1, 0285 %7 = insertvalue [5 x i32] %6, i32 %2, 1286 %8 = insertvalue [5 x i32] %7, i32 %3, 2287 %9 = insertvalue [5 x i32] %8, i32 %4, 3288 %oldret = insertvalue [5 x i32] %9, i32 %5, 4289 ret [5 x i32] %oldret290}291 292define dso_local void @caller_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {293 ; CHECK-LABEL: .visible .func caller_St4x6(294 ; CHECK: .param .align 4 .b8 caller_St4x6_param_0[24],295 ; CHECK: .param .b64 caller_St4x6_param_1296 ; CHECK: )297 ; CHECK: .param .align 16 .b8 param0[24];298 ; CHECK: .param .align 16 .b8 retval0[24];299 ; CHECK-DAG: st.param.v4.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};300 ; CHECK-DAG: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};301 ; CHECK: call.uni (retval0), callee_St4x6, (param0);302 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];303 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];304 %call = tail call fastcc [6 x i32] @callee_St4x6(ptr noundef nonnull byval(%struct.St4x6) align 4 %in) #2305 %.fca.0.extract = extractvalue [6 x i32] %call, 0306 %.fca.1.extract = extractvalue [6 x i32] %call, 1307 %.fca.2.extract = extractvalue [6 x i32] %call, 2308 %.fca.3.extract = extractvalue [6 x i32] %call, 3309 %.fca.4.extract = extractvalue [6 x i32] %call, 4310 %.fca.5.extract = extractvalue [6 x i32] %call, 5311 store i32 %.fca.0.extract, ptr %ret, align 4312 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4313 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4314 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8315 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4316 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12317 store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4318 %ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16319 store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4320 %ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20321 store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4322 ret void323}324 325define internal fastcc [6 x i32] @callee_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in) {326 ; CHECK: .func (.param .align 16 .b8 func_retval0[24])327 ; CHECK-LABEL: callee_St4x6(328 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x6_param_0[24]329 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x6_param_0];330 ; CHECK: ld.param.v2.b32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x6_param_0+16];331 ; CHECK-DAG: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};332 ; CHECK-DAG: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]};333 ; CHECK-NEXT: ret;334 %1 = load i32, ptr %in, align 4335 %arrayidx.1 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 1336 %2 = load i32, ptr %arrayidx.1, align 4337 %arrayidx.2 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 2338 %3 = load i32, ptr %arrayidx.2, align 4339 %arrayidx.3 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 3340 %4 = load i32, ptr %arrayidx.3, align 4341 %arrayidx.4 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 4342 %5 = load i32, ptr %arrayidx.4, align 4343 %arrayidx.5 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 5344 %6 = load i32, ptr %arrayidx.5, align 4345 %7 = insertvalue [6 x i32] poison, i32 %1, 0346 %8 = insertvalue [6 x i32] %7, i32 %2, 1347 %9 = insertvalue [6 x i32] %8, i32 %3, 2348 %10 = insertvalue [6 x i32] %9, i32 %4, 3349 %11 = insertvalue [6 x i32] %10, i32 %5, 4350 %oldret = insertvalue [6 x i32] %11, i32 %6, 5351 ret [6 x i32] %oldret352}353 354define dso_local void @caller_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {355 ; CHECK-LABEL: .visible .func caller_St4x7(356 ; CHECK: .param .align 4 .b8 caller_St4x7_param_0[28],357 ; CHECK: .param .b64 caller_St4x7_param_1358 ; CHECK: )359 ; CHECK: .param .align 16 .b8 param0[28];360 ; CHECK: .param .align 16 .b8 retval0[28];361 ; CHECK-DAG: st.param.v4.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};362 ; CHECK-DAG: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};363 ; CHECK-DAG: st.param.b32 [param0+24], {{%r[0-9]+}};364 ; CHECK: call.uni (retval0), callee_St4x7, (param0);365 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];366 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];367 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+24];368 %call = tail call fastcc [7 x i32] @callee_St4x7(ptr noundef nonnull byval(%struct.St4x7) align 4 %in) #2369 %.fca.0.extract = extractvalue [7 x i32] %call, 0370 %.fca.1.extract = extractvalue [7 x i32] %call, 1371 %.fca.2.extract = extractvalue [7 x i32] %call, 2372 %.fca.3.extract = extractvalue [7 x i32] %call, 3373 %.fca.4.extract = extractvalue [7 x i32] %call, 4374 %.fca.5.extract = extractvalue [7 x i32] %call, 5375 %.fca.6.extract = extractvalue [7 x i32] %call, 6376 store i32 %.fca.0.extract, ptr %ret, align 4377 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4378 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4379 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8380 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4381 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12382 store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4383 %ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16384 store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4385 %ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20386 store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4387 %ref.tmp.sroa.9.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24388 store i32 %.fca.6.extract, ptr %ref.tmp.sroa.9.0..sroa_idx, align 4389 ret void390}391 392define internal fastcc [7 x i32] @callee_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in) {393 ; CHECK: .func (.param .align 16 .b8 func_retval0[28])394 ; CHECK-LABEL: callee_St4x7(395 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x7_param_0[28]396 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x7_param_0];397 ; CHECK: ld.param.v2.b32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x7_param_0+16];398 ; CHECK: ld.param.b32 [[R7:%r[0-9]+]], [callee_St4x7_param_0+24];399 ; CHECK-DAG: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};400 ; CHECK-DAG: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]};401 ; CHECK-DAG: st.param.b32 [func_retval0+24], [[R7]];402 ; CHECK-NEXT: ret;403 %1 = load i32, ptr %in, align 4404 %arrayidx.1 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 1405 %2 = load i32, ptr %arrayidx.1, align 4406 %arrayidx.2 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 2407 %3 = load i32, ptr %arrayidx.2, align 4408 %arrayidx.3 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 3409 %4 = load i32, ptr %arrayidx.3, align 4410 %arrayidx.4 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 4411 %5 = load i32, ptr %arrayidx.4, align 4412 %arrayidx.5 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 5413 %6 = load i32, ptr %arrayidx.5, align 4414 %arrayidx.6 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 6415 %7 = load i32, ptr %arrayidx.6, align 4416 %8 = insertvalue [7 x i32] poison, i32 %1, 0417 %9 = insertvalue [7 x i32] %8, i32 %2, 1418 %10 = insertvalue [7 x i32] %9, i32 %3, 2419 %11 = insertvalue [7 x i32] %10, i32 %4, 3420 %12 = insertvalue [7 x i32] %11, i32 %5, 4421 %13 = insertvalue [7 x i32] %12, i32 %6, 5422 %oldret = insertvalue [7 x i32] %13, i32 %7, 6423 ret [7 x i32] %oldret424}425 426define dso_local void @caller_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {427 ; CHECK-LABEL: .visible .func caller_St4x8(428 ; CHECK: .param .align 4 .b8 caller_St4x8_param_0[32],429 ; CHECK: .param .b64 caller_St4x8_param_1430 ; CHECK: )431 ; CHECK: .param .align 16 .b8 param0[32];432 ; CHECK: .param .align 16 .b8 retval0[32];433 ; CHECK-DAG: st.param.v4.b32 [param0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};434 ; CHECK-DAG: st.param.v4.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};435 ; CHECK: call.uni (retval0), callee_St4x8, (param0);436 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0];437 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16];438 %call = tail call fastcc [8 x i32] @callee_St4x8(ptr noundef nonnull byval(%struct.St4x8) align 4 %in) #2439 %.fca.0.extract = extractvalue [8 x i32] %call, 0440 %.fca.1.extract = extractvalue [8 x i32] %call, 1441 %.fca.2.extract = extractvalue [8 x i32] %call, 2442 %.fca.3.extract = extractvalue [8 x i32] %call, 3443 %.fca.4.extract = extractvalue [8 x i32] %call, 4444 %.fca.5.extract = extractvalue [8 x i32] %call, 5445 %.fca.6.extract = extractvalue [8 x i32] %call, 6446 %.fca.7.extract = extractvalue [8 x i32] %call, 7447 store i32 %.fca.0.extract, ptr %ret, align 4448 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 4449 store i32 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 4450 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8451 store i32 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 4452 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 12453 store i32 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 4454 %ref.tmp.sroa.7.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16455 store i32 %.fca.4.extract, ptr %ref.tmp.sroa.7.0..sroa_idx, align 4456 %ref.tmp.sroa.8.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 20457 store i32 %.fca.5.extract, ptr %ref.tmp.sroa.8.0..sroa_idx, align 4458 %ref.tmp.sroa.9.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24459 store i32 %.fca.6.extract, ptr %ref.tmp.sroa.9.0..sroa_idx, align 4460 %ref.tmp.sroa.10.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 28461 store i32 %.fca.7.extract, ptr %ref.tmp.sroa.10.0..sroa_idx, align 4462 ret void463}464 465define internal fastcc [8 x i32] @callee_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in) {466 ; CHECK: .func (.param .align 16 .b8 func_retval0[32])467 ; CHECK-LABEL: callee_St4x8(468 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x8_param_0[32]469 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x8_param_0];470 ; CHECK: ld.param.v4.b32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]], [[R7:%r[0-9]+]], [[R8:%r[0-9]+]]}, [callee_St4x8_param_0+16];471 ; CHECK-DAG: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};472 ; CHECK-DAG: st.param.v4.b32 [func_retval0+16], {[[R5]], [[R6]], [[R7]], [[R8]]};473 ; CHECK-NEXT: ret;474 %1 = load i32, ptr %in, align 4475 %arrayidx.1 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 1476 %2 = load i32, ptr %arrayidx.1, align 4477 %arrayidx.2 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 2478 %3 = load i32, ptr %arrayidx.2, align 4479 %arrayidx.3 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 3480 %4 = load i32, ptr %arrayidx.3, align 4481 %arrayidx.4 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 4482 %5 = load i32, ptr %arrayidx.4, align 4483 %arrayidx.5 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 5484 %6 = load i32, ptr %arrayidx.5, align 4485 %arrayidx.6 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 6486 %7 = load i32, ptr %arrayidx.6, align 4487 %arrayidx.7 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 7488 %8 = load i32, ptr %arrayidx.7, align 4489 %9 = insertvalue [8 x i32] poison, i32 %1, 0490 %10 = insertvalue [8 x i32] %9, i32 %2, 1491 %11 = insertvalue [8 x i32] %10, i32 %3, 2492 %12 = insertvalue [8 x i32] %11, i32 %4, 3493 %13 = insertvalue [8 x i32] %12, i32 %5, 4494 %14 = insertvalue [8 x i32] %13, i32 %6, 5495 %15 = insertvalue [8 x i32] %14, i32 %7, 6496 %oldret = insertvalue [8 x i32] %15, i32 %8, 7497 ret [8 x i32] %oldret498}499 500define dso_local void @caller_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) {501 ; CHECK-LABEL: .visible .func caller_St8x1(502 ; CHECK: .param .align 8 .b8 caller_St8x1_param_0[8],503 ; CHECK: .param .b64 caller_St8x1_param_1504 ; CHECK: )505 ; CHECK: .param .b64 param0;506 ; CHECK: .param .align 16 .b8 retval0[8];507 ; CHECK: st.param.b64 [param0], {{%rd[0-9]+}};508 ; CHECK: call.uni (retval0), callee_St8x1, (param0);509 ; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0];510 %1 = load i64, ptr %in, align 8511 %call = tail call fastcc [1 x i64] @callee_St8x1(i64 %1) #2512 %.fca.0.extract = extractvalue [1 x i64] %call, 0513 store i64 %.fca.0.extract, ptr %ret, align 8514 ret void515}516 517define internal fastcc [1 x i64] @callee_St8x1(i64 %in.0.val) {518 ; CHECK: .func (.param .align 16 .b8 func_retval0[8])519 ; CHECK-LABEL: callee_St8x1(520 ; CHECK-NEXT: .param .b64 callee_St8x1_param_0521 ; CHECK: ld.param.b64 [[RD1:%rd[0-9]+]], [callee_St8x1_param_0];522 ; CHECK: st.param.b64 [func_retval0], [[RD1]];523 ; CHECK-NEXT: ret;524 %oldret = insertvalue [1 x i64] poison, i64 %in.0.val, 0525 ret [1 x i64] %oldret526}527 528define dso_local void @caller_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) {529 ; CHECK-LABEL: .visible .func caller_St8x2(530 ; CHECK: .param .align 8 .b8 caller_St8x2_param_0[16],531 ; CHECK: .param .b64 caller_St8x2_param_1532 ; CHECK: )533 ; CHECK: .param .align 16 .b8 param0[16];534 ; CHECK: .param .align 16 .b8 retval0[16];535 ; CHECK: st.param.v2.b64 [param0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};536 ; CHECK: call.uni (retval0), callee_St8x2, (param0);537 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0];538 %call = tail call fastcc [2 x i64] @callee_St8x2(ptr noundef nonnull byval(%struct.St8x2) align 8 %in) #2539 %.fca.0.extract = extractvalue [2 x i64] %call, 0540 %.fca.1.extract = extractvalue [2 x i64] %call, 1541 store i64 %.fca.0.extract, ptr %ret, align 8542 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8543 store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8544 ret void545}546 547define internal fastcc [2 x i64] @callee_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in) {548 ; CHECK: .func (.param .align 16 .b8 func_retval0[16])549 ; CHECK-LABEL: callee_St8x2(550 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x2_param_0[16]551 ; CHECK: ld.param.v2.b64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x2_param_0];552 ; CHECK: st.param.v2.b64 [func_retval0], {[[RD1]], [[RD2]]};553 ; CHECK-NEXT: ret;554 %1 = load i64, ptr %in, align 8555 %arrayidx.1 = getelementptr inbounds [2 x i64], ptr %in, i64 0, i64 1556 %2 = load i64, ptr %arrayidx.1, align 8557 %3 = insertvalue [2 x i64] poison, i64 %1, 0558 %oldret = insertvalue [2 x i64] %3, i64 %2, 1559 ret [2 x i64] %oldret560}561 562define dso_local void @caller_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) {563 ; CHECK-LABEL: .visible .func caller_St8x3(564 ; CHECK: .param .align 8 .b8 caller_St8x3_param_0[24],565 ; CHECK: .param .b64 caller_St8x3_param_1566 ; CHECK: )567 ; CHECK: .param .align 16 .b8 param0[24];568 ; CHECK: .param .align 16 .b8 retval0[24];569 ; CHECK-DAG: st.param.v2.b64 [param0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};570 ; CHECK-DAG: st.param.b64 [param0+16], {{%rd[0-9]+}};571 ; CHECK: call.uni (retval0), callee_St8x3, (param0);572 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0];573 ; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0+16];574 %call = tail call fastcc [3 x i64] @callee_St8x3(ptr noundef nonnull byval(%struct.St8x3) align 8 %in) #2575 %.fca.0.extract = extractvalue [3 x i64] %call, 0576 %.fca.1.extract = extractvalue [3 x i64] %call, 1577 %.fca.2.extract = extractvalue [3 x i64] %call, 2578 store i64 %.fca.0.extract, ptr %ret, align 8579 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8580 store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8581 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16582 store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 8583 ret void584}585 586define internal fastcc [3 x i64] @callee_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in) {587 ; CHECK: .func (.param .align 16 .b8 func_retval0[24])588 ; CHECK-LABEL: callee_St8x3(589 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x3_param_0[24]590 ; CHECK: ld.param.v2.b64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x3_param_0];591 ; CHECK: ld.param.b64 [[RD3:%rd[0-9]+]], [callee_St8x3_param_0+16];592 ; CHECK-DAG: st.param.v2.b64 [func_retval0], {[[RD1]], [[RD2]]};593 ; CHECK-DAG: st.param.b64 [func_retval0+16], [[RD3]];594 ; CHECK-NEXT: ret;595 %1 = load i64, ptr %in, align 8596 %arrayidx.1 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 1597 %2 = load i64, ptr %arrayidx.1, align 8598 %arrayidx.2 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 2599 %3 = load i64, ptr %arrayidx.2, align 8600 %4 = insertvalue [3 x i64] poison, i64 %1, 0601 %5 = insertvalue [3 x i64] %4, i64 %2, 1602 %oldret = insertvalue [3 x i64] %5, i64 %3, 2603 ret [3 x i64] %oldret604}605 606define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {607 ; CHECK-LABEL: .visible .func caller_St8x4(608 ; CHECK: .param .align 8 .b8 caller_St8x4_param_0[32],609 ; CHECK: .param .b64 caller_St8x4_param_1610 ; CHECK: )611 ; CHECK: .param .align 16 .b8 param0[32];612 ; CHECK: .param .align 16 .b8 retval0[32];613 ; CHECK-DAG: st.param.v2.b64 [param0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};614 ; CHECK-DAG: st.param.v2.b64 [param0+16], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};615 ; CHECK: call.uni (retval0), callee_St8x4, (param0);616 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0];617 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+16];618 %call = tail call fastcc [4 x i64] @callee_St8x4(ptr noundef nonnull byval(%struct.St8x4) align 8 %in) #2619 %.fca.0.extract = extractvalue [4 x i64] %call, 0620 %.fca.1.extract = extractvalue [4 x i64] %call, 1621 %.fca.2.extract = extractvalue [4 x i64] %call, 2622 %.fca.3.extract = extractvalue [4 x i64] %call, 3623 store i64 %.fca.0.extract, ptr %ret, align 8624 %ref.tmp.sroa.4.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 8625 store i64 %.fca.1.extract, ptr %ref.tmp.sroa.4.0..sroa_idx, align 8626 %ref.tmp.sroa.5.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 16627 store i64 %.fca.2.extract, ptr %ref.tmp.sroa.5.0..sroa_idx, align 8628 %ref.tmp.sroa.6.0..sroa_idx = getelementptr inbounds i8, ptr %ret, i64 24629 store i64 %.fca.3.extract, ptr %ref.tmp.sroa.6.0..sroa_idx, align 8630 ret void631}632 633define internal fastcc [4 x i64] @callee_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in) {634 ; CHECK: .func (.param .align 16 .b8 func_retval0[32])635 ; CHECK-LABEL: callee_St8x4(636 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x4_param_0[32]637 ; CHECK: ld.param.v2.b64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x4_param_0];638 ; CHECK: ld.param.v2.b64 {[[RD3:%rd[0-9]+]], [[RD4:%rd[0-9]+]]}, [callee_St8x4_param_0+16];639 ; CHECK-DAG: st.param.v2.b64 [func_retval0], {[[RD1]], [[RD2]]};640 ; CHECK-DAG: st.param.v2.b64 [func_retval0+16], {[[RD3]], [[RD4]]};641 ; CHECK-NEXT: ret;642 %1 = load i64, ptr %in, align 8643 %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 1644 %2 = load i64, ptr %arrayidx.1, align 8645 %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 2646 %3 = load i64, ptr %arrayidx.2, align 8647 %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 3648 %4 = load i64, ptr %arrayidx.3, align 8649 %5 = insertvalue [4 x i64] poison, i64 %1, 0650 %6 = insertvalue [4 x i64] %5, i64 %2, 1651 %7 = insertvalue [4 x i64] %6, i64 %3, 2652 %oldret = insertvalue [4 x i64] %7, i64 %4, 3653 ret [4 x i64] %oldret654}655 656; Section 2 - checking that function argument (including retval) vectorization is done with private linkage.657 658define private fastcc [4 x i32] @callee_St4x4_private(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {659 ; CHECK: .func (.param .align 16 .b8 func_retval0[16])660 ; CHECK-LABEL: callee_St4x4_private(661 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_private_param_0[16]662 ; CHECK: ld.param.v4.b32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]], [[R3:%r[0-9]+]], [[R4:%r[0-9]+]]}, [callee_St4x4_private_param_0];663 ; CHECK: st.param.v4.b32 [func_retval0], {[[R1]], [[R2]], [[R3]], [[R4]]};664 ; CHECK-NEXT: ret;665 %1 = load i32, ptr %in, align 4666 %arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1667 %2 = load i32, ptr %arrayidx.1, align 4668 %arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2669 %3 = load i32, ptr %arrayidx.2, align 4670 %arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3671 %4 = load i32, ptr %arrayidx.3, align 4672 %5 = insertvalue [4 x i32] poison, i32 %1, 0673 %6 = insertvalue [4 x i32] %5, i32 %2, 1674 %7 = insertvalue [4 x i32] %6, i32 %3, 2675 %oldret = insertvalue [4 x i32] %7, i32 %4, 3676 ret [4 x i32] %oldret677}678 679; Section 3 - checking that function argument (including retval) vectorization680; is NOT done with linkage types other than internal and private.681 682define external fastcc [4 x i32] @callee_St4x4_external(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {683 ; CHECK: .func (.param .align 4 .b8 func_retval0[16])684 ; CHECK-LABEL: callee_St4x4_external(685 ; CHECK-NEXT: .param .align 4 .b8 callee_St4x4_external_param_0[16]686 ; CHECK: ld.param.b32 [[R1:%r[0-9]+]], [callee_St4x4_external_param_0];687 ; CHECK: ld.param.b32 [[R2:%r[0-9]+]], [callee_St4x4_external_param_0+4];688 ; CHECK: ld.param.b32 [[R3:%r[0-9]+]], [callee_St4x4_external_param_0+8];689 ; CHECK: ld.param.b32 [[R4:%r[0-9]+]], [callee_St4x4_external_param_0+12];690 ; CHECK-DAG: st.param.b32 [func_retval0], [[R1]];691 ; CHECK-DAG: st.param.b32 [func_retval0+4], [[R2]];692 ; CHECK-DAG: st.param.b32 [func_retval0+8], [[R3]];693 ; CHECK-DAG: st.param.b32 [func_retval0+12], [[R4]];694 ; CHECK-NEXT: ret;695 %1 = load i32, ptr %in, align 4696 %arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1697 %2 = load i32, ptr %arrayidx.1, align 4698 %arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2699 %3 = load i32, ptr %arrayidx.2, align 4700 %arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3701 %4 = load i32, ptr %arrayidx.3, align 4702 %5 = insertvalue [4 x i32] poison, i32 %1, 0703 %6 = insertvalue [4 x i32] %5, i32 %2, 1704 %7 = insertvalue [4 x i32] %6, i32 %3, 2705 %oldret = insertvalue [4 x i32] %7, i32 %4, 3706 ret [4 x i32] %oldret707}708 709