1; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s 2; 3; Check that parameters of a __device__ function with private or internal 4; linkage called from a __global__ (kernel) function get increased alignment, 5; and additional vectorization is performed on loads/stores with that 6; parameters. 7; 8; Test IR is a minimized version of IR generated with the following command 9; from the source code below: 10; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu 11; 12; ---------------------------------------------------------------------------- 13; #include <stdint.h> 14; 15; struct St4x1 { uint32_t field[1]; }; 16; struct St4x2 { uint32_t field[2]; }; 17; struct St4x3 { uint32_t field[3]; }; 18; struct St4x4 { uint32_t field[4]; }; 19; struct St4x5 { uint32_t field[5]; }; 20; struct St4x6 { uint32_t field[6]; }; 21; struct St4x7 { uint32_t field[7]; }; 22; struct St4x8 { uint32_t field[8]; }; 23; struct St8x1 { uint64_t field[1]; }; 24; struct St8x2 { uint64_t field[2]; }; 25; struct St8x3 { uint64_t field[3]; }; 26; struct St8x4 { uint64_t field[4]; }; 27; 28; #define DECLARE_CALLEE(StName) \ 29; static __device__ __attribute__((noinline)) \ 30; struct StName callee_##StName(struct StName in) { \ 31; struct StName ret; \ 32; const unsigned size = sizeof(ret.field) / sizeof(*ret.field); \ 33; for (unsigned i = 0; i != size; ++i) \ 34; ret.field[i] = in.field[i]; \ 35; return ret; \ 36; } \ 37 38; #define DECLARE_CALLER(StName) \ 39; __global__ \ 40; void caller_##StName(struct StName in, struct StName* ret) \ 41; { \ 42; *ret = callee_##StName(in); \ 43; } \ 44; 45; #define DECLARE_CALL(StName) \ 46; DECLARE_CALLEE(StName) \ 47; DECLARE_CALLER(StName) \ 48; 49; DECLARE_CALL(St4x1) 50; DECLARE_CALL(St4x2) 51; DECLARE_CALL(St4x3) 52; DECLARE_CALL(St4x4) 53; DECLARE_CALL(St4x5) 54; DECLARE_CALL(St4x6) 55; DECLARE_CALL(St4x7) 56; DECLARE_CALL(St4x8) 57; DECLARE_CALL(St8x1) 58; DECLARE_CALL(St8x2) 59; DECLARE_CALL(St8x3) 60; DECLARE_CALL(St8x4) 61; ---------------------------------------------------------------------------- 62 63%struct.St4x1 = type { [1 x i32] } 64%struct.St4x2 = type { [2 x i32] } 65%struct.St4x3 = type { [3 x i32] } 66%struct.St4x4 = type { [4 x i32] } 67%struct.St4x5 = type { [5 x i32] } 68%struct.St4x6 = type { [6 x i32] } 69%struct.St4x7 = type { [7 x i32] } 70%struct.St4x8 = type { [8 x i32] } 71%struct.St8x1 = type { [1 x i64] } 72%struct.St8x2 = type { [2 x i64] } 73%struct.St8x3 = type { [3 x i64] } 74%struct.St8x4 = type { [4 x i64] } 75 76; Section 1 - checking that: 77; - function argument (including retval) vectorization is done with internal linkage; 78; - caller and callee specify correct alignment for callee's params. 79 80define dso_local void @caller_St4x1(%struct.St4x1* nocapture noundef readonly byval(%struct.St4x1) align 4 %in, %struct.St4x1* nocapture noundef writeonly %ret) { 81 ; CHECK-LABEL: .visible .func caller_St4x1( 82 ; CHECK: .param .align 4 .b8 caller_St4x1_param_0[4], 83 ; CHECK: .param .b32 caller_St4x1_param_1 84 ; CHECK: ) 85 ; CHECK: .param .b32 param0; 86 ; CHECK: st.param.b32 [param0+0], {{%r[0-9]+}}; 87 ; CHECK: .param .align 16 .b8 retval0[4]; 88 ; CHECK: call.uni (retval0), 89 ; CHECK-NEXT: callee_St4x1, 90 ; CHECK-NEXT: ( 91 ; CHECK-NEXT: param0 92 ; CHECK-NEXT: ); 93 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+0]; 94 %1 = getelementptr inbounds %struct.St4x1, %struct.St4x1* %in, i64 0, i32 0, i64 0 95 %2 = load i32, i32* %1, align 4 96 %call = tail call fastcc [1 x i32] @callee_St4x1(i32 %2) 97 %.fca.0.extract = extractvalue [1 x i32] %call, 0 98 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x1, %struct.St4x1* %ret, i64 0, i32 0, i64 0 99 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 100 ret void 101} 102 103define internal fastcc [1 x i32] @callee_St4x1(i32 %in.0.val) { 104 ; CHECK: .func (.param .align 16 .b8 func_retval0[4]) 105 ; CHECK-LABEL: callee_St4x1( 106 ; CHECK-NEXT: .param .b32 callee_St4x1_param_0 107 ; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [callee_St4x1_param_0]; 108 ; CHECK: st.param.b32 [func_retval0+0], [[R1]]; 109 ; CHECK-NEXT: ret; 110 %oldret = insertvalue [1 x i32] poison, i32 %in.0.val, 0 111 ret [1 x i32] %oldret 112} 113 114define dso_local void @caller_St4x2(%struct.St4x2* nocapture noundef readonly byval(%struct.St4x2) align 4 %in, %struct.St4x2* nocapture noundef writeonly %ret) { 115 ; CHECK-LABEL: .visible .func caller_St4x2( 116 ; CHECK: .param .align 4 .b8 caller_St4x2_param_0[8], 117 ; CHECK: .param .b32 caller_St4x2_param_1 118 ; CHECK: ) 119 ; CHECK: .param .align 16 .b8 param0[8]; 120 ; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}}; 121 ; CHECK: .param .align 16 .b8 retval0[8]; 122 ; CHECK: call.uni (retval0), 123 ; CHECK-NEXT: callee_St4x2, 124 ; CHECK-NEXT: ( 125 ; CHECK-NEXT: param0 126 ; CHECK-NEXT: ); 127 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 128 %agg.tmp = alloca i64, align 8 129 %tmpcast = bitcast i64* %agg.tmp to %struct.St4x2* 130 %1 = bitcast %struct.St4x2* %in to i64* 131 %2 = load i64, i64* %1, align 4 132 store i64 %2, i64* %agg.tmp, align 8 133 %call = tail call fastcc [2 x i32] @callee_St4x2(%struct.St4x2* noundef nonnull byval(%struct.St4x2) align 4 %tmpcast) 134 %.fca.0.extract = extractvalue [2 x i32] %call, 0 135 %.fca.1.extract = extractvalue [2 x i32] %call, 1 136 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x2, %struct.St4x2* %ret, i64 0, i32 0, i64 0 137 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 138 %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St4x2, %struct.St4x2* %ret, i64 0, i32 0, i64 1 139 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx3, align 4 140 ret void 141} 142 143define internal fastcc [2 x i32] @callee_St4x2(%struct.St4x2* nocapture noundef readonly byval(%struct.St4x2) align 4 %in) { 144 ; CHECK: .func (.param .align 16 .b8 func_retval0[8]) 145 ; CHECK-LABEL: callee_St4x2( 146 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x2_param_0[8] 147 ; CHECK: ld.param.v2.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x2_param_0]; 148 ; CHECK: st.param.v2.b32 [func_retval0+0], {[[R1]], [[R2]]}; 149 ; CHECK-NEXT: ret; 150 %arrayidx = getelementptr inbounds %struct.St4x2, %struct.St4x2* %in, i64 0, i32 0, i64 0 151 %1 = load i32, i32* %arrayidx, align 4 152 %arrayidx.1 = getelementptr inbounds %struct.St4x2, %struct.St4x2* %in, i64 0, i32 0, i64 1 153 %2 = load i32, i32* %arrayidx.1, align 4 154 %3 = insertvalue [2 x i32] poison, i32 %1, 0 155 %oldret = insertvalue [2 x i32] %3, i32 %2, 1 156 ret [2 x i32] %oldret 157} 158 159define dso_local void @caller_St4x3(%struct.St4x3* nocapture noundef readonly byval(%struct.St4x3) align 4 %in, %struct.St4x3* nocapture noundef writeonly %ret) { 160 ; CHECK-LABEL: .visible .func caller_St4x3( 161 ; CHECK: .param .align 4 .b8 caller_St4x3_param_0[12], 162 ; CHECK: .param .b32 caller_St4x3_param_1 163 ; CHECK: ) 164 ; CHECK: .param .align 16 .b8 param0[12]; 165 ; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}}; 166 ; CHECK: st.param.b32 [param0+8], {{%r[0-9]+}}; 167 ; CHECK: .param .align 16 .b8 retval0[12]; 168 ; CHECK: call.uni (retval0), 169 ; CHECK-NEXT: callee_St4x3, 170 ; CHECK-NEXT: ( 171 ; CHECK-NEXT: param0 172 ; CHECK-NEXT: ); 173 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 174 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+8]; 175 %call = tail call fastcc [3 x i32] @callee_St4x3(%struct.St4x3* noundef nonnull byval(%struct.St4x3) align 4 %in) 176 %.fca.0.extract = extractvalue [3 x i32] %call, 0 177 %.fca.1.extract = extractvalue [3 x i32] %call, 1 178 %.fca.2.extract = extractvalue [3 x i32] %call, 2 179 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x3, %struct.St4x3* %ret, i64 0, i32 0, i64 0 180 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 181 %ref.tmp.sroa.4.0..sroa_idx2 = getelementptr inbounds %struct.St4x3, %struct.St4x3* %ret, i64 0, i32 0, i64 1 182 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx2, align 4 183 %ref.tmp.sroa.5.0..sroa_idx4 = getelementptr inbounds %struct.St4x3, %struct.St4x3* %ret, i64 0, i32 0, i64 2 184 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx4, align 4 185 ret void 186} 187 188 189define internal fastcc [3 x i32] @callee_St4x3(%struct.St4x3* nocapture noundef readonly byval(%struct.St4x3) align 4 %in) { 190 ; CHECK: .func (.param .align 16 .b8 func_retval0[12]) 191 ; CHECK-LABEL: callee_St4x3( 192 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x3_param_0[12] 193 ; CHECK: ld.param.v2.u32 {[[R1:%r[0-9]+]], [[R2:%r[0-9]+]]}, [callee_St4x3_param_0]; 194 ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [callee_St4x3_param_0+8]; 195 ; CHECK: st.param.v2.b32 [func_retval0+0], {[[R1]], [[R2]]}; 196 ; CHECK: st.param.b32 [func_retval0+8], [[R3]]; 197 ; CHECK-NEXT: ret; 198 %arrayidx = getelementptr inbounds %struct.St4x3, %struct.St4x3* %in, i64 0, i32 0, i64 0 199 %1 = load i32, i32* %arrayidx, align 4 200 %arrayidx.1 = getelementptr inbounds %struct.St4x3, %struct.St4x3* %in, i64 0, i32 0, i64 1 201 %2 = load i32, i32* %arrayidx.1, align 4 202 %arrayidx.2 = getelementptr inbounds %struct.St4x3, %struct.St4x3* %in, i64 0, i32 0, i64 2 203 %3 = load i32, i32* %arrayidx.2, align 4 204 %4 = insertvalue [3 x i32] poison, i32 %1, 0 205 %5 = insertvalue [3 x i32] %4, i32 %2, 1 206 %oldret = insertvalue [3 x i32] %5, i32 %3, 2 207 ret [3 x i32] %oldret 208} 209 210 211define dso_local void @caller_St4x4(%struct.St4x4* nocapture noundef readonly byval(%struct.St4x4) align 4 %in, %struct.St4x4* nocapture noundef writeonly %ret) { 212 ; CHECK-LABEL: .visible .func caller_St4x4( 213 ; CHECK: .param .align 4 .b8 caller_St4x4_param_0[16], 214 ; CHECK: .param .b32 caller_St4x4_param_1 215 ; CHECK: ) 216 ; CHECK: .param .align 16 .b8 param0[16]; 217 ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 218 ; CHECK: .param .align 16 .b8 retval0[16]; 219 ; CHECK: call.uni (retval0), 220 ; CHECK-NEXT: callee_St4x4, 221 ; CHECK-NEXT: ( 222 ; CHECK-NEXT: param0 223 ; CHECK-NEXT: ); 224 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 225 %call = tail call fastcc [4 x i32] @callee_St4x4(%struct.St4x4* noundef nonnull byval(%struct.St4x4) align 4 %in) 226 %.fca.0.extract = extractvalue [4 x i32] %call, 0 227 %.fca.1.extract = extractvalue [4 x i32] %call, 1 228 %.fca.2.extract = extractvalue [4 x i32] %call, 2 229 %.fca.3.extract = extractvalue [4 x i32] %call, 3 230 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x4, %struct.St4x4* %ret, i64 0, i32 0, i64 0 231 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 232 %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %ret, i64 0, i32 0, i64 1 233 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx3, align 4 234 %ref.tmp.sroa.5.0..sroa_idx5 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %ret, i64 0, i32 0, i64 2 235 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx5, align 4 236 %ref.tmp.sroa.6.0..sroa_idx7 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %ret, i64 0, i32 0, i64 3 237 store i32 %.fca.3.extract, i32* %ref.tmp.sroa.6.0..sroa_idx7, align 4 238 ret void 239} 240 241 242define internal fastcc [4 x i32] @callee_St4x4(%struct.St4x4* nocapture noundef readonly byval(%struct.St4x4) align 4 %in) { 243 ; CHECK: .func (.param .align 16 .b8 func_retval0[16]) 244 ; CHECK-LABEL: callee_St4x4( 245 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_param_0[16] 246 ; 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]; 247 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 248 ; CHECK-NEXT: ret; 249 %arrayidx = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 0 250 %1 = load i32, i32* %arrayidx, align 4 251 %arrayidx.1 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 1 252 %2 = load i32, i32* %arrayidx.1, align 4 253 %arrayidx.2 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 2 254 %3 = load i32, i32* %arrayidx.2, align 4 255 %arrayidx.3 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 3 256 %4 = load i32, i32* %arrayidx.3, align 4 257 %5 = insertvalue [4 x i32] poison, i32 %1, 0 258 %6 = insertvalue [4 x i32] %5, i32 %2, 1 259 %7 = insertvalue [4 x i32] %6, i32 %3, 2 260 %oldret = insertvalue [4 x i32] %7, i32 %4, 3 261 ret [4 x i32] %oldret 262} 263 264 265define dso_local void @caller_St4x5(%struct.St4x5* nocapture noundef readonly byval(%struct.St4x5) align 4 %in, %struct.St4x5* nocapture noundef writeonly %ret) { 266 ; CHECK-LABEL: .visible .func caller_St4x5( 267 ; CHECK: .param .align 4 .b8 caller_St4x5_param_0[20], 268 ; CHECK: .param .b32 caller_St4x5_param_1 269 ; CHECK: ) 270 ; CHECK: .param .align 16 .b8 param0[20]; 271 ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 272 ; CHECK: st.param.b32 [param0+16], {{%r[0-9]+}}; 273 ; CHECK: .param .align 16 .b8 retval0[20]; 274 ; CHECK: call.uni (retval0), 275 ; CHECK-NEXT: callee_St4x5, 276 ; CHECK-NEXT: ( 277 ; CHECK-NEXT: param0 278 ; CHECK-NEXT: ); 279 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 280 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+16]; 281 %call = tail call fastcc [5 x i32] @callee_St4x5(%struct.St4x5* noundef nonnull byval(%struct.St4x5) align 4 %in) 282 %.fca.0.extract = extractvalue [5 x i32] %call, 0 283 %.fca.1.extract = extractvalue [5 x i32] %call, 1 284 %.fca.2.extract = extractvalue [5 x i32] %call, 2 285 %.fca.3.extract = extractvalue [5 x i32] %call, 3 286 %.fca.4.extract = extractvalue [5 x i32] %call, 4 287 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x5, %struct.St4x5* %ret, i64 0, i32 0, i64 0 288 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 289 %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %ret, i64 0, i32 0, i64 1 290 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx3, align 4 291 %ref.tmp.sroa.5.0..sroa_idx5 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %ret, i64 0, i32 0, i64 2 292 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx5, align 4 293 %ref.tmp.sroa.6.0..sroa_idx7 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %ret, i64 0, i32 0, i64 3 294 store i32 %.fca.3.extract, i32* %ref.tmp.sroa.6.0..sroa_idx7, align 4 295 %ref.tmp.sroa.7.0..sroa_idx9 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %ret, i64 0, i32 0, i64 4 296 store i32 %.fca.4.extract, i32* %ref.tmp.sroa.7.0..sroa_idx9, align 4 297 ret void 298} 299 300 301define internal fastcc [5 x i32] @callee_St4x5(%struct.St4x5* nocapture noundef readonly byval(%struct.St4x5) align 4 %in) { 302 ; CHECK: .func (.param .align 16 .b8 func_retval0[20]) 303 ; CHECK-LABEL: callee_St4x5( 304 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x5_param_0[20] 305 ; 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]; 306 ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [callee_St4x5_param_0+16]; 307 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 308 ; CHECK: st.param.b32 [func_retval0+16], [[R5]]; 309 ; CHECK-NEXT: ret; 310 %arrayidx = getelementptr inbounds %struct.St4x5, %struct.St4x5* %in, i64 0, i32 0, i64 0 311 %1 = load i32, i32* %arrayidx, align 4 312 %arrayidx.1 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %in, i64 0, i32 0, i64 1 313 %2 = load i32, i32* %arrayidx.1, align 4 314 %arrayidx.2 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %in, i64 0, i32 0, i64 2 315 %3 = load i32, i32* %arrayidx.2, align 4 316 %arrayidx.3 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %in, i64 0, i32 0, i64 3 317 %4 = load i32, i32* %arrayidx.3, align 4 318 %arrayidx.4 = getelementptr inbounds %struct.St4x5, %struct.St4x5* %in, i64 0, i32 0, i64 4 319 %5 = load i32, i32* %arrayidx.4, align 4 320 %6 = insertvalue [5 x i32] poison, i32 %1, 0 321 %7 = insertvalue [5 x i32] %6, i32 %2, 1 322 %8 = insertvalue [5 x i32] %7, i32 %3, 2 323 %9 = insertvalue [5 x i32] %8, i32 %4, 3 324 %oldret = insertvalue [5 x i32] %9, i32 %5, 4 325 ret [5 x i32] %oldret 326} 327 328 329define dso_local void @caller_St4x6(%struct.St4x6* nocapture noundef readonly byval(%struct.St4x6) align 4 %in, %struct.St4x6* nocapture noundef writeonly %ret) { 330 ; CHECK-LABEL: .visible .func caller_St4x6( 331 ; CHECK: .param .align 4 .b8 caller_St4x6_param_0[24], 332 ; CHECK: .param .b32 caller_St4x6_param_1 333 ; CHECK: ) 334 ; CHECK: .param .align 16 .b8 param0[24]; 335 ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 336 ; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}}; 337 ; CHECK: .param .align 16 .b8 retval0[24]; 338 ; CHECK: call.uni (retval0), 339 ; CHECK-NEXT: callee_St4x6, 340 ; CHECK-NEXT: ( 341 ; CHECK-NEXT: param0 342 ; CHECK-NEXT: ); 343 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 344 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16]; 345 %call = tail call fastcc [6 x i32] @callee_St4x6(%struct.St4x6* noundef nonnull byval(%struct.St4x6) align 4 %in) 346 %.fca.0.extract = extractvalue [6 x i32] %call, 0 347 %.fca.1.extract = extractvalue [6 x i32] %call, 1 348 %.fca.2.extract = extractvalue [6 x i32] %call, 2 349 %.fca.3.extract = extractvalue [6 x i32] %call, 3 350 %.fca.4.extract = extractvalue [6 x i32] %call, 4 351 %.fca.5.extract = extractvalue [6 x i32] %call, 5 352 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 0 353 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 354 %ref.tmp.sroa.4.0..sroa_idx2 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 1 355 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx2, align 4 356 %ref.tmp.sroa.5.0..sroa_idx4 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 2 357 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx4, align 4 358 %ref.tmp.sroa.6.0..sroa_idx6 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 3 359 store i32 %.fca.3.extract, i32* %ref.tmp.sroa.6.0..sroa_idx6, align 4 360 %ref.tmp.sroa.7.0..sroa_idx8 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 4 361 store i32 %.fca.4.extract, i32* %ref.tmp.sroa.7.0..sroa_idx8, align 4 362 %ref.tmp.sroa.8.0..sroa_idx10 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %ret, i64 0, i32 0, i64 5 363 store i32 %.fca.5.extract, i32* %ref.tmp.sroa.8.0..sroa_idx10, align 4 364 ret void 365} 366 367 368define internal fastcc [6 x i32] @callee_St4x6(%struct.St4x6* nocapture noundef readonly byval(%struct.St4x6) align 4 %in) { 369 ; CHECK: .func (.param .align 16 .b8 func_retval0[24]) 370 ; CHECK-LABEL: callee_St4x6( 371 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x6_param_0[24] 372 ; 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]; 373 ; CHECK: ld.param.v2.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x6_param_0+16]; 374 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 375 ; CHECK: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]}; 376 ; CHECK-NEXT: ret; 377 %arrayidx = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 0 378 %1 = load i32, i32* %arrayidx, align 4 379 %arrayidx.1 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 1 380 %2 = load i32, i32* %arrayidx.1, align 4 381 %arrayidx.2 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 2 382 %3 = load i32, i32* %arrayidx.2, align 4 383 %arrayidx.3 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 3 384 %4 = load i32, i32* %arrayidx.3, align 4 385 %arrayidx.4 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 4 386 %5 = load i32, i32* %arrayidx.4, align 4 387 %arrayidx.5 = getelementptr inbounds %struct.St4x6, %struct.St4x6* %in, i64 0, i32 0, i64 5 388 %6 = load i32, i32* %arrayidx.5, align 4 389 %7 = insertvalue [6 x i32] poison, i32 %1, 0 390 %8 = insertvalue [6 x i32] %7, i32 %2, 1 391 %9 = insertvalue [6 x i32] %8, i32 %3, 2 392 %10 = insertvalue [6 x i32] %9, i32 %4, 3 393 %11 = insertvalue [6 x i32] %10, i32 %5, 4 394 %oldret = insertvalue [6 x i32] %11, i32 %6, 5 395 ret [6 x i32] %oldret 396} 397 398 399define dso_local void @caller_St4x7(%struct.St4x7* nocapture noundef readonly byval(%struct.St4x7) align 4 %in, %struct.St4x7* nocapture noundef writeonly %ret) { 400 ; CHECK-LABEL: .visible .func caller_St4x7( 401 ; CHECK: .param .align 4 .b8 caller_St4x7_param_0[28], 402 ; CHECK: .param .b32 caller_St4x7_param_1 403 ; CHECK: ) 404 ; CHECK: .param .align 16 .b8 param0[28]; 405 ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 406 ; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}}; 407 ; CHECK: st.param.b32 [param0+24], {{%r[0-9]+}}; 408 ; CHECK: .param .align 16 .b8 retval0[28]; 409 ; CHECK: call.uni (retval0), 410 ; CHECK-NEXT: callee_St4x7, 411 ; CHECK-NEXT: ( 412 ; CHECK-NEXT: param0 413 ; CHECK-NEXT: ); 414 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 415 ; CHECK: ld.param.v2.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16]; 416 ; CHECK: ld.param.b32 {{%r[0-9]+}}, [retval0+24]; 417 %call = tail call fastcc [7 x i32] @callee_St4x7(%struct.St4x7* noundef nonnull byval(%struct.St4x7) align 4 %in) 418 %.fca.0.extract = extractvalue [7 x i32] %call, 0 419 %.fca.1.extract = extractvalue [7 x i32] %call, 1 420 %.fca.2.extract = extractvalue [7 x i32] %call, 2 421 %.fca.3.extract = extractvalue [7 x i32] %call, 3 422 %.fca.4.extract = extractvalue [7 x i32] %call, 4 423 %.fca.5.extract = extractvalue [7 x i32] %call, 5 424 %.fca.6.extract = extractvalue [7 x i32] %call, 6 425 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 0 426 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 427 %ref.tmp.sroa.4.0..sroa_idx2 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 1 428 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx2, align 4 429 %ref.tmp.sroa.5.0..sroa_idx4 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 2 430 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx4, align 4 431 %ref.tmp.sroa.6.0..sroa_idx6 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 3 432 store i32 %.fca.3.extract, i32* %ref.tmp.sroa.6.0..sroa_idx6, align 4 433 %ref.tmp.sroa.7.0..sroa_idx8 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 4 434 store i32 %.fca.4.extract, i32* %ref.tmp.sroa.7.0..sroa_idx8, align 4 435 %ref.tmp.sroa.8.0..sroa_idx10 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 5 436 store i32 %.fca.5.extract, i32* %ref.tmp.sroa.8.0..sroa_idx10, align 4 437 %ref.tmp.sroa.9.0..sroa_idx12 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %ret, i64 0, i32 0, i64 6 438 store i32 %.fca.6.extract, i32* %ref.tmp.sroa.9.0..sroa_idx12, align 4 439 ret void 440} 441 442 443define internal fastcc [7 x i32] @callee_St4x7(%struct.St4x7* nocapture noundef readonly byval(%struct.St4x7) align 4 %in) { 444 ; CHECK: .func (.param .align 16 .b8 func_retval0[28]) 445 ; CHECK-LABEL: callee_St4x7( 446 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x7_param_0[28] 447 ; 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]; 448 ; CHECK: ld.param.v2.u32 {[[R5:%r[0-9]+]], [[R6:%r[0-9]+]]}, [callee_St4x7_param_0+16]; 449 ; CHECK: ld.param.u32 [[R7:%r[0-9]+]], [callee_St4x7_param_0+24]; 450 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 451 ; CHECK: st.param.v2.b32 [func_retval0+16], {[[R5]], [[R6]]}; 452 ; CHECK: st.param.b32 [func_retval0+24], [[R7]]; 453 ; CHECK-NEXT: ret; 454 %arrayidx = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 0 455 %1 = load i32, i32* %arrayidx, align 4 456 %arrayidx.1 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 1 457 %2 = load i32, i32* %arrayidx.1, align 4 458 %arrayidx.2 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 2 459 %3 = load i32, i32* %arrayidx.2, align 4 460 %arrayidx.3 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 3 461 %4 = load i32, i32* %arrayidx.3, align 4 462 %arrayidx.4 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 4 463 %5 = load i32, i32* %arrayidx.4, align 4 464 %arrayidx.5 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 5 465 %6 = load i32, i32* %arrayidx.5, align 4 466 %arrayidx.6 = getelementptr inbounds %struct.St4x7, %struct.St4x7* %in, i64 0, i32 0, i64 6 467 %7 = load i32, i32* %arrayidx.6, align 4 468 %8 = insertvalue [7 x i32] poison, i32 %1, 0 469 %9 = insertvalue [7 x i32] %8, i32 %2, 1 470 %10 = insertvalue [7 x i32] %9, i32 %3, 2 471 %11 = insertvalue [7 x i32] %10, i32 %4, 3 472 %12 = insertvalue [7 x i32] %11, i32 %5, 4 473 %13 = insertvalue [7 x i32] %12, i32 %6, 5 474 %oldret = insertvalue [7 x i32] %13, i32 %7, 6 475 ret [7 x i32] %oldret 476} 477 478 479define dso_local void @caller_St4x8(%struct.St4x8* nocapture noundef readonly byval(%struct.St4x8) align 4 %in, %struct.St4x8* nocapture noundef writeonly %ret) { 480 ; CHECK-LABEL: .visible .func caller_St4x8( 481 ; CHECK: .param .align 4 .b8 caller_St4x8_param_0[32], 482 ; CHECK: .param .b32 caller_St4x8_param_1 483 ; CHECK: ) 484 ; CHECK: .param .align 16 .b8 param0[32]; 485 ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 486 ; CHECK: st.param.v4.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; 487 ; CHECK: .param .align 16 .b8 retval0[32]; 488 ; CHECK: call.uni (retval0), 489 ; CHECK-NEXT: callee_St4x8, 490 ; CHECK-NEXT: ( 491 ; CHECK-NEXT: param0 492 ; CHECK-NEXT: ); 493 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+0]; 494 ; CHECK: ld.param.v4.b32 {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}, [retval0+16]; 495 %call = tail call fastcc [8 x i32] @callee_St4x8(%struct.St4x8* noundef nonnull byval(%struct.St4x8) align 4 %in) 496 %.fca.0.extract = extractvalue [8 x i32] %call, 0 497 %.fca.1.extract = extractvalue [8 x i32] %call, 1 498 %.fca.2.extract = extractvalue [8 x i32] %call, 2 499 %.fca.3.extract = extractvalue [8 x i32] %call, 3 500 %.fca.4.extract = extractvalue [8 x i32] %call, 4 501 %.fca.5.extract = extractvalue [8 x i32] %call, 5 502 %.fca.6.extract = extractvalue [8 x i32] %call, 6 503 %.fca.7.extract = extractvalue [8 x i32] %call, 7 504 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 0 505 store i32 %.fca.0.extract, i32* %ref.tmp.sroa.0.0..sroa_idx, align 4 506 %ref.tmp.sroa.4.0..sroa_idx2 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 1 507 store i32 %.fca.1.extract, i32* %ref.tmp.sroa.4.0..sroa_idx2, align 4 508 %ref.tmp.sroa.5.0..sroa_idx4 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 2 509 store i32 %.fca.2.extract, i32* %ref.tmp.sroa.5.0..sroa_idx4, align 4 510 %ref.tmp.sroa.6.0..sroa_idx6 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 3 511 store i32 %.fca.3.extract, i32* %ref.tmp.sroa.6.0..sroa_idx6, align 4 512 %ref.tmp.sroa.7.0..sroa_idx8 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 4 513 store i32 %.fca.4.extract, i32* %ref.tmp.sroa.7.0..sroa_idx8, align 4 514 %ref.tmp.sroa.8.0..sroa_idx10 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 5 515 store i32 %.fca.5.extract, i32* %ref.tmp.sroa.8.0..sroa_idx10, align 4 516 %ref.tmp.sroa.9.0..sroa_idx12 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 6 517 store i32 %.fca.6.extract, i32* %ref.tmp.sroa.9.0..sroa_idx12, align 4 518 %ref.tmp.sroa.10.0..sroa_idx14 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %ret, i64 0, i32 0, i64 7 519 store i32 %.fca.7.extract, i32* %ref.tmp.sroa.10.0..sroa_idx14, align 4 520 ret void 521} 522 523 524define internal fastcc [8 x i32] @callee_St4x8(%struct.St4x8* nocapture noundef readonly byval(%struct.St4x8) align 4 %in) { 525 ; CHECK: .func (.param .align 16 .b8 func_retval0[32]) 526 ; CHECK-LABEL: callee_St4x8( 527 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x8_param_0[32] 528 ; 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]; 529 ; 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]; 530 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 531 ; CHECK: st.param.v4.b32 [func_retval0+16], {[[R5]], [[R6]], [[R7]], [[R8]]}; 532 ; CHECK-NEXT: ret; 533 %arrayidx = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 0 534 %1 = load i32, i32* %arrayidx, align 4 535 %arrayidx.1 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 1 536 %2 = load i32, i32* %arrayidx.1, align 4 537 %arrayidx.2 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 2 538 %3 = load i32, i32* %arrayidx.2, align 4 539 %arrayidx.3 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 3 540 %4 = load i32, i32* %arrayidx.3, align 4 541 %arrayidx.4 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 4 542 %5 = load i32, i32* %arrayidx.4, align 4 543 %arrayidx.5 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 5 544 %6 = load i32, i32* %arrayidx.5, align 4 545 %arrayidx.6 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 6 546 %7 = load i32, i32* %arrayidx.6, align 4 547 %arrayidx.7 = getelementptr inbounds %struct.St4x8, %struct.St4x8* %in, i64 0, i32 0, i64 7 548 %8 = load i32, i32* %arrayidx.7, align 4 549 %9 = insertvalue [8 x i32] poison, i32 %1, 0 550 %10 = insertvalue [8 x i32] %9, i32 %2, 1 551 %11 = insertvalue [8 x i32] %10, i32 %3, 2 552 %12 = insertvalue [8 x i32] %11, i32 %4, 3 553 %13 = insertvalue [8 x i32] %12, i32 %5, 4 554 %14 = insertvalue [8 x i32] %13, i32 %6, 5 555 %15 = insertvalue [8 x i32] %14, i32 %7, 6 556 %oldret = insertvalue [8 x i32] %15, i32 %8, 7 557 ret [8 x i32] %oldret 558} 559 560 561define dso_local void @caller_St8x1(%struct.St8x1* nocapture noundef readonly byval(%struct.St8x1) align 8 %in, %struct.St8x1* nocapture noundef writeonly %ret) { 562 ; CHECK-LABEL: .visible .func caller_St8x1( 563 ; CHECK: .param .align 8 .b8 caller_St8x1_param_0[8], 564 ; CHECK: .param .b32 caller_St8x1_param_1 565 ; CHECK: ) 566 ; CHECK: .param .b64 param0; 567 ; CHECK: st.param.b64 [param0+0], {{%rd[0-9]+}}; 568 ; CHECK: .param .align 16 .b8 retval0[8]; 569 ; CHECK: call.uni (retval0), 570 ; CHECK-NEXT: callee_St8x1, 571 ; CHECK-NEXT: ( 572 ; CHECK-NEXT: param0 573 ; CHECK-NEXT: ); 574 ; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0+0]; 575 %1 = getelementptr inbounds %struct.St8x1, %struct.St8x1* %in, i64 0, i32 0, i64 0 576 %2 = load i64, i64* %1, align 8 577 %call = tail call fastcc [1 x i64] @callee_St8x1(i64 %2) 578 %.fca.0.extract = extractvalue [1 x i64] %call, 0 579 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x1, %struct.St8x1* %ret, i64 0, i32 0, i64 0 580 store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 581 ret void 582} 583 584 585define internal fastcc [1 x i64] @callee_St8x1(i64 %in.0.val) { 586 ; CHECK: .func (.param .align 16 .b8 func_retval0[8]) 587 ; CHECK-LABEL: callee_St8x1( 588 ; CHECK-NEXT: .param .b64 callee_St8x1_param_0 589 ; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [callee_St8x1_param_0]; 590 ; CHECK: st.param.b64 [func_retval0+0], [[RD1]]; 591 ; CHECK-NEXT: ret; 592 %oldret = insertvalue [1 x i64] poison, i64 %in.0.val, 0 593 ret [1 x i64] %oldret 594} 595 596 597define dso_local void @caller_St8x2(%struct.St8x2* nocapture noundef readonly byval(%struct.St8x2) align 8 %in, %struct.St8x2* nocapture noundef writeonly %ret) { 598 ; CHECK-LABEL: .visible .func caller_St8x2( 599 ; CHECK: .param .align 8 .b8 caller_St8x2_param_0[16], 600 ; CHECK: .param .b32 caller_St8x2_param_1 601 ; CHECK: ) 602 ; CHECK: .param .align 16 .b8 param0[16]; 603 ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; 604 ; CHECK: .param .align 16 .b8 retval0[16]; 605 ; CHECK: call.uni (retval0), 606 ; CHECK-NEXT: callee_St8x2, 607 ; CHECK-NEXT: ( 608 ; CHECK-NEXT: param0 609 ; CHECK-NEXT: ); 610 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0]; 611 %call = tail call fastcc [2 x i64] @callee_St8x2(%struct.St8x2* noundef nonnull byval(%struct.St8x2) align 8 %in) 612 %.fca.0.extract = extractvalue [2 x i64] %call, 0 613 %.fca.1.extract = extractvalue [2 x i64] %call, 1 614 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x2, %struct.St8x2* %ret, i64 0, i32 0, i64 0 615 store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 616 %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St8x2, %struct.St8x2* %ret, i64 0, i32 0, i64 1 617 store i64 %.fca.1.extract, i64* %ref.tmp.sroa.4.0..sroa_idx3, align 8 618 ret void 619} 620 621 622define internal fastcc [2 x i64] @callee_St8x2(%struct.St8x2* nocapture noundef readonly byval(%struct.St8x2) align 8 %in) { 623 ; CHECK: .func (.param .align 16 .b8 func_retval0[16]) 624 ; CHECK-LABEL: callee_St8x2( 625 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x2_param_0[16] 626 ; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x2_param_0]; 627 ; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]}; 628 ; CHECK-NEXT: ret; 629 %arrayidx = getelementptr inbounds %struct.St8x2, %struct.St8x2* %in, i64 0, i32 0, i64 0 630 %1 = load i64, i64* %arrayidx, align 8 631 %arrayidx.1 = getelementptr inbounds %struct.St8x2, %struct.St8x2* %in, i64 0, i32 0, i64 1 632 %2 = load i64, i64* %arrayidx.1, align 8 633 %3 = insertvalue [2 x i64] poison, i64 %1, 0 634 %oldret = insertvalue [2 x i64] %3, i64 %2, 1 635 ret [2 x i64] %oldret 636} 637 638 639define dso_local void @caller_St8x3(%struct.St8x3* nocapture noundef readonly byval(%struct.St8x3) align 8 %in, %struct.St8x3* nocapture noundef writeonly %ret) { 640 ; CHECK-LABEL: .visible .func caller_St8x3( 641 ; CHECK: .param .align 8 .b8 caller_St8x3_param_0[24], 642 ; CHECK: .param .b32 caller_St8x3_param_1 643 ; CHECK: ) 644 ; CHECK: .param .align 16 .b8 param0[24]; 645 ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; 646 ; CHECK: st.param.b64 [param0+16], {{%rd[0-9]+}}; 647 ; CHECK: .param .align 16 .b8 retval0[24]; 648 ; CHECK: call.uni (retval0), 649 ; CHECK-NEXT: callee_St8x3, 650 ; CHECK-NEXT: ( 651 ; CHECK-NEXT: param0 652 ; CHECK-NEXT: ); 653 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0]; 654 ; CHECK: ld.param.b64 {{%rd[0-9]+}}, [retval0+16]; 655 %call = tail call fastcc [3 x i64] @callee_St8x3(%struct.St8x3* noundef nonnull byval(%struct.St8x3) align 8 %in) 656 %.fca.0.extract = extractvalue [3 x i64] %call, 0 657 %.fca.1.extract = extractvalue [3 x i64] %call, 1 658 %.fca.2.extract = extractvalue [3 x i64] %call, 2 659 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x3, %struct.St8x3* %ret, i64 0, i32 0, i64 0 660 store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 661 %ref.tmp.sroa.4.0..sroa_idx2 = getelementptr inbounds %struct.St8x3, %struct.St8x3* %ret, i64 0, i32 0, i64 1 662 store i64 %.fca.1.extract, i64* %ref.tmp.sroa.4.0..sroa_idx2, align 8 663 %ref.tmp.sroa.5.0..sroa_idx4 = getelementptr inbounds %struct.St8x3, %struct.St8x3* %ret, i64 0, i32 0, i64 2 664 store i64 %.fca.2.extract, i64* %ref.tmp.sroa.5.0..sroa_idx4, align 8 665 ret void 666} 667 668 669define internal fastcc [3 x i64] @callee_St8x3(%struct.St8x3* nocapture noundef readonly byval(%struct.St8x3) align 8 %in) { 670 ; CHECK: .func (.param .align 16 .b8 func_retval0[24]) 671 ; CHECK-LABEL: callee_St8x3( 672 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x3_param_0[24] 673 ; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x3_param_0]; 674 ; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [callee_St8x3_param_0+16]; 675 ; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]}; 676 ; CHECK: st.param.b64 [func_retval0+16], [[RD3]]; 677 ; CHECK-NEXT: ret; 678 %arrayidx = getelementptr inbounds %struct.St8x3, %struct.St8x3* %in, i64 0, i32 0, i64 0 679 %1 = load i64, i64* %arrayidx, align 8 680 %arrayidx.1 = getelementptr inbounds %struct.St8x3, %struct.St8x3* %in, i64 0, i32 0, i64 1 681 %2 = load i64, i64* %arrayidx.1, align 8 682 %arrayidx.2 = getelementptr inbounds %struct.St8x3, %struct.St8x3* %in, i64 0, i32 0, i64 2 683 %3 = load i64, i64* %arrayidx.2, align 8 684 %4 = insertvalue [3 x i64] poison, i64 %1, 0 685 %5 = insertvalue [3 x i64] %4, i64 %2, 1 686 %oldret = insertvalue [3 x i64] %5, i64 %3, 2 687 ret [3 x i64] %oldret 688} 689 690 691define dso_local void @caller_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in, %struct.St8x4* nocapture noundef writeonly %ret) { 692 ; CHECK-LABEL: .visible .func caller_St8x4( 693 ; CHECK: .param .align 8 .b8 caller_St8x4_param_0[32], 694 ; CHECK: .param .b32 caller_St8x4_param_1 695 ; CHECK: ) 696 ; CHECK: .param .align 16 .b8 param0[32]; 697 ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; 698 ; CHECK: st.param.v2.b64 [param0+16], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; 699 ; CHECK: .param .align 16 .b8 retval0[32]; 700 ; CHECK: call.uni (retval0), 701 ; CHECK-NEXT: callee_St8x4, 702 ; CHECK-NEXT: ( 703 ; CHECK-NEXT: param0 704 ; CHECK-NEXT: ); 705 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+0]; 706 ; CHECK: ld.param.v2.b64 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [retval0+16]; 707 %call = tail call fastcc [4 x i64] @callee_St8x4(%struct.St8x4* noundef nonnull byval(%struct.St8x4) align 8 %in) 708 %.fca.0.extract = extractvalue [4 x i64] %call, 0 709 %.fca.1.extract = extractvalue [4 x i64] %call, 1 710 %.fca.2.extract = extractvalue [4 x i64] %call, 2 711 %.fca.3.extract = extractvalue [4 x i64] %call, 3 712 %ref.tmp.sroa.0.0..sroa_idx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 0 713 store i64 %.fca.0.extract, i64* %ref.tmp.sroa.0.0..sroa_idx, align 8 714 %ref.tmp.sroa.4.0..sroa_idx3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 1 715 store i64 %.fca.1.extract, i64* %ref.tmp.sroa.4.0..sroa_idx3, align 8 716 %ref.tmp.sroa.5.0..sroa_idx5 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 2 717 store i64 %.fca.2.extract, i64* %ref.tmp.sroa.5.0..sroa_idx5, align 8 718 %ref.tmp.sroa.6.0..sroa_idx7 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %ret, i64 0, i32 0, i64 3 719 store i64 %.fca.3.extract, i64* %ref.tmp.sroa.6.0..sroa_idx7, align 8 720 ret void 721} 722 723 724define internal fastcc [4 x i64] @callee_St8x4(%struct.St8x4* nocapture noundef readonly byval(%struct.St8x4) align 8 %in) { 725 ; CHECK: .func (.param .align 16 .b8 func_retval0[32]) 726 ; CHECK-LABEL: callee_St8x4( 727 ; CHECK-NEXT: .param .align 16 .b8 callee_St8x4_param_0[32] 728 ; CHECK: ld.param.v2.u64 {[[RD1:%rd[0-9]+]], [[RD2:%rd[0-9]+]]}, [callee_St8x4_param_0]; 729 ; CHECK: ld.param.v2.u64 {[[RD3:%rd[0-9]+]], [[RD4:%rd[0-9]+]]}, [callee_St8x4_param_0+16]; 730 ; CHECK: st.param.v2.b64 [func_retval0+0], {[[RD1]], [[RD2]]}; 731 ; CHECK: st.param.v2.b64 [func_retval0+16], {[[RD3]], [[RD4]]}; 732 ; CHECK-NEXT: ret; 733 %arrayidx = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 0 734 %1 = load i64, i64* %arrayidx, align 8 735 %arrayidx.1 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 1 736 %2 = load i64, i64* %arrayidx.1, align 8 737 %arrayidx.2 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 2 738 %3 = load i64, i64* %arrayidx.2, align 8 739 %arrayidx.3 = getelementptr inbounds %struct.St8x4, %struct.St8x4* %in, i64 0, i32 0, i64 3 740 %4 = load i64, i64* %arrayidx.3, align 8 741 %5 = insertvalue [4 x i64] poison, i64 %1, 0 742 %6 = insertvalue [4 x i64] %5, i64 %2, 1 743 %7 = insertvalue [4 x i64] %6, i64 %3, 2 744 %oldret = insertvalue [4 x i64] %7, i64 %4, 3 745 ret [4 x i64] %oldret 746} 747 748; Section 2 - checking that function argument (including retval) vectorization is done with private linkage. 749 750define private fastcc [4 x i32] @callee_St4x4_private(%struct.St4x4* nocapture noundef readonly byval(%struct.St4x4) align 4 %in) { 751 ; CHECK: .func (.param .align 16 .b8 func_retval0[16]) 752 ; CHECK-LABEL: callee_St4x4_private( 753 ; CHECK-NEXT: .param .align 16 .b8 callee_St4x4_private_param_0[16] 754 ; 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]; 755 ; CHECK: st.param.v4.b32 [func_retval0+0], {[[R1]], [[R2]], [[R3]], [[R4]]}; 756 ; CHECK-NEXT: ret; 757 %arrayidx = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 0 758 %1 = load i32, i32* %arrayidx, align 4 759 %arrayidx.1 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 1 760 %2 = load i32, i32* %arrayidx.1, align 4 761 %arrayidx.2 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 2 762 %3 = load i32, i32* %arrayidx.2, align 4 763 %arrayidx.3 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 3 764 %4 = load i32, i32* %arrayidx.3, align 4 765 %5 = insertvalue [4 x i32] poison, i32 %1, 0 766 %6 = insertvalue [4 x i32] %5, i32 %2, 1 767 %7 = insertvalue [4 x i32] %6, i32 %3, 2 768 %oldret = insertvalue [4 x i32] %7, i32 %4, 3 769 ret [4 x i32] %oldret 770} 771 772; Section 3 - checking that function argument (including retval) vectorization 773; is NOT done with linkage types other than internal and private. 774 775define external fastcc [4 x i32] @callee_St4x4_external(%struct.St4x4* nocapture noundef readonly byval(%struct.St4x4) align 4 %in) { 776 ; CHECK: .func (.param .align 4 .b8 func_retval0[16]) 777 ; CHECK-LABEL: callee_St4x4_external( 778 ; CHECK-NEXT: .param .align 4 .b8 callee_St4x4_external_param_0[16] 779 ; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [callee_St4x4_external_param_0]; 780 ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [callee_St4x4_external_param_0+4]; 781 ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [callee_St4x4_external_param_0+8]; 782 ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [callee_St4x4_external_param_0+12]; 783 ; CHECK: st.param.b32 [func_retval0+0], [[R1]]; 784 ; CHECK: st.param.b32 [func_retval0+4], [[R2]]; 785 ; CHECK: st.param.b32 [func_retval0+8], [[R3]]; 786 ; CHECK: st.param.b32 [func_retval0+12], [[R4]]; 787 ; CHECK-NEXT: ret; 788 %arrayidx = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 0 789 %1 = load i32, i32* %arrayidx, align 4 790 %arrayidx.1 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 1 791 %2 = load i32, i32* %arrayidx.1, align 4 792 %arrayidx.2 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 2 793 %3 = load i32, i32* %arrayidx.2, align 4 794 %arrayidx.3 = getelementptr inbounds %struct.St4x4, %struct.St4x4* %in, i64 0, i32 0, i64 3 795 %4 = load i32, i32* %arrayidx.3, align 4 796 %5 = insertvalue [4 x i32] poison, i32 %1, 0 797 %6 = insertvalue [4 x i32] %5, i32 %2, 1 798 %7 = insertvalue [4 x i32] %6, i32 %3, 2 799 %oldret = insertvalue [4 x i32] %7, i32 %4, 3 800 ret [4 x i32] %oldret 801} 802