1; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s
2;
3; Check that parameters of a __global__ (kernel) function do not get increased
4; alignment, and no additional vectorization is performed on loads/stores with
5; that parameters.
6;
7; Test IR is a minimized version of IR generated with the following command
8; from the source code below:
9; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu
10;
11; ----------------------------------------------------------------------------
12; #include <stdint.h>
13;
14; struct St4x1 { uint32_t field[1]; };
15; struct St4x2 { uint32_t field[2]; };
16; struct St4x3 { uint32_t field[3]; };
17; struct St4x4 { uint32_t field[4]; };
18; struct St4x5 { uint32_t field[5]; };
19; struct St4x6 { uint32_t field[6]; };
20; struct St4x7 { uint32_t field[7]; };
21; struct St4x8 { uint32_t field[8]; };
22; struct St8x1 { uint64_t field[1]; };
23; struct St8x2 { uint64_t field[2]; };
24; struct St8x3 { uint64_t field[3]; };
25; struct St8x4 { uint64_t field[4]; };
26;
27; #define DECLARE_FUNCTION(StName)                                    \
28; static __global__  __attribute__((noinline))                        \
29; void foo_##StName(struct StName in, struct StName* ret) {           \
30;   const unsigned size = sizeof(ret->field) / sizeof(*ret->field);   \
31;   for (unsigned i = 0; i != size; ++i)                              \
32;     ret->field[i] = in.field[i];                                    \
33; }                                                                   \
34;
35; DECLARE_FUNCTION(St4x1)
36; DECLARE_FUNCTION(St4x2)
37; DECLARE_FUNCTION(St4x3)
38; DECLARE_FUNCTION(St4x4)
39; DECLARE_FUNCTION(St4x5)
40; DECLARE_FUNCTION(St4x6)
41; DECLARE_FUNCTION(St4x7)
42; DECLARE_FUNCTION(St4x8)
43; DECLARE_FUNCTION(St8x1)
44; DECLARE_FUNCTION(St8x2)
45; DECLARE_FUNCTION(St8x3)
46; DECLARE_FUNCTION(St8x4)
47; ----------------------------------------------------------------------------
48
49%struct.St4x1 = type { [1 x i32] }
50%struct.St4x2 = type { [2 x i32] }
51%struct.St4x3 = type { [3 x i32] }
52%struct.St4x4 = type { [4 x i32] }
53%struct.St4x5 = type { [5 x i32] }
54%struct.St4x6 = type { [6 x i32] }
55%struct.St4x7 = type { [7 x i32] }
56%struct.St4x8 = type { [8 x i32] }
57%struct.St8x1 = type { [1 x i64] }
58%struct.St8x2 = type { [2 x i64] }
59%struct.St8x3 = type { [3 x i64] }
60%struct.St8x4 = type { [4 x i64] }
61
62define dso_local void @foo_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) {
63  ; CHECK-LABEL: .visible .func foo_St4x1(
64  ; CHECK:               .param .align 4 .b8 foo_St4x1_param_0[4],
65  ; CHECK:               .param .b32 foo_St4x1_param_1
66  ; CHECK:       )
67  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x1_param_1];
68  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x1_param_0];
69  ; CHECK:       st.u32  [[[R1]]], [[R2]];
70  ; CHECK:       ret;
71  %1 = load i32, ptr %in, align 4
72  store i32 %1, ptr %ret, align 4
73  ret void
74}
75
76define dso_local void @foo_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) {
77  ; CHECK-LABEL: .visible .func foo_St4x2(
78  ; CHECK:               .param .align 4 .b8 foo_St4x2_param_0[8],
79  ; CHECK:               .param .b32 foo_St4x2_param_1
80  ; CHECK:       )
81  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x2_param_1];
82  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x2_param_0];
83  ; CHECK:       st.u32  [[[R1]]], [[R2]];
84  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x2_param_0+4];
85  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
86  ; CHECK:       ret;
87  %1 = load i32, ptr %in, align 4
88  store i32 %1, ptr %ret, align 4
89  %arrayidx.1 = getelementptr inbounds [2 x i32], ptr %in, i64 0, i64 1
90  %2 = load i32, ptr %arrayidx.1, align 4
91  %arrayidx3.1 = getelementptr inbounds [2 x i32], ptr %ret, i64 0, i64 1
92  store i32 %2, ptr %arrayidx3.1, align 4
93  ret void
94}
95
96define dso_local void @foo_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) {
97  ; CHECK-LABEL: .visible .func foo_St4x3(
98  ; CHECK:               .param .align 4 .b8 foo_St4x3_param_0[12],
99  ; CHECK:               .param .b32 foo_St4x3_param_1
100  ; CHECK:       )
101  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x3_param_1];
102  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x3_param_0];
103  ; CHECK:       st.u32  [[[R1]]], [[R2]];
104  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x3_param_0+4];
105  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
106  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x3_param_0+8];
107  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
108  ; CHECK:       ret;
109  %1 = load i32, ptr %in, align 4
110  store i32 %1, ptr %ret, align 4
111  %arrayidx.1 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 1
112  %2 = load i32, ptr %arrayidx.1, align 4
113  %arrayidx3.1 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 1
114  store i32 %2, ptr %arrayidx3.1, align 4
115  %arrayidx.2 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 2
116  %3 = load i32, ptr %arrayidx.2, align 4
117  %arrayidx3.2 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 2
118  store i32 %3, ptr %arrayidx3.2, align 4
119  ret void
120}
121
122define dso_local void @foo_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {
123  ; CHECK-LABEL: .visible .func foo_St4x4(
124  ; CHECK:               .param .align 4 .b8 foo_St4x4_param_0[16],
125  ; CHECK:               .param .b32 foo_St4x4_param_1
126  ; CHECK:       )
127  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x4_param_1];
128  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x4_param_0];
129  ; CHECK:       st.u32  [[[R1]]], [[R2]];
130  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x4_param_0+4];
131  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
132  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x4_param_0+8];
133  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
134  ; CHECK:       ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x4_param_0+12];
135  ; CHECK:       st.u32  [[[R1]]+12], [[R5]];
136  ; CHECK:       ret;
137  %1 = load i32, ptr %in, align 4
138  store i32 %1, ptr %ret, align 4
139  %arrayidx.1 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 1
140  %2 = load i32, ptr %arrayidx.1, align 4
141  %arrayidx3.1 = getelementptr inbounds [4 x i32], ptr %ret, i64 0, i64 1
142  store i32 %2, ptr %arrayidx3.1, align 4
143  %arrayidx.2 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 2
144  %3 = load i32, ptr %arrayidx.2, align 4
145  %arrayidx3.2 = getelementptr inbounds [4 x i32], ptr %ret, i64 0, i64 2
146  store i32 %3, ptr %arrayidx3.2, align 4
147  %arrayidx.3 = getelementptr inbounds [4 x i32], ptr %in, i64 0, i64 3
148  %4 = load i32, ptr %arrayidx.3, align 4
149  %arrayidx3.3 = getelementptr inbounds [4 x i32], ptr %ret, i64 0, i64 3
150  store i32 %4, ptr %arrayidx3.3, align 4
151  ret void
152}
153
154define dso_local void @foo_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {
155  ; CHECK-LABEL: .visible .func foo_St4x5(
156  ; CHECK:               .param .align 4 .b8 foo_St4x5_param_0[20],
157  ; CHECK:               .param .b32 foo_St4x5_param_1
158  ; CHECK:       )
159  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x5_param_1];
160  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x5_param_0];
161  ; CHECK:       st.u32  [[[R1]]], [[R2]];
162  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x5_param_0+4];
163  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
164  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x5_param_0+8];
165  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
166  ; CHECK:       ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x5_param_0+12];
167  ; CHECK:       st.u32  [[[R1]]+12], [[R5]];
168  ; CHECK:       ld.param.u32 [[R6:%r[0-9]+]], [foo_St4x5_param_0+16];
169  ; CHECK:       st.u32  [[[R1]]+16], [[R6]];
170  ; CHECK:       ret;
171  %1 = load i32, ptr %in, align 4
172  store i32 %1, ptr %ret, align 4
173  %arrayidx.1 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 1
174  %2 = load i32, ptr %arrayidx.1, align 4
175  %arrayidx3.1 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 1
176  store i32 %2, ptr %arrayidx3.1, align 4
177  %arrayidx.2 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 2
178  %3 = load i32, ptr %arrayidx.2, align 4
179  %arrayidx3.2 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 2
180  store i32 %3, ptr %arrayidx3.2, align 4
181  %arrayidx.3 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 3
182  %4 = load i32, ptr %arrayidx.3, align 4
183  %arrayidx3.3 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 3
184  store i32 %4, ptr %arrayidx3.3, align 4
185  %arrayidx.4 = getelementptr inbounds [5 x i32], ptr %in, i64 0, i64 4
186  %5 = load i32, ptr %arrayidx.4, align 4
187  %arrayidx3.4 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 4
188  store i32 %5, ptr %arrayidx3.4, align 4
189  ret void
190}
191
192define dso_local void @foo_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {
193  ; CHECK-LABEL: .visible .func foo_St4x6(
194  ; CHECK:               .param .align 4 .b8 foo_St4x6_param_0[24],
195  ; CHECK:               .param .b32 foo_St4x6_param_1
196  ; CHECK:       )
197  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x6_param_1];
198  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x6_param_0];
199  ; CHECK:       st.u32  [[[R1]]], [[R2]];
200  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x6_param_0+4];
201  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
202  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x6_param_0+8];
203  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
204  ; CHECK:       ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x6_param_0+12];
205  ; CHECK:       st.u32  [[[R1]]+12], [[R5]];
206  ; CHECK:       ld.param.u32 [[R6:%r[0-9]+]], [foo_St4x6_param_0+16];
207  ; CHECK:       st.u32  [[[R1]]+16], [[R6]];
208  ; CHECK:       ld.param.u32 [[R7:%r[0-9]+]], [foo_St4x6_param_0+20];
209  ; CHECK:       st.u32  [[[R1]]+20], [[R7]];
210  ; CHECK:       ret;
211  %1 = load i32, ptr %in, align 4
212  store i32 %1, ptr %ret, align 4
213  %arrayidx.1 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 1
214  %2 = load i32, ptr %arrayidx.1, align 4
215  %arrayidx3.1 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 1
216  store i32 %2, ptr %arrayidx3.1, align 4
217  %arrayidx.2 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 2
218  %3 = load i32, ptr %arrayidx.2, align 4
219  %arrayidx3.2 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 2
220  store i32 %3, ptr %arrayidx3.2, align 4
221  %arrayidx.3 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 3
222  %4 = load i32, ptr %arrayidx.3, align 4
223  %arrayidx3.3 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 3
224  store i32 %4, ptr %arrayidx3.3, align 4
225  %arrayidx.4 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 4
226  %5 = load i32, ptr %arrayidx.4, align 4
227  %arrayidx3.4 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 4
228  store i32 %5, ptr %arrayidx3.4, align 4
229  %arrayidx.5 = getelementptr inbounds [6 x i32], ptr %in, i64 0, i64 5
230  %6 = load i32, ptr %arrayidx.5, align 4
231  %arrayidx3.5 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 5
232  store i32 %6, ptr %arrayidx3.5, align 4
233  ret void
234}
235
236define dso_local void @foo_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {
237  ; CHECK-LABEL: .visible .func foo_St4x7(
238  ; CHECK:               .param .align 4 .b8 foo_St4x7_param_0[28],
239  ; CHECK:               .param .b32 foo_St4x7_param_1
240  ; CHECK:       )
241  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x7_param_1];
242  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x7_param_0];
243  ; CHECK:       st.u32  [[[R1]]], [[R2]];
244  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x7_param_0+4];
245  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
246  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x7_param_0+8];
247  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
248  ; CHECK:       ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x7_param_0+12];
249  ; CHECK:       st.u32  [[[R1]]+12], [[R5]];
250  ; CHECK:       ld.param.u32 [[R6:%r[0-9]+]], [foo_St4x7_param_0+16];
251  ; CHECK:       st.u32  [[[R1]]+16], [[R6]];
252  ; CHECK:       ld.param.u32 [[R7:%r[0-9]+]], [foo_St4x7_param_0+20];
253  ; CHECK:       st.u32  [[[R1]]+20], [[R7]];
254  ; CHECK:       ld.param.u32 [[R8:%r[0-9]+]], [foo_St4x7_param_0+24];
255  ; CHECK:       st.u32  [[[R1]]+24], [[R8]];
256  ; CHECK:       ret;
257  %1 = load i32, ptr %in, align 4
258  store i32 %1, ptr %ret, align 4
259  %arrayidx.1 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 1
260  %2 = load i32, ptr %arrayidx.1, align 4
261  %arrayidx3.1 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 1
262  store i32 %2, ptr %arrayidx3.1, align 4
263  %arrayidx.2 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 2
264  %3 = load i32, ptr %arrayidx.2, align 4
265  %arrayidx3.2 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 2
266  store i32 %3, ptr %arrayidx3.2, align 4
267  %arrayidx.3 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 3
268  %4 = load i32, ptr %arrayidx.3, align 4
269  %arrayidx3.3 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 3
270  store i32 %4, ptr %arrayidx3.3, align 4
271  %arrayidx.4 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 4
272  %5 = load i32, ptr %arrayidx.4, align 4
273  %arrayidx3.4 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 4
274  store i32 %5, ptr %arrayidx3.4, align 4
275  %arrayidx.5 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 5
276  %6 = load i32, ptr %arrayidx.5, align 4
277  %arrayidx3.5 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 5
278  store i32 %6, ptr %arrayidx3.5, align 4
279  %arrayidx.6 = getelementptr inbounds [7 x i32], ptr %in, i64 0, i64 6
280  %7 = load i32, ptr %arrayidx.6, align 4
281  %arrayidx3.6 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 6
282  store i32 %7, ptr %arrayidx3.6, align 4
283  ret void
284}
285
286define dso_local void @foo_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {
287  ; CHECK-LABEL: .visible .func foo_St4x8(
288  ; CHECK:               .param .align 4 .b8 foo_St4x8_param_0[32],
289  ; CHECK:               .param .b32 foo_St4x8_param_1
290  ; CHECK:       )
291  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x8_param_1];
292  ; CHECK:       ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x8_param_0];
293  ; CHECK:       st.u32  [[[R1]]], [[R2]];
294  ; CHECK:       ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x8_param_0+4];
295  ; CHECK:       st.u32  [[[R1]]+4], [[R3]];
296  ; CHECK:       ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x8_param_0+8];
297  ; CHECK:       st.u32  [[[R1]]+8], [[R4]];
298  ; CHECK:       ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x8_param_0+12];
299  ; CHECK:       st.u32  [[[R1]]+12], [[R5]];
300  ; CHECK:       ld.param.u32 [[R6:%r[0-9]+]], [foo_St4x8_param_0+16];
301  ; CHECK:       st.u32  [[[R1]]+16], [[R6]];
302  ; CHECK:       ld.param.u32 [[R7:%r[0-9]+]], [foo_St4x8_param_0+20];
303  ; CHECK:       st.u32  [[[R1]]+20], [[R7]];
304  ; CHECK:       ld.param.u32 [[R8:%r[0-9]+]], [foo_St4x8_param_0+24];
305  ; CHECK:       st.u32  [[[R1]]+24], [[R8]];
306  ; CHECK:       ld.param.u32 [[R9:%r[0-9]+]], [foo_St4x8_param_0+28];
307  ; CHECK:       st.u32  [[[R1]]+28], [[R9]];
308  ; CHECK:       ret;
309  %1 = load i32, ptr %in, align 4
310  store i32 %1, ptr %ret, align 4
311  %arrayidx.1 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 1
312  %2 = load i32, ptr %arrayidx.1, align 4
313  %arrayidx3.1 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 1
314  store i32 %2, ptr %arrayidx3.1, align 4
315  %arrayidx.2 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 2
316  %3 = load i32, ptr %arrayidx.2, align 4
317  %arrayidx3.2 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 2
318  store i32 %3, ptr %arrayidx3.2, align 4
319  %arrayidx.3 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 3
320  %4 = load i32, ptr %arrayidx.3, align 4
321  %arrayidx3.3 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 3
322  store i32 %4, ptr %arrayidx3.3, align 4
323  %arrayidx.4 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 4
324  %5 = load i32, ptr %arrayidx.4, align 4
325  %arrayidx3.4 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 4
326  store i32 %5, ptr %arrayidx3.4, align 4
327  %arrayidx.5 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 5
328  %6 = load i32, ptr %arrayidx.5, align 4
329  %arrayidx3.5 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 5
330  store i32 %6, ptr %arrayidx3.5, align 4
331  %arrayidx.6 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 6
332  %7 = load i32, ptr %arrayidx.6, align 4
333  %arrayidx3.6 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 6
334  store i32 %7, ptr %arrayidx3.6, align 4
335  %arrayidx.7 = getelementptr inbounds [8 x i32], ptr %in, i64 0, i64 7
336  %8 = load i32, ptr %arrayidx.7, align 4
337  %arrayidx3.7 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 7
338  store i32 %8, ptr %arrayidx3.7, align 4
339  ret void
340}
341
342define dso_local void @foo_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) {
343  ; CHECK-LABEL: .visible .func foo_St8x1(
344  ; CHECK:               .param .align 8 .b8 foo_St8x1_param_0[8],
345  ; CHECK:               .param .b32 foo_St8x1_param_1
346  ; CHECK:       )
347  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x1_param_1];
348  ; CHECK:       ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x1_param_0];
349  ; CHECK:       st.u64 [[[R1]]], [[RD1]];
350  ; CHECK:       ret;
351  %1 = load i64, ptr %in, align 8
352  store i64 %1, ptr %ret, align 8
353  ret void
354}
355
356define dso_local void @foo_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) {
357  ; CHECK-LABEL: .visible .func foo_St8x2(
358  ; CHECK:               .param .align 8 .b8 foo_St8x2_param_0[16],
359  ; CHECK:               .param .b32 foo_St8x2_param_1
360  ; CHECK:       )
361  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x2_param_1];
362  ; CHECK:       ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x2_param_0];
363  ; CHECK:       st.u64 [[[R1]]], [[RD1]];
364  ; CHECK:       ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x2_param_0+8];
365  ; CHECK:       st.u64 [[[R1]]+8], [[RD2]];
366  ; CHECK:       ret;
367  %1 = load i64, ptr %in, align 8
368  store i64 %1, ptr %ret, align 8
369  %arrayidx.1 = getelementptr inbounds [2 x i64], ptr %in, i64 0, i64 1
370  %2 = load i64, ptr %arrayidx.1, align 8
371  %arrayidx3.1 = getelementptr inbounds [2 x i64], ptr %ret, i64 0, i64 1
372  store i64 %2, ptr %arrayidx3.1, align 8
373  ret void
374}
375
376define dso_local void @foo_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) {
377  ; CHECK-LABEL: .visible .func foo_St8x3(
378  ; CHECK:               .param .align 8 .b8 foo_St8x3_param_0[24],
379  ; CHECK:               .param .b32 foo_St8x3_param_1
380  ; CHECK:       )
381  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x3_param_1];
382  ; CHECK:       ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x3_param_0];
383  ; CHECK:       st.u64 [[[R1]]], [[RD1]];
384  ; CHECK:       ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x3_param_0+8];
385  ; CHECK:       st.u64 [[[R1]]+8], [[RD2]];
386  ; CHECK:       ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x3_param_0+16];
387  ; CHECK:       st.u64 [[[R1]]+16], [[RD3]];
388  ; CHECK:       ret;
389  %1 = load i64, ptr %in, align 8
390  store i64 %1, ptr %ret, align 8
391  %arrayidx.1 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 1
392  %2 = load i64, ptr %arrayidx.1, align 8
393  %arrayidx3.1 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 1
394  store i64 %2, ptr %arrayidx3.1, align 8
395  %arrayidx.2 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 2
396  %3 = load i64, ptr %arrayidx.2, align 8
397  %arrayidx3.2 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 2
398  store i64 %3, ptr %arrayidx3.2, align 8
399  ret void
400}
401
402define dso_local void @foo_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {
403  ; CHECK-LABEL: .visible .func foo_St8x4(
404  ; CHECK:               .param .align 8 .b8 foo_St8x4_param_0[32],
405  ; CHECK:               .param .b32 foo_St8x4_param_1
406  ; CHECK:       )
407  ; CHECK:       ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x4_param_1];
408  ; CHECK:       ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x4_param_0];
409  ; CHECK:       st.u64 [[[R1]]], [[RD1]];
410  ; CHECK:       ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x4_param_0+8];
411  ; CHECK:       st.u64 [[[R1]]+8], [[RD2]];
412  ; CHECK:       ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x4_param_0+16];
413  ; CHECK:       st.u64 [[[R1]]+16], [[RD3]];
414  ; CHECK:       ld.param.u64 [[RD4:%rd[0-9]+]], [foo_St8x4_param_0+24];
415  ; CHECK:       st.u64 [[[R1]]+24], [[RD4]];
416  ; CHECK:       ret;
417  %1 = load i64, ptr %in, align 8
418  store i64 %1, ptr %ret, align 8
419  %arrayidx.1 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 1
420  %2 = load i64, ptr %arrayidx.1, align 8
421  %arrayidx3.1 = getelementptr inbounds [4 x i64], ptr %ret, i64 0, i64 1
422  store i64 %2, ptr %arrayidx3.1, align 8
423  %arrayidx.2 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 2
424  %3 = load i64, ptr %arrayidx.2, align 8
425  %arrayidx3.2 = getelementptr inbounds [4 x i64], ptr %ret, i64 0, i64 2
426  store i64 %3, ptr %arrayidx3.2, align 8
427  %arrayidx.3 = getelementptr inbounds [4 x i64], ptr %in, i64 0, i64 3
428  %4 = load i64, ptr %arrayidx.3, align 8
429  %arrayidx3.3 = getelementptr inbounds [4 x i64], ptr %ret, i64 0, i64 3
430  store i64 %4, ptr %arrayidx3.3, align 8
431  ret void
432}
433