1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2 ; ## Support i16x2 instructions
3 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 \
4 ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
5 ; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,I16x2 %s
7 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 \
8 ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
9 ; RUN: | %ptxas-verify -arch=sm_90 \
11 ; ## No support for i16x2 instructions
12 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \
13 ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
14 ; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,NO-I16x2 %s
16 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \
17 ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
18 ; RUN: | %ptxas-verify -arch=sm_53 \
21 target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
23 define <2 x i16> @test_ret_const() #0 {
24 ; COMMON-LABEL: test_ret_const(
26 ; COMMON-NEXT: .reg .b32 %r<2>;
28 ; COMMON-NEXT: // %bb.0:
29 ; COMMON-NEXT: mov.b32 %r1, 131073;
30 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
32 ret <2 x i16> <i16 1, i16 2>
35 define i16 @test_extract_0(<2 x i16> %a) #0 {
36 ; COMMON-LABEL: test_extract_0(
38 ; COMMON-NEXT: .reg .b16 %rs<2>;
39 ; COMMON-NEXT: .reg .b32 %r<3>;
41 ; COMMON-NEXT: // %bb.0:
42 ; COMMON-NEXT: ld.param.u32 %r1, [test_extract_0_param_0];
43 ; COMMON-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }
44 ; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
45 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
47 %e = extractelement <2 x i16> %a, i32 0
51 define i16 @test_extract_1(<2 x i16> %a) #0 {
52 ; COMMON-LABEL: test_extract_1(
54 ; COMMON-NEXT: .reg .b16 %rs<2>;
55 ; COMMON-NEXT: .reg .b32 %r<3>;
57 ; COMMON-NEXT: // %bb.0:
58 ; COMMON-NEXT: ld.param.u32 %r1, [test_extract_1_param_0];
59 ; COMMON-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }
60 ; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
61 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
63 %e = extractelement <2 x i16> %a, i32 1
67 define i16 @test_extract_i(<2 x i16> %a, i64 %idx) #0 {
68 ; COMMON-LABEL: test_extract_i(
70 ; COMMON-NEXT: .reg .pred %p<2>;
71 ; COMMON-NEXT: .reg .b16 %rs<4>;
72 ; COMMON-NEXT: .reg .b32 %r<3>;
73 ; COMMON-NEXT: .reg .b64 %rd<2>;
75 ; COMMON-NEXT: // %bb.0:
76 ; COMMON-NEXT: ld.param.u64 %rd1, [test_extract_i_param_1];
77 ; COMMON-NEXT: ld.param.u32 %r1, [test_extract_i_param_0];
78 ; COMMON-NEXT: setp.eq.s64 %p1, %rd1, 0;
79 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
80 ; COMMON-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
81 ; COMMON-NEXT: cvt.u32.u16 %r2, %rs3;
82 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
84 %e = extractelement <2 x i16> %a, i64 %idx
88 define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) #0 {
89 ; I16x2-LABEL: test_add(
91 ; I16x2-NEXT: .reg .b32 %r<4>;
93 ; I16x2-NEXT: // %bb.0:
94 ; I16x2-NEXT: ld.param.u32 %r2, [test_add_param_1];
95 ; I16x2-NEXT: ld.param.u32 %r1, [test_add_param_0];
96 ; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;
97 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
100 ; NO-I16x2-LABEL: test_add(
102 ; NO-I16x2-NEXT: .reg .b16 %rs<7>;
103 ; NO-I16x2-NEXT: .reg .b32 %r<4>;
105 ; NO-I16x2-NEXT: // %bb.0:
106 ; NO-I16x2-NEXT: ld.param.u32 %r2, [test_add_param_1];
107 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_add_param_0];
108 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
109 ; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
110 ; NO-I16x2-NEXT: add.s16 %rs5, %rs4, %rs2;
111 ; NO-I16x2-NEXT: add.s16 %rs6, %rs3, %rs1;
112 ; NO-I16x2-NEXT: mov.b32 %r3, {%rs6, %rs5};
113 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;
114 ; NO-I16x2-NEXT: ret;
115 %r = add <2 x i16> %a, %b
119 ; Check that we can lower add with immediate arguments.
120 define <2 x i16> @test_add_imm_0(<2 x i16> %a) #0 {
121 ; I16x2-LABEL: test_add_imm_0(
123 ; I16x2-NEXT: .reg .b32 %r<4>;
125 ; I16x2-NEXT: // %bb.0:
126 ; I16x2-NEXT: ld.param.u32 %r1, [test_add_imm_0_param_0];
127 ; I16x2-NEXT: mov.b32 %r2, 131073;
128 ; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;
129 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
132 ; NO-I16x2-LABEL: test_add_imm_0(
134 ; NO-I16x2-NEXT: .reg .b16 %rs<5>;
135 ; NO-I16x2-NEXT: .reg .b32 %r<3>;
137 ; NO-I16x2-NEXT: // %bb.0:
138 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_add_imm_0_param_0];
139 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;
140 ; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;
141 ; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;
142 ; NO-I16x2-NEXT: mov.b32 %r2, {%rs4, %rs3};
143 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;
144 ; NO-I16x2-NEXT: ret;
145 %r = add <2 x i16> <i16 1, i16 2>, %a
149 define <2 x i16> @test_add_imm_1(<2 x i16> %a) #0 {
150 ; I16x2-LABEL: test_add_imm_1(
152 ; I16x2-NEXT: .reg .b32 %r<4>;
154 ; I16x2-NEXT: // %bb.0:
155 ; I16x2-NEXT: ld.param.u32 %r1, [test_add_imm_1_param_0];
156 ; I16x2-NEXT: mov.b32 %r2, 131073;
157 ; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;
158 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
161 ; NO-I16x2-LABEL: test_add_imm_1(
163 ; NO-I16x2-NEXT: .reg .b16 %rs<5>;
164 ; NO-I16x2-NEXT: .reg .b32 %r<3>;
166 ; NO-I16x2-NEXT: // %bb.0:
167 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_add_imm_1_param_0];
168 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;
169 ; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;
170 ; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;
171 ; NO-I16x2-NEXT: mov.b32 %r2, {%rs4, %rs3};
172 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;
173 ; NO-I16x2-NEXT: ret;
174 %r = add <2 x i16> %a, <i16 1, i16 2>
178 define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) #0 {
179 ; COMMON-LABEL: test_sub(
181 ; COMMON-NEXT: .reg .b16 %rs<7>;
182 ; COMMON-NEXT: .reg .b32 %r<4>;
184 ; COMMON-NEXT: // %bb.0:
185 ; COMMON-NEXT: ld.param.u32 %r2, [test_sub_param_1];
186 ; COMMON-NEXT: ld.param.u32 %r1, [test_sub_param_0];
187 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
188 ; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
189 ; COMMON-NEXT: sub.s16 %rs5, %rs4, %rs2;
190 ; COMMON-NEXT: sub.s16 %rs6, %rs3, %rs1;
191 ; COMMON-NEXT: mov.b32 %r3, {%rs6, %rs5};
192 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
194 %r = sub <2 x i16> %a, %b
198 define <2 x i16> @test_smax(<2 x i16> %a, <2 x i16> %b) #0 {
199 ; I16x2-LABEL: test_smax(
201 ; I16x2-NEXT: .reg .b32 %r<4>;
203 ; I16x2-NEXT: // %bb.0:
204 ; I16x2-NEXT: ld.param.u32 %r2, [test_smax_param_1];
205 ; I16x2-NEXT: ld.param.u32 %r1, [test_smax_param_0];
206 ; I16x2-NEXT: max.s16x2 %r3, %r1, %r2;
207 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
210 ; NO-I16x2-LABEL: test_smax(
212 ; NO-I16x2-NEXT: .reg .b16 %rs<7>;
213 ; NO-I16x2-NEXT: .reg .b32 %r<4>;
215 ; NO-I16x2-NEXT: // %bb.0:
216 ; NO-I16x2-NEXT: ld.param.u32 %r2, [test_smax_param_1];
217 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_smax_param_0];
218 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
219 ; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
220 ; NO-I16x2-NEXT: max.s16 %rs5, %rs4, %rs2;
221 ; NO-I16x2-NEXT: max.s16 %rs6, %rs3, %rs1;
222 ; NO-I16x2-NEXT: mov.b32 %r3, {%rs6, %rs5};
223 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;
224 ; NO-I16x2-NEXT: ret;
225 %cmp = icmp sgt <2 x i16> %a, %b
226 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b
230 define <2 x i16> @test_umax(<2 x i16> %a, <2 x i16> %b) #0 {
231 ; I16x2-LABEL: test_umax(
233 ; I16x2-NEXT: .reg .b32 %r<4>;
235 ; I16x2-NEXT: // %bb.0:
236 ; I16x2-NEXT: ld.param.u32 %r2, [test_umax_param_1];
237 ; I16x2-NEXT: ld.param.u32 %r1, [test_umax_param_0];
238 ; I16x2-NEXT: max.u16x2 %r3, %r1, %r2;
239 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
242 ; NO-I16x2-LABEL: test_umax(
244 ; NO-I16x2-NEXT: .reg .b16 %rs<7>;
245 ; NO-I16x2-NEXT: .reg .b32 %r<4>;
247 ; NO-I16x2-NEXT: // %bb.0:
248 ; NO-I16x2-NEXT: ld.param.u32 %r2, [test_umax_param_1];
249 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_umax_param_0];
250 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
251 ; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
252 ; NO-I16x2-NEXT: max.u16 %rs5, %rs4, %rs2;
253 ; NO-I16x2-NEXT: max.u16 %rs6, %rs3, %rs1;
254 ; NO-I16x2-NEXT: mov.b32 %r3, {%rs6, %rs5};
255 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;
256 ; NO-I16x2-NEXT: ret;
257 %cmp = icmp ugt <2 x i16> %a, %b
258 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b
262 define <2 x i16> @test_smin(<2 x i16> %a, <2 x i16> %b) #0 {
263 ; I16x2-LABEL: test_smin(
265 ; I16x2-NEXT: .reg .b32 %r<4>;
267 ; I16x2-NEXT: // %bb.0:
268 ; I16x2-NEXT: ld.param.u32 %r2, [test_smin_param_1];
269 ; I16x2-NEXT: ld.param.u32 %r1, [test_smin_param_0];
270 ; I16x2-NEXT: min.s16x2 %r3, %r1, %r2;
271 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
274 ; NO-I16x2-LABEL: test_smin(
276 ; NO-I16x2-NEXT: .reg .b16 %rs<7>;
277 ; NO-I16x2-NEXT: .reg .b32 %r<4>;
279 ; NO-I16x2-NEXT: // %bb.0:
280 ; NO-I16x2-NEXT: ld.param.u32 %r2, [test_smin_param_1];
281 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_smin_param_0];
282 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
283 ; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
284 ; NO-I16x2-NEXT: min.s16 %rs5, %rs4, %rs2;
285 ; NO-I16x2-NEXT: min.s16 %rs6, %rs3, %rs1;
286 ; NO-I16x2-NEXT: mov.b32 %r3, {%rs6, %rs5};
287 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;
288 ; NO-I16x2-NEXT: ret;
289 %cmp = icmp sle <2 x i16> %a, %b
290 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b
294 define <2 x i16> @test_umin(<2 x i16> %a, <2 x i16> %b) #0 {
295 ; I16x2-LABEL: test_umin(
297 ; I16x2-NEXT: .reg .b32 %r<4>;
299 ; I16x2-NEXT: // %bb.0:
300 ; I16x2-NEXT: ld.param.u32 %r2, [test_umin_param_1];
301 ; I16x2-NEXT: ld.param.u32 %r1, [test_umin_param_0];
302 ; I16x2-NEXT: min.u16x2 %r3, %r1, %r2;
303 ; I16x2-NEXT: st.param.b32 [func_retval0], %r3;
306 ; NO-I16x2-LABEL: test_umin(
308 ; NO-I16x2-NEXT: .reg .b16 %rs<7>;
309 ; NO-I16x2-NEXT: .reg .b32 %r<4>;
311 ; NO-I16x2-NEXT: // %bb.0:
312 ; NO-I16x2-NEXT: ld.param.u32 %r2, [test_umin_param_1];
313 ; NO-I16x2-NEXT: ld.param.u32 %r1, [test_umin_param_0];
314 ; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
315 ; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
316 ; NO-I16x2-NEXT: min.u16 %rs5, %rs4, %rs2;
317 ; NO-I16x2-NEXT: min.u16 %rs6, %rs3, %rs1;
318 ; NO-I16x2-NEXT: mov.b32 %r3, {%rs6, %rs5};
319 ; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;
320 ; NO-I16x2-NEXT: ret;
321 %cmp = icmp ule <2 x i16> %a, %b
322 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b
326 define <2 x i16> @test_mul(<2 x i16> %a, <2 x i16> %b) #0 {
327 ; COMMON-LABEL: test_mul(
329 ; COMMON-NEXT: .reg .b16 %rs<7>;
330 ; COMMON-NEXT: .reg .b32 %r<4>;
332 ; COMMON-NEXT: // %bb.0:
333 ; COMMON-NEXT: ld.param.u32 %r2, [test_mul_param_1];
334 ; COMMON-NEXT: ld.param.u32 %r1, [test_mul_param_0];
335 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
336 ; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
337 ; COMMON-NEXT: mul.lo.s16 %rs5, %rs4, %rs2;
338 ; COMMON-NEXT: mul.lo.s16 %rs6, %rs3, %rs1;
339 ; COMMON-NEXT: mov.b32 %r3, {%rs6, %rs5};
340 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
342 %r = mul <2 x i16> %a, %b
346 ;; Logical ops are available on all GPUs as regular 32-bit logical ops
347 define <2 x i16> @test_or(<2 x i16> %a, <2 x i16> %b) #0 {
348 ; COMMON-LABEL: test_or(
350 ; COMMON-NEXT: .reg .b32 %r<7>;
352 ; COMMON-NEXT: // %bb.0:
353 ; COMMON-NEXT: ld.param.u32 %r3, [test_or_param_1];
354 ; COMMON-NEXT: ld.param.u32 %r4, [test_or_param_0];
355 ; COMMON-NEXT: or.b32 %r5, %r4, %r3;
356 ; COMMON-NEXT: st.param.b32 [func_retval0], %r5;
358 %r = or <2 x i16> %a, %b
362 ; Ops that operate on computed arguments go though a different lowering path.
363 ; compared to the ones that operate on loaded data. So we test them separately.
364 define <2 x i16> @test_or_computed(i16 %a) {
365 ; COMMON-LABEL: test_or_computed(
367 ; COMMON-NEXT: .reg .b16 %rs<4>;
368 ; COMMON-NEXT: .reg .b32 %r<4>;
370 ; COMMON-NEXT: // %bb.0:
371 ; COMMON-NEXT: ld.param.u16 %rs1, [test_or_computed_param_0];
372 ; COMMON-NEXT: mov.u16 %rs2, 0;
373 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
374 ; COMMON-NEXT: mov.u16 %rs3, 5;
375 ; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};
376 ; COMMON-NEXT: or.b32 %r3, %r2, %r1;
377 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
379 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0
380 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1
381 %r = or <2 x i16> %ins.1, %ins.0
385 ; Check that we can lower or with immediate arguments.
386 define <2 x i16> @test_or_imm_0(<2 x i16> %a) #0 {
387 ; COMMON-LABEL: test_or_imm_0(
389 ; COMMON-NEXT: .reg .b32 %r<3>;
391 ; COMMON-NEXT: // %bb.0:
392 ; COMMON-NEXT: ld.param.u32 %r1, [test_or_imm_0_param_0];
393 ; COMMON-NEXT: or.b32 %r2, %r1, 131073;
394 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
396 %r = or <2 x i16> <i16 1, i16 2>, %a
400 define <2 x i16> @test_or_imm_1(<2 x i16> %a) #0 {
401 ; COMMON-LABEL: test_or_imm_1(
403 ; COMMON-NEXT: .reg .b32 %r<3>;
405 ; COMMON-NEXT: // %bb.0:
406 ; COMMON-NEXT: ld.param.u32 %r1, [test_or_imm_1_param_0];
407 ; COMMON-NEXT: or.b32 %r2, %r1, 131073;
408 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
410 %r = or <2 x i16> %a, <i16 1, i16 2>
414 define <2 x i16> @test_xor(<2 x i16> %a, <2 x i16> %b) #0 {
415 ; COMMON-LABEL: test_xor(
417 ; COMMON-NEXT: .reg .b32 %r<7>;
419 ; COMMON-NEXT: // %bb.0:
420 ; COMMON-NEXT: ld.param.u32 %r3, [test_xor_param_1];
421 ; COMMON-NEXT: ld.param.u32 %r4, [test_xor_param_0];
422 ; COMMON-NEXT: xor.b32 %r5, %r4, %r3;
423 ; COMMON-NEXT: st.param.b32 [func_retval0], %r5;
425 %r = xor <2 x i16> %a, %b
429 define <2 x i16> @test_xor_computed(i16 %a) {
430 ; COMMON-LABEL: test_xor_computed(
432 ; COMMON-NEXT: .reg .b16 %rs<4>;
433 ; COMMON-NEXT: .reg .b32 %r<4>;
435 ; COMMON-NEXT: // %bb.0:
436 ; COMMON-NEXT: ld.param.u16 %rs1, [test_xor_computed_param_0];
437 ; COMMON-NEXT: mov.u16 %rs2, 0;
438 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
439 ; COMMON-NEXT: mov.u16 %rs3, 5;
440 ; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};
441 ; COMMON-NEXT: xor.b32 %r3, %r2, %r1;
442 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
444 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0
445 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1
446 %r = xor <2 x i16> %ins.1, %ins.0
450 ; Check that we can lower xor with immediate arguments.
451 define <2 x i16> @test_xor_imm_0(<2 x i16> %a) #0 {
452 ; COMMON-LABEL: test_xor_imm_0(
454 ; COMMON-NEXT: .reg .b32 %r<3>;
456 ; COMMON-NEXT: // %bb.0:
457 ; COMMON-NEXT: ld.param.u32 %r1, [test_xor_imm_0_param_0];
458 ; COMMON-NEXT: xor.b32 %r2, %r1, 131073;
459 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
461 %r = xor <2 x i16> <i16 1, i16 2>, %a
465 define <2 x i16> @test_xor_imm_1(<2 x i16> %a) #0 {
466 ; COMMON-LABEL: test_xor_imm_1(
468 ; COMMON-NEXT: .reg .b32 %r<3>;
470 ; COMMON-NEXT: // %bb.0:
471 ; COMMON-NEXT: ld.param.u32 %r1, [test_xor_imm_1_param_0];
472 ; COMMON-NEXT: xor.b32 %r2, %r1, 131073;
473 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
475 %r = xor <2 x i16> %a, <i16 1, i16 2>
479 define <2 x i16> @test_and(<2 x i16> %a, <2 x i16> %b) #0 {
480 ; COMMON-LABEL: test_and(
482 ; COMMON-NEXT: .reg .b32 %r<7>;
484 ; COMMON-NEXT: // %bb.0:
485 ; COMMON-NEXT: ld.param.u32 %r3, [test_and_param_1];
486 ; COMMON-NEXT: ld.param.u32 %r4, [test_and_param_0];
487 ; COMMON-NEXT: and.b32 %r5, %r4, %r3;
488 ; COMMON-NEXT: st.param.b32 [func_retval0], %r5;
490 %r = and <2 x i16> %a, %b
494 ; Ops that operate on computed arguments go though a different lowering path.
495 ; compared to the ones that operate on loaded data. So we test them separately.
496 define <2 x i16> @test_and_computed(i16 %a) {
497 ; COMMON-LABEL: test_and_computed(
499 ; COMMON-NEXT: .reg .b16 %rs<4>;
500 ; COMMON-NEXT: .reg .b32 %r<4>;
502 ; COMMON-NEXT: // %bb.0:
503 ; COMMON-NEXT: ld.param.u16 %rs1, [test_and_computed_param_0];
504 ; COMMON-NEXT: mov.u16 %rs2, 0;
505 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
506 ; COMMON-NEXT: mov.u16 %rs3, 5;
507 ; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};
508 ; COMMON-NEXT: and.b32 %r3, %r2, %r1;
509 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
511 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0
512 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1
513 %r = and <2 x i16> %ins.1, %ins.0
517 ; Check that we can lower and with immediate arguments.
518 define <2 x i16> @test_and_imm_0(<2 x i16> %a) #0 {
519 ; COMMON-LABEL: test_and_imm_0(
521 ; COMMON-NEXT: .reg .b32 %r<3>;
523 ; COMMON-NEXT: // %bb.0:
524 ; COMMON-NEXT: ld.param.u32 %r1, [test_and_imm_0_param_0];
525 ; COMMON-NEXT: and.b32 %r2, %r1, 131073;
526 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
528 %r = and <2 x i16> <i16 1, i16 2>, %a
532 define <2 x i16> @test_and_imm_1(<2 x i16> %a) #0 {
533 ; COMMON-LABEL: test_and_imm_1(
535 ; COMMON-NEXT: .reg .b32 %r<3>;
537 ; COMMON-NEXT: // %bb.0:
538 ; COMMON-NEXT: ld.param.u32 %r1, [test_and_imm_1_param_0];
539 ; COMMON-NEXT: and.b32 %r2, %r1, 131073;
540 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
542 %r = and <2 x i16> %a, <i16 1, i16 2>
546 define void @test_ldst_v2i16(ptr %a, ptr %b) {
547 ; COMMON-LABEL: test_ldst_v2i16(
549 ; COMMON-NEXT: .reg .b32 %r<2>;
550 ; COMMON-NEXT: .reg .b64 %rd<3>;
552 ; COMMON-NEXT: // %bb.0:
553 ; COMMON-NEXT: ld.param.u64 %rd2, [test_ldst_v2i16_param_1];
554 ; COMMON-NEXT: ld.param.u64 %rd1, [test_ldst_v2i16_param_0];
555 ; COMMON-NEXT: ld.u32 %r1, [%rd1];
556 ; COMMON-NEXT: st.u32 [%rd2], %r1;
558 %t1 = load <2 x i16>, ptr %a
559 store <2 x i16> %t1, ptr %b, align 16
563 ; -- v3 is inconvenient to capture as it's lowered as ld.b64 + fair
564 ; number of bitshifting instructions that may change at llvm's whim.
565 ; So we only verify that we only issue correct number of writes using
566 ; correct offset, but not the values we write.
567 define void @test_ldst_v3i16(ptr %a, ptr %b) {
568 ; COMMON-LABEL: test_ldst_v3i16(
570 ; COMMON-NEXT: .reg .b64 %rd<5>;
572 ; COMMON-NEXT: // %bb.0:
573 ; COMMON-NEXT: ld.param.u64 %rd2, [test_ldst_v3i16_param_1];
574 ; COMMON-NEXT: ld.param.u64 %rd1, [test_ldst_v3i16_param_0];
575 ; COMMON-NEXT: ld.u64 %rd3, [%rd1];
576 ; COMMON-NEXT: shr.u64 %rd4, %rd3, 32;
577 ; COMMON-NEXT: st.u32 [%rd2], %rd3;
578 ; COMMON-NEXT: st.u16 [%rd2+4], %rd4;
580 %t1 = load <3 x i16>, ptr %a
581 store <3 x i16> %t1, ptr %b, align 16
585 define void @test_ldst_v4i16(ptr %a, ptr %b) {
586 ; COMMON-LABEL: test_ldst_v4i16(
588 ; COMMON-NEXT: .reg .b16 %rs<5>;
589 ; COMMON-NEXT: .reg .b64 %rd<3>;
591 ; COMMON-NEXT: // %bb.0:
592 ; COMMON-NEXT: ld.param.u64 %rd2, [test_ldst_v4i16_param_1];
593 ; COMMON-NEXT: ld.param.u64 %rd1, [test_ldst_v4i16_param_0];
594 ; COMMON-NEXT: ld.v4.u16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];
595 ; COMMON-NEXT: st.v4.u16 [%rd2], {%rs1, %rs2, %rs3, %rs4};
597 %t1 = load <4 x i16>, ptr %a
598 store <4 x i16> %t1, ptr %b, align 16
602 define void @test_ldst_v8i16(ptr %a, ptr %b) {
603 ; COMMON-LABEL: test_ldst_v8i16(
605 ; COMMON-NEXT: .reg .b32 %r<5>;
606 ; COMMON-NEXT: .reg .b64 %rd<3>;
608 ; COMMON-NEXT: // %bb.0:
609 ; COMMON-NEXT: ld.param.u64 %rd2, [test_ldst_v8i16_param_1];
610 ; COMMON-NEXT: ld.param.u64 %rd1, [test_ldst_v8i16_param_0];
611 ; COMMON-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
612 ; COMMON-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};
614 %t1 = load <8 x i16>, ptr %a
615 store <8 x i16> %t1, ptr %b, align 16
619 declare <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b) #0
621 define <2 x i16> @test_call(<2 x i16> %a, <2 x i16> %b) #0 {
622 ; COMMON-LABEL: test_call(
624 ; COMMON-NEXT: .reg .b32 %r<5>;
626 ; COMMON-NEXT: // %bb.0:
627 ; COMMON-NEXT: ld.param.u32 %r2, [test_call_param_1];
628 ; COMMON-NEXT: ld.param.u32 %r1, [test_call_param_0];
629 ; COMMON-NEXT: { // callseq 0, 0
630 ; COMMON-NEXT: .param .align 4 .b8 param0[4];
631 ; COMMON-NEXT: st.param.b32 [param0], %r1;
632 ; COMMON-NEXT: .param .align 4 .b8 param1[4];
633 ; COMMON-NEXT: st.param.b32 [param1], %r2;
634 ; COMMON-NEXT: .param .align 4 .b8 retval0[4];
635 ; COMMON-NEXT: call.uni (retval0),
636 ; COMMON-NEXT: test_callee,
638 ; COMMON-NEXT: param0,
639 ; COMMON-NEXT: param1
641 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
642 ; COMMON-NEXT: } // callseq 0
643 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
645 %r = call <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b)
649 define <2 x i16> @test_call_flipped(<2 x i16> %a, <2 x i16> %b) #0 {
650 ; COMMON-LABEL: test_call_flipped(
652 ; COMMON-NEXT: .reg .b32 %r<5>;
654 ; COMMON-NEXT: // %bb.0:
655 ; COMMON-NEXT: ld.param.u32 %r2, [test_call_flipped_param_1];
656 ; COMMON-NEXT: ld.param.u32 %r1, [test_call_flipped_param_0];
657 ; COMMON-NEXT: { // callseq 1, 0
658 ; COMMON-NEXT: .param .align 4 .b8 param0[4];
659 ; COMMON-NEXT: st.param.b32 [param0], %r2;
660 ; COMMON-NEXT: .param .align 4 .b8 param1[4];
661 ; COMMON-NEXT: st.param.b32 [param1], %r1;
662 ; COMMON-NEXT: .param .align 4 .b8 retval0[4];
663 ; COMMON-NEXT: call.uni (retval0),
664 ; COMMON-NEXT: test_callee,
666 ; COMMON-NEXT: param0,
667 ; COMMON-NEXT: param1
669 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
670 ; COMMON-NEXT: } // callseq 1
671 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
673 %r = call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a)
677 define <2 x i16> @test_tailcall_flipped(<2 x i16> %a, <2 x i16> %b) #0 {
678 ; COMMON-LABEL: test_tailcall_flipped(
680 ; COMMON-NEXT: .reg .b32 %r<5>;
682 ; COMMON-NEXT: // %bb.0:
683 ; COMMON-NEXT: ld.param.u32 %r2, [test_tailcall_flipped_param_1];
684 ; COMMON-NEXT: ld.param.u32 %r1, [test_tailcall_flipped_param_0];
685 ; COMMON-NEXT: { // callseq 2, 0
686 ; COMMON-NEXT: .param .align 4 .b8 param0[4];
687 ; COMMON-NEXT: st.param.b32 [param0], %r2;
688 ; COMMON-NEXT: .param .align 4 .b8 param1[4];
689 ; COMMON-NEXT: st.param.b32 [param1], %r1;
690 ; COMMON-NEXT: .param .align 4 .b8 retval0[4];
691 ; COMMON-NEXT: call.uni (retval0),
692 ; COMMON-NEXT: test_callee,
694 ; COMMON-NEXT: param0,
695 ; COMMON-NEXT: param1
697 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
698 ; COMMON-NEXT: } // callseq 2
699 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
701 %r = tail call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a)
705 define <2 x i16> @test_select(<2 x i16> %a, <2 x i16> %b, i1 zeroext %c) #0 {
706 ; COMMON-LABEL: test_select(
708 ; COMMON-NEXT: .reg .pred %p<2>;
709 ; COMMON-NEXT: .reg .b16 %rs<3>;
710 ; COMMON-NEXT: .reg .b32 %r<4>;
712 ; COMMON-NEXT: // %bb.0:
713 ; COMMON-NEXT: ld.param.u8 %rs1, [test_select_param_2];
714 ; COMMON-NEXT: and.b16 %rs2, %rs1, 1;
715 ; COMMON-NEXT: setp.eq.b16 %p1, %rs2, 1;
716 ; COMMON-NEXT: ld.param.u32 %r2, [test_select_param_1];
717 ; COMMON-NEXT: ld.param.u32 %r1, [test_select_param_0];
718 ; COMMON-NEXT: selp.b32 %r3, %r1, %r2, %p1;
719 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
721 %r = select i1 %c, <2 x i16> %a, <2 x i16> %b
725 define <2 x i16> @test_select_cc(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> %d) #0 {
726 ; COMMON-LABEL: test_select_cc(
728 ; COMMON-NEXT: .reg .pred %p<3>;
729 ; COMMON-NEXT: .reg .b16 %rs<11>;
730 ; COMMON-NEXT: .reg .b32 %r<6>;
732 ; COMMON-NEXT: // %bb.0:
733 ; COMMON-NEXT: ld.param.u32 %r4, [test_select_cc_param_3];
734 ; COMMON-NEXT: ld.param.u32 %r3, [test_select_cc_param_2];
735 ; COMMON-NEXT: ld.param.u32 %r2, [test_select_cc_param_1];
736 ; COMMON-NEXT: ld.param.u32 %r1, [test_select_cc_param_0];
737 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r4;
738 ; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r3;
739 ; COMMON-NEXT: setp.ne.s16 %p1, %rs3, %rs1;
740 ; COMMON-NEXT: setp.ne.s16 %p2, %rs4, %rs2;
741 ; COMMON-NEXT: mov.b32 {%rs5, %rs6}, %r2;
742 ; COMMON-NEXT: mov.b32 {%rs7, %rs8}, %r1;
743 ; COMMON-NEXT: selp.b16 %rs9, %rs8, %rs6, %p2;
744 ; COMMON-NEXT: selp.b16 %rs10, %rs7, %rs5, %p1;
745 ; COMMON-NEXT: mov.b32 %r5, {%rs10, %rs9};
746 ; COMMON-NEXT: st.param.b32 [func_retval0], %r5;
748 %cc = icmp ne <2 x i16> %c, %d
749 %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b
753 define <2 x i32> @test_select_cc_i32_i16(<2 x i32> %a, <2 x i32> %b,
754 ; COMMON-LABEL: test_select_cc_i32_i16(
756 ; COMMON-NEXT: .reg .pred %p<3>;
757 ; COMMON-NEXT: .reg .b16 %rs<5>;
758 ; COMMON-NEXT: .reg .b32 %r<9>;
760 ; COMMON-NEXT: // %bb.0:
761 ; COMMON-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_select_cc_i32_i16_param_1];
762 ; COMMON-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_select_cc_i32_i16_param_0];
763 ; COMMON-NEXT: ld.param.u32 %r6, [test_select_cc_i32_i16_param_3];
764 ; COMMON-NEXT: ld.param.u32 %r5, [test_select_cc_i32_i16_param_2];
765 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r6;
766 ; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r5;
767 ; COMMON-NEXT: setp.ne.s16 %p1, %rs3, %rs1;
768 ; COMMON-NEXT: setp.ne.s16 %p2, %rs4, %rs2;
769 ; COMMON-NEXT: selp.b32 %r7, %r2, %r4, %p2;
770 ; COMMON-NEXT: selp.b32 %r8, %r1, %r3, %p1;
771 ; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
773 <2 x i16> %c, <2 x i16> %d) #0 {
774 %cc = icmp ne <2 x i16> %c, %d
775 %r = select <2 x i1> %cc, <2 x i32> %a, <2 x i32> %b
779 define <2 x i16> @test_select_cc_i16_i32(<2 x i16> %a, <2 x i16> %b,
780 ; COMMON-LABEL: test_select_cc_i16_i32(
782 ; COMMON-NEXT: .reg .pred %p<3>;
783 ; COMMON-NEXT: .reg .b16 %rs<7>;
784 ; COMMON-NEXT: .reg .b32 %r<8>;
786 ; COMMON-NEXT: // %bb.0:
787 ; COMMON-NEXT: ld.param.v2.u32 {%r5, %r6}, [test_select_cc_i16_i32_param_3];
788 ; COMMON-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_select_cc_i16_i32_param_2];
789 ; COMMON-NEXT: ld.param.u32 %r2, [test_select_cc_i16_i32_param_1];
790 ; COMMON-NEXT: ld.param.u32 %r1, [test_select_cc_i16_i32_param_0];
791 ; COMMON-NEXT: setp.ne.s32 %p1, %r3, %r5;
792 ; COMMON-NEXT: setp.ne.s32 %p2, %r4, %r6;
793 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
794 ; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
795 ; COMMON-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;
796 ; COMMON-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;
797 ; COMMON-NEXT: mov.b32 %r7, {%rs6, %rs5};
798 ; COMMON-NEXT: st.param.b32 [func_retval0], %r7;
800 <2 x i32> %c, <2 x i32> %d) #0 {
801 %cc = icmp ne <2 x i32> %c, %d
802 %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b
807 define <2 x i16> @test_trunc_2xi32(<2 x i32> %a) #0 {
808 ; COMMON-LABEL: test_trunc_2xi32(
810 ; COMMON-NEXT: .reg .b16 %rs<3>;
811 ; COMMON-NEXT: .reg .b32 %r<4>;
813 ; COMMON-NEXT: // %bb.0:
814 ; COMMON-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_param_0];
815 ; COMMON-NEXT: cvt.u16.u32 %rs1, %r2;
816 ; COMMON-NEXT: cvt.u16.u32 %rs2, %r1;
817 ; COMMON-NEXT: mov.b32 %r3, {%rs2, %rs1};
818 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
820 %r = trunc <2 x i32> %a to <2 x i16>
824 define <2 x i16> @test_trunc_2xi64(<2 x i64> %a) #0 {
825 ; COMMON-LABEL: test_trunc_2xi64(
827 ; COMMON-NEXT: .reg .b16 %rs<3>;
828 ; COMMON-NEXT: .reg .b32 %r<2>;
829 ; COMMON-NEXT: .reg .b64 %rd<3>;
831 ; COMMON-NEXT: // %bb.0:
832 ; COMMON-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
833 ; COMMON-NEXT: cvt.u16.u64 %rs1, %rd2;
834 ; COMMON-NEXT: cvt.u16.u64 %rs2, %rd1;
835 ; COMMON-NEXT: mov.b32 %r1, {%rs2, %rs1};
836 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
838 %r = trunc <2 x i64> %a to <2 x i16>
842 define <2 x i32> @test_zext_2xi32(<2 x i16> %a) #0 {
843 ; COMMON-LABEL: test_zext_2xi32(
845 ; COMMON-NEXT: .reg .b16 %rs<3>;
846 ; COMMON-NEXT: .reg .b32 %r<4>;
848 ; COMMON-NEXT: // %bb.0:
849 ; COMMON-NEXT: ld.param.u32 %r1, [test_zext_2xi32_param_0];
850 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
851 ; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
852 ; COMMON-NEXT: cvt.u32.u16 %r3, %rs2;
853 ; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r3};
855 %r = zext <2 x i16> %a to <2 x i32>
859 define <2 x i64> @test_zext_2xi64(<2 x i16> %a) #0 {
860 ; COMMON-LABEL: test_zext_2xi64(
862 ; COMMON-NEXT: .reg .b16 %rs<3>;
863 ; COMMON-NEXT: .reg .b32 %r<2>;
864 ; COMMON-NEXT: .reg .b64 %rd<3>;
866 ; COMMON-NEXT: // %bb.0:
867 ; COMMON-NEXT: ld.param.u32 %r1, [test_zext_2xi64_param_0];
868 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
869 ; COMMON-NEXT: cvt.u64.u16 %rd1, %rs2;
870 ; COMMON-NEXT: cvt.u64.u16 %rd2, %rs1;
871 ; COMMON-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
873 %r = zext <2 x i16> %a to <2 x i64>
877 define <2 x i16> @test_bitcast_i32_to_2xi16(i32 %a) #0 {
878 ; COMMON-LABEL: test_bitcast_i32_to_2xi16(
880 ; COMMON-NEXT: .reg .b32 %r<3>;
882 ; COMMON-NEXT: // %bb.0:
883 ; COMMON-NEXT: ld.param.u32 %r1, [test_bitcast_i32_to_2xi16_param_0];
884 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
886 %r = bitcast i32 %a to <2 x i16>
890 define i32 @test_bitcast_2xi16_to_i32(<2 x i16> %a) #0 {
891 ; COMMON-LABEL: test_bitcast_2xi16_to_i32(
893 ; COMMON-NEXT: .reg .b32 %r<3>;
895 ; COMMON-NEXT: // %bb.0:
896 ; COMMON-NEXT: ld.param.u32 %r2, [test_bitcast_2xi16_to_i32_param_0];
897 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
899 %r = bitcast <2 x i16> %a to i32
903 define <2 x half> @test_bitcast_2xi16_to_2xhalf(i16 %a) #0 {
904 ; COMMON-LABEL: test_bitcast_2xi16_to_2xhalf(
906 ; COMMON-NEXT: .reg .b16 %rs<3>;
907 ; COMMON-NEXT: .reg .b32 %r<3>;
909 ; COMMON-NEXT: // %bb.0:
910 ; COMMON-NEXT: ld.param.u16 %rs1, [test_bitcast_2xi16_to_2xhalf_param_0];
911 ; COMMON-NEXT: mov.u16 %rs2, 5;
912 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
913 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
915 %ins.0 = insertelement <2 x i16> undef, i16 %a, i32 0
916 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1
917 %r = bitcast <2 x i16> %ins.1 to <2 x half>
922 define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 {
923 ; COMMON-LABEL: test_shufflevector(
925 ; COMMON-NEXT: .reg .b16 %rs<3>;
926 ; COMMON-NEXT: .reg .b32 %r<3>;
928 ; COMMON-NEXT: // %bb.0:
929 ; COMMON-NEXT: ld.param.u32 %r1, [test_shufflevector_param_0];
930 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
931 ; COMMON-NEXT: mov.b32 %r2, {%rs2, %rs1};
932 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
934 %s = shufflevector <2 x i16> %a, <2 x i16> undef, <2 x i32> <i32 1, i32 0>
938 define <2 x i16> @test_insertelement(<2 x i16> %a, i16 %x) #0 {
939 ; COMMON-LABEL: test_insertelement(
941 ; COMMON-NEXT: .reg .b16 %rs<3>;
942 ; COMMON-NEXT: .reg .b32 %r<3>;
944 ; COMMON-NEXT: // %bb.0:
945 ; COMMON-NEXT: ld.param.u16 %rs1, [test_insertelement_param_1];
946 ; COMMON-NEXT: ld.param.u32 %r1, [test_insertelement_param_0];
947 ; COMMON-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }
948 ; COMMON-NEXT: mov.b32 %r2, {%rs2, %rs1};
949 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
951 %i = insertelement <2 x i16> %a, i16 %x, i64 1
955 define <2 x i16> @test_fptosi_2xhalf_to_2xi16(<2 x half> %a) #0 {
956 ; COMMON-LABEL: test_fptosi_2xhalf_to_2xi16(
958 ; COMMON-NEXT: .reg .b16 %rs<5>;
959 ; COMMON-NEXT: .reg .b32 %r<3>;
961 ; COMMON-NEXT: // %bb.0:
962 ; COMMON-NEXT: ld.param.b32 %r1, [test_fptosi_2xhalf_to_2xi16_param_0];
963 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
964 ; COMMON-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;
965 ; COMMON-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;
966 ; COMMON-NEXT: mov.b32 %r2, {%rs4, %rs3};
967 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
969 %r = fptosi <2 x half> %a to <2 x i16>
973 define <2 x i16> @test_fptoui_2xhalf_to_2xi16(<2 x half> %a) #0 {
974 ; COMMON-LABEL: test_fptoui_2xhalf_to_2xi16(
976 ; COMMON-NEXT: .reg .b16 %rs<5>;
977 ; COMMON-NEXT: .reg .b32 %r<3>;
979 ; COMMON-NEXT: // %bb.0:
980 ; COMMON-NEXT: ld.param.b32 %r1, [test_fptoui_2xhalf_to_2xi16_param_0];
981 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
982 ; COMMON-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;
983 ; COMMON-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;
984 ; COMMON-NEXT: mov.b32 %r2, {%rs4, %rs3};
985 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
987 %r = fptoui <2 x half> %a to <2 x i16>
991 attributes #0 = { nounwind }