Lines Matching full:common
5 ; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,I16x2 %s
14 ; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,NO-I16x2 %s
24 ; COMMON-LABEL: test_ret_const(
25 ; COMMON: {
26 ; COMMON-NEXT: .reg .b32 %r<2>;
27 ; COMMON-EMPTY:
28 ; COMMON-NEXT: // %bb.0:
29 ; COMMON-NEXT: mov.b32 %r1, 131073;
30 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
31 ; COMMON-NEXT: ret;
36 ; COMMON-LABEL: test_extract_0(
37 ; COMMON: {
38 ; COMMON-NEXT: .reg .b16 %rs<2>;
39 ; COMMON-NEXT: .reg .b32 %r<3>;
40 ; COMMON-EMPTY:
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;
46 ; COMMON-NEXT: ret;
52 ; COMMON-LABEL: test_extract_1(
53 ; COMMON: {
54 ; COMMON-NEXT: .reg .b16 %rs<2>;
55 ; COMMON-NEXT: .reg .b32 %r<3>;
56 ; COMMON-EMPTY:
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;
62 ; COMMON-NEXT: ret;
68 ; COMMON-LABEL: test_extract_i(
69 ; COMMON: {
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>;
74 ; COMMON-EMPTY:
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;
83 ; COMMON-NEXT: ret;
179 ; COMMON-LABEL: test_sub(
180 ; COMMON: {
181 ; COMMON-NEXT: .reg .b16 %rs<7>;
182 ; COMMON-NEXT: .reg .b32 %r<4>;
183 ; COMMON-EMPTY:
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;
193 ; COMMON-NEXT: ret;
327 ; COMMON-LABEL: test_mul(
328 ; COMMON: {
329 ; COMMON-NEXT: .reg .b16 %rs<7>;
330 ; COMMON-NEXT: .reg .b32 %r<4>;
331 ; COMMON-EMPTY:
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;
341 ; COMMON-NEXT: ret;
348 ; COMMON-LABEL: test_or(
349 ; COMMON: {
350 ; COMMON-NEXT: .reg .b32 %r<4>;
351 ; COMMON-EMPTY:
352 ; COMMON-NEXT: // %bb.0:
353 ; COMMON-NEXT: ld.param.u32 %r2, [test_or_param_1];
354 ; COMMON-NEXT: ld.param.u32 %r1, [test_or_param_0];
355 ; COMMON-NEXT: or.b32 %r3, %r1, %r2;
356 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
357 ; COMMON-NEXT: ret;
365 ; COMMON-LABEL: test_or_computed(
366 ; COMMON: {
367 ; COMMON-NEXT: .reg .b16 %rs<4>;
368 ; COMMON-NEXT: .reg .b32 %r<4>;
369 ; COMMON-EMPTY:
370 ; COMMON-NEXT: // %bb.0:
371 ; COMMON-NEXT: ld.param.u16 %rs1, [test_or_computed_param_0];
372 ; COMMON-NEXT: mov.b16 %rs2, 0;
373 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
374 ; COMMON-NEXT: mov.b16 %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;
378 ; COMMON-NEXT: ret;
387 ; COMMON-LABEL: test_or_imm_0(
388 ; COMMON: {
389 ; COMMON-NEXT: .reg .b32 %r<3>;
390 ; COMMON-EMPTY:
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;
395 ; COMMON-NEXT: ret;
401 ; COMMON-LABEL: test_or_imm_1(
402 ; COMMON: {
403 ; COMMON-NEXT: .reg .b32 %r<3>;
404 ; COMMON-EMPTY:
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;
409 ; COMMON-NEXT: ret;
415 ; COMMON-LABEL: test_xor(
416 ; COMMON: {
417 ; COMMON-NEXT: .reg .b32 %r<4>;
418 ; COMMON-EMPTY:
419 ; COMMON-NEXT: // %bb.0:
420 ; COMMON-NEXT: ld.param.u32 %r2, [test_xor_param_1];
421 ; COMMON-NEXT: ld.param.u32 %r1, [test_xor_param_0];
422 ; COMMON-NEXT: xor.b32 %r3, %r1, %r2;
423 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
424 ; COMMON-NEXT: ret;
430 ; COMMON-LABEL: test_xor_computed(
431 ; COMMON: {
432 ; COMMON-NEXT: .reg .b16 %rs<4>;
433 ; COMMON-NEXT: .reg .b32 %r<4>;
434 ; COMMON-EMPTY:
435 ; COMMON-NEXT: // %bb.0:
436 ; COMMON-NEXT: ld.param.u16 %rs1, [test_xor_computed_param_0];
437 ; COMMON-NEXT: mov.b16 %rs2, 0;
438 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
439 ; COMMON-NEXT: mov.b16 %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;
443 ; COMMON-NEXT: ret;
452 ; COMMON-LABEL: test_xor_imm_0(
453 ; COMMON: {
454 ; COMMON-NEXT: .reg .b32 %r<3>;
455 ; COMMON-EMPTY:
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;
460 ; COMMON-NEXT: ret;
466 ; COMMON-LABEL: test_xor_imm_1(
467 ; COMMON: {
468 ; COMMON-NEXT: .reg .b32 %r<3>;
469 ; COMMON-EMPTY:
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;
474 ; COMMON-NEXT: ret;
480 ; COMMON-LABEL: test_and(
481 ; COMMON: {
482 ; COMMON-NEXT: .reg .b32 %r<4>;
483 ; COMMON-EMPTY:
484 ; COMMON-NEXT: // %bb.0:
485 ; COMMON-NEXT: ld.param.u32 %r2, [test_and_param_1];
486 ; COMMON-NEXT: ld.param.u32 %r1, [test_and_param_0];
487 ; COMMON-NEXT: and.b32 %r3, %r1, %r2;
488 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
489 ; COMMON-NEXT: ret;
497 ; COMMON-LABEL: test_and_computed(
498 ; COMMON: {
499 ; COMMON-NEXT: .reg .b16 %rs<4>;
500 ; COMMON-NEXT: .reg .b32 %r<4>;
501 ; COMMON-EMPTY:
502 ; COMMON-NEXT: // %bb.0:
503 ; COMMON-NEXT: ld.param.u16 %rs1, [test_and_computed_param_0];
504 ; COMMON-NEXT: mov.b16 %rs2, 0;
505 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
506 ; COMMON-NEXT: mov.b16 %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;
510 ; COMMON-NEXT: ret;
519 ; COMMON-LABEL: test_and_imm_0(
520 ; COMMON: {
521 ; COMMON-NEXT: .reg .b32 %r<3>;
522 ; COMMON-EMPTY:
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;
527 ; COMMON-NEXT: ret;
533 ; COMMON-LABEL: test_and_imm_1(
534 ; COMMON: {
535 ; COMMON-NEXT: .reg .b32 %r<3>;
536 ; COMMON-EMPTY:
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;
541 ; COMMON-NEXT: ret;
547 ; COMMON-LABEL: test_ldst_v2i16(
548 ; COMMON: {
549 ; COMMON-NEXT: .reg .b32 %r<2>;
550 ; COMMON-NEXT: .reg .b64 %rd<3>;
551 ; COMMON-EMPTY:
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;
557 ; COMMON-NEXT: ret;
568 ; COMMON-LABEL: test_ldst_v3i16(
569 ; COMMON: {
570 ; COMMON-NEXT: .reg .b64 %rd<5>;
571 ; COMMON-EMPTY:
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;
579 ; COMMON-NEXT: ret;
586 ; COMMON-LABEL: test_ldst_v4i16(
587 ; COMMON: {
588 ; COMMON-NEXT: .reg .b16 %rs<5>;
589 ; COMMON-NEXT: .reg .b64 %rd<3>;
590 ; COMMON-EMPTY:
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};
596 ; COMMON-NEXT: ret;
603 ; COMMON-LABEL: test_ldst_v8i16(
604 ; COMMON: {
605 ; COMMON-NEXT: .reg .b32 %r<5>;
606 ; COMMON-NEXT: .reg .b64 %rd<3>;
607 ; COMMON-EMPTY:
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};
613 ; COMMON-NEXT: ret;
622 ; COMMON-LABEL: test_call(
623 ; COMMON: {
624 ; COMMON-NEXT: .reg .b32 %r<5>;
625 ; COMMON-EMPTY:
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,
637 ; COMMON-NEXT: (
638 ; COMMON-NEXT: param0,
639 ; COMMON-NEXT: param1
640 ; COMMON-NEXT: );
641 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
642 ; COMMON-NEXT: } // callseq 0
643 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
644 ; COMMON-NEXT: ret;
650 ; COMMON-LABEL: test_call_flipped(
651 ; COMMON: {
652 ; COMMON-NEXT: .reg .b32 %r<5>;
653 ; COMMON-EMPTY:
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,
665 ; COMMON-NEXT: (
666 ; COMMON-NEXT: param0,
667 ; COMMON-NEXT: param1
668 ; COMMON-NEXT: );
669 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
670 ; COMMON-NEXT: } // callseq 1
671 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
672 ; COMMON-NEXT: ret;
678 ; COMMON-LABEL: test_tailcall_flipped(
679 ; COMMON: {
680 ; COMMON-NEXT: .reg .b32 %r<5>;
681 ; COMMON-EMPTY:
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,
693 ; COMMON-NEXT: (
694 ; COMMON-NEXT: param0,
695 ; COMMON-NEXT: param1
696 ; COMMON-NEXT: );
697 ; COMMON-NEXT: ld.param.b32 %r3, [retval0];
698 ; COMMON-NEXT: } // callseq 2
699 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
700 ; COMMON-NEXT: ret;
706 ; COMMON-LABEL: test_select(
707 ; COMMON: {
708 ; COMMON-NEXT: .reg .pred %p<2>;
709 ; COMMON-NEXT: .reg .b16 %rs<3>;
710 ; COMMON-NEXT: .reg .b32 %r<4>;
711 ; COMMON-EMPTY:
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;
720 ; COMMON-NEXT: ret;
726 ; COMMON-LABEL: test_select_cc(
727 ; COMMON: {
728 ; COMMON-NEXT: .reg .pred %p<3>;
729 ; COMMON-NEXT: .reg .b16 %rs<11>;
730 ; COMMON-NEXT: .reg .b32 %r<6>;
731 ; COMMON-EMPTY:
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;
747 ; COMMON-NEXT: ret;
754 ; COMMON-LABEL: test_select_cc_i32_i16(
755 ; COMMON: {
756 ; COMMON-NEXT: .reg .pred %p<3>;
757 ; COMMON-NEXT: .reg .b16 %rs<5>;
758 ; COMMON-NEXT: .reg .b32 %r<9>;
759 ; COMMON-EMPTY:
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};
772 ; COMMON-NEXT: ret;
780 ; COMMON-LABEL: test_select_cc_i16_i32(
781 ; COMMON: {
782 ; COMMON-NEXT: .reg .pred %p<3>;
783 ; COMMON-NEXT: .reg .b16 %rs<7>;
784 ; COMMON-NEXT: .reg .b32 %r<8>;
785 ; COMMON-EMPTY:
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;
799 ; COMMON-NEXT: ret;
808 ; COMMON-LABEL: test_trunc_2xi32(
809 ; COMMON: {
810 ; COMMON-NEXT: .reg .b32 %r<4>;
811 ; COMMON-EMPTY:
812 ; COMMON-NEXT: // %bb.0:
813 ; COMMON-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_param_0];
814 ; COMMON-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;
815 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
816 ; COMMON-NEXT: ret;
864 ; COMMON-LABEL: test_trunc_2xi32_muliple_use1(
865 ; COMMON: {
866 ; COMMON-NEXT: .reg .b32 %r<6>;
867 ; COMMON-NEXT: .reg .b64 %rd<2>;
868 ; COMMON-EMPTY:
869 ; COMMON-NEXT: // %bb.0:
870 ; COMMON-NEXT: ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_muliple_use1_param_0];
871 ; COMMON-NEXT: ld.param.u64 %rd1, [test_trunc_2xi32_muliple_use1_param_1];
872 ; COMMON-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;
873 ; COMMON-NEXT: add.s32 %r4, %r2, 1;
874 ; COMMON-NEXT: add.s32 %r5, %r1, 1;
875 ; COMMON-NEXT: st.v2.u32 [%rd1], {%r5, %r4};
876 ; COMMON-NEXT: st.param.b32 [func_retval0], %r3;
877 ; COMMON-NEXT: ret;
887 ; COMMON-LABEL: test_trunc_2xi64(
888 ; COMMON: {
889 ; COMMON-NEXT: .reg .b16 %rs<3>;
890 ; COMMON-NEXT: .reg .b32 %r<2>;
891 ; COMMON-NEXT: .reg .b64 %rd<3>;
892 ; COMMON-EMPTY:
893 ; COMMON-NEXT: // %bb.0:
894 ; COMMON-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
895 ; COMMON-NEXT: cvt.u16.u64 %rs1, %rd2;
896 ; COMMON-NEXT: cvt.u16.u64 %rs2, %rd1;
897 ; COMMON-NEXT: mov.b32 %r1, {%rs2, %rs1};
898 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
899 ; COMMON-NEXT: ret;
905 ; COMMON-LABEL: test_zext_2xi32(
906 ; COMMON: {
907 ; COMMON-NEXT: .reg .b16 %rs<3>;
908 ; COMMON-NEXT: .reg .b32 %r<4>;
909 ; COMMON-EMPTY:
910 ; COMMON-NEXT: // %bb.0:
911 ; COMMON-NEXT: ld.param.u32 %r1, [test_zext_2xi32_param_0];
912 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
913 ; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
914 ; COMMON-NEXT: cvt.u32.u16 %r3, %rs2;
915 ; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r3};
916 ; COMMON-NEXT: ret;
922 ; COMMON-LABEL: test_zext_2xi64(
923 ; COMMON: {
924 ; COMMON-NEXT: .reg .b16 %rs<3>;
925 ; COMMON-NEXT: .reg .b32 %r<2>;
926 ; COMMON-NEXT: .reg .b64 %rd<3>;
927 ; COMMON-EMPTY:
928 ; COMMON-NEXT: // %bb.0:
929 ; COMMON-NEXT: ld.param.u32 %r1, [test_zext_2xi64_param_0];
930 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
931 ; COMMON-NEXT: cvt.u64.u16 %rd1, %rs2;
932 ; COMMON-NEXT: cvt.u64.u16 %rd2, %rs1;
933 ; COMMON-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
934 ; COMMON-NEXT: ret;
940 ; COMMON-LABEL: test_bitcast_i32_to_2xi16(
941 ; COMMON: {
942 ; COMMON-NEXT: .reg .b32 %r<2>;
943 ; COMMON-EMPTY:
944 ; COMMON-NEXT: // %bb.0:
945 ; COMMON-NEXT: ld.param.u32 %r1, [test_bitcast_i32_to_2xi16_param_0];
946 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
947 ; COMMON-NEXT: ret;
953 ; COMMON-LABEL: test_bitcast_2xi16_to_i32(
954 ; COMMON: {
955 ; COMMON-NEXT: .reg .b32 %r<2>;
956 ; COMMON-EMPTY:
957 ; COMMON-NEXT: // %bb.0:
958 ; COMMON-NEXT: ld.param.u32 %r1, [test_bitcast_2xi16_to_i32_param_0];
959 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
960 ; COMMON-NEXT: ret;
966 ; COMMON-LABEL: test_bitcast_2xi16_to_2xhalf(
967 ; COMMON: {
968 ; COMMON-NEXT: .reg .b16 %rs<3>;
969 ; COMMON-NEXT: .reg .b32 %r<2>;
970 ; COMMON-EMPTY:
971 ; COMMON-NEXT: // %bb.0:
972 ; COMMON-NEXT: ld.param.u16 %rs1, [test_bitcast_2xi16_to_2xhalf_param_0];
973 ; COMMON-NEXT: mov.b16 %rs2, 5;
974 ; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
975 ; COMMON-NEXT: st.param.b32 [func_retval0], %r1;
976 ; COMMON-NEXT: ret;
985 ; COMMON-LABEL: test_shufflevector(
986 ; COMMON: {
987 ; COMMON-NEXT: .reg .b16 %rs<3>;
988 ; COMMON-NEXT: .reg .b32 %r<3>;
989 ; COMMON-EMPTY:
990 ; COMMON-NEXT: // %bb.0:
991 ; COMMON-NEXT: ld.param.u32 %r1, [test_shufflevector_param_0];
992 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
993 ; COMMON-NEXT: mov.b32 %r2, {%rs2, %rs1};
994 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
995 ; COMMON-NEXT: ret;
1001 ; COMMON-LABEL: test_insertelement(
1002 ; COMMON: {
1003 ; COMMON-NEXT: .reg .b16 %rs<3>;
1004 ; COMMON-NEXT: .reg .b32 %r<3>;
1005 ; COMMON-EMPTY:
1006 ; COMMON-NEXT: // %bb.0:
1007 ; COMMON-NEXT: ld.param.u16 %rs1, [test_insertelement_param_1];
1008 ; COMMON-NEXT: ld.param.u32 %r1, [test_insertelement_param_0];
1009 ; COMMON-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }
1010 ; COMMON-NEXT: mov.b32 %r2, {%rs2, %rs1};
1011 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
1012 ; COMMON-NEXT: ret;
1018 ; COMMON-LABEL: test_fptosi_2xhalf_to_2xi16(
1019 ; COMMON: {
1020 ; COMMON-NEXT: .reg .b16 %rs<5>;
1021 ; COMMON-NEXT: .reg .b32 %r<3>;
1022 ; COMMON-EMPTY:
1023 ; COMMON-NEXT: // %bb.0:
1024 ; COMMON-NEXT: ld.param.b32 %r1, [test_fptosi_2xhalf_to_2xi16_param_0];
1025 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
1026 ; COMMON-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;
1027 ; COMMON-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;
1028 ; COMMON-NEXT: mov.b32 %r2, {%rs4, %rs3};
1029 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
1030 ; COMMON-NEXT: ret;
1036 ; COMMON-LABEL: test_fptoui_2xhalf_to_2xi16(
1037 ; COMMON: {
1038 ; COMMON-NEXT: .reg .b16 %rs<5>;
1039 ; COMMON-NEXT: .reg .b32 %r<3>;
1040 ; COMMON-EMPTY:
1041 ; COMMON-NEXT: // %bb.0:
1042 ; COMMON-NEXT: ld.param.b32 %r1, [test_fptoui_2xhalf_to_2xi16_param_0];
1043 ; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
1044 ; COMMON-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;
1045 ; COMMON-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;
1046 ; COMMON-NEXT: mov.b32 %r2, {%rs4, %rs3};
1047 ; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
1048 ; COMMON-NEXT: ret;