| ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 3 |
| ; ## Support i16x2 instructions |
| ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 \ |
| ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ |
| ; RUN: | FileCheck -allow-deprecated-dag-overlap %s |
| ; RUN: %if ptxas %{ \ |
| ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 \ |
| ; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ |
| ; RUN: | %ptxas-verify -arch=sm_90 \ |
| ; RUN: %} |
| |
| target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128" |
| |
| define <4 x i8> @test_ret_const() #0 { |
| ; CHECK-LABEL: test_ret_const( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<2>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: mov.b32 %r1, -66911489; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; |
| ; CHECK-NEXT: ret; |
| ret <4 x i8> <i8 -1, i8 2, i8 3, i8 -4> |
| } |
| |
| define i8 @test_extract_0(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_extract_0( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_0_param_0]; |
| ; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %e = extractelement <4 x i8> %a, i32 0 |
| ret i8 %e |
| } |
| |
| define i8 @test_extract_1(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_extract_1( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_1_param_0]; |
| ; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %e = extractelement <4 x i8> %a, i32 1 |
| ret i8 %e |
| } |
| |
| define i8 @test_extract_2(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_extract_2( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_2_param_0]; |
| ; CHECK-NEXT: bfe.u32 %r2, %r1, 16, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %e = extractelement <4 x i8> %a, i32 2 |
| ret i8 %e |
| } |
| |
| define i8 @test_extract_3(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_extract_3( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_3_param_0]; |
| ; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %e = extractelement <4 x i8> %a, i32 3 |
| ret i8 %e |
| } |
| |
| define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 { |
| ; CHECK-LABEL: test_extract_i( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-NEXT: .reg .b64 %rd<2>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_extract_i_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_i_param_0]; |
| ; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; |
| ; CHECK-NEXT: shl.b32 %r3, %r2, 3; |
| ; CHECK-NEXT: bfe.u32 %r4, %r1, %r3, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; |
| ; CHECK-NEXT: ret; |
| %e = extractelement <4 x i8> %a, i64 %idx |
| ret i8 %e |
| } |
| |
| define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_add( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_add_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_add_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; |
| ; CHECK-NEXT: add.s16 %rs3, %rs2, %rs1; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; |
| ; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; |
| ; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; |
| ; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; |
| ; CHECK-NEXT: add.s16 %rs12, %rs11, %rs10; |
| ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %r = add <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_add_imm_0( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<9>; |
| ; CHECK-NEXT: .reg .b32 %r<14>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_0_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r2, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; |
| ; CHECK-NEXT: add.s16 %rs2, %rs1, 1; |
| ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; |
| ; CHECK-NEXT: add.s16 %rs4, %rs3, 2; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; |
| ; CHECK-NEXT: bfi.b32 %r6, %r5, %r3, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; |
| ; CHECK-NEXT: add.s16 %rs6, %rs5, 3; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r6, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; |
| ; CHECK-NEXT: add.s16 %rs8, %rs7, 4; |
| ; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; |
| ; CHECK-NEXT: bfi.b32 %r12, %r11, %r9, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; |
| ; CHECK-NEXT: ret; |
| %r = add <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_add_imm_1( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<9>; |
| ; CHECK-NEXT: .reg .b32 %r<14>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_1_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r2, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; |
| ; CHECK-NEXT: add.s16 %rs2, %rs1, 1; |
| ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; |
| ; CHECK-NEXT: add.s16 %rs4, %rs3, 2; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; |
| ; CHECK-NEXT: bfi.b32 %r6, %r5, %r3, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; |
| ; CHECK-NEXT: add.s16 %rs6, %rs5, 3; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r6, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; |
| ; CHECK-NEXT: add.s16 %rs8, %rs7, 4; |
| ; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; |
| ; CHECK-NEXT: bfi.b32 %r12, %r11, %r9, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; |
| ; CHECK-NEXT: ret; |
| %r = add <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_sub( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_sub_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_sub_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; |
| ; CHECK-NEXT: sub.s16 %rs3, %rs2, %rs1; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; |
| ; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; |
| ; CHECK-NEXT: sub.s16 %rs9, %rs8, %rs7; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; |
| ; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; |
| ; CHECK-NEXT: sub.s16 %rs12, %rs11, %rs10; |
| ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %r = sub <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_smax( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_smax_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_smax_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r3, %r4; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: setp.gt.s32 %p2, %r5, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: setp.gt.s32 %p3, %r7, %r8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: setp.gt.s32 %p4, %r9, %r10; |
| ; CHECK-NEXT: selp.b32 %r11, %r9, %r10, %p4; |
| ; CHECK-NEXT: selp.b32 %r12, %r7, %r8, %p3; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r14, %r5, %r6, %p2; |
| ; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r16, %r3, %r4, %p1; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp sgt <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_umax( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_umax_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_umax_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p1, %r3, %r4; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p2, %r5, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p3, %r7, %r8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p4, %r9, %r10; |
| ; CHECK-NEXT: selp.b32 %r11, %r9, %r10, %p4; |
| ; CHECK-NEXT: selp.b32 %r12, %r7, %r8, %p3; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r14, %r5, %r6, %p2; |
| ; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r16, %r3, %r4, %p1; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp ugt <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_smin( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_smin_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_smin_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: setp.le.s32 %p1, %r3, %r4; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: setp.le.s32 %p2, %r5, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: setp.le.s32 %p3, %r7, %r8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: setp.le.s32 %p4, %r9, %r10; |
| ; CHECK-NEXT: selp.b32 %r11, %r9, %r10, %p4; |
| ; CHECK-NEXT: selp.b32 %r12, %r7, %r8, %p3; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r14, %r5, %r6, %p2; |
| ; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r16, %r3, %r4, %p1; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp sle <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_umin( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_umin_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_umin_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: setp.ls.u32 %p1, %r3, %r4; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: setp.ls.u32 %p2, %r5, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: setp.ls.u32 %p3, %r7, %r8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: setp.ls.u32 %p4, %r9, %r10; |
| ; CHECK-NEXT: selp.b32 %r11, %r9, %r10, %p4; |
| ; CHECK-NEXT: selp.b32 %r12, %r7, %r8, %p3; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r14, %r5, %r6, %p2; |
| ; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r16, %r3, %r4, %p1; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp ule <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { |
| ; CHECK-LABEL: test_eq( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<24>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_eq_param_2]; |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_eq_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_eq_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 24, 8; |
| ; CHECK-NEXT: setp.eq.u32 %p1, %r5, %r4; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 16, 8; |
| ; CHECK-NEXT: setp.eq.u32 %p2, %r7, %r6; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 8, 8; |
| ; CHECK-NEXT: setp.eq.u32 %p3, %r9, %r8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 0, 8; |
| ; CHECK-NEXT: setp.eq.u32 %p4, %r11, %r10; |
| ; CHECK-NEXT: bfe.s32 %r12, %r3, 0, 8; |
| ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; |
| ; CHECK-NEXT: bfe.s32 %r14, %r3, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; |
| ; CHECK-NEXT: bfi.b32 %r16, %r15, %r13, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r17, %r3, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; |
| ; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r20, %r3, 24, 8; |
| ; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; |
| ; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp eq <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { |
| ; CHECK-LABEL: test_ne( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<24>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_ne_param_2]; |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_ne_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_ne_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r4, %r2, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 24, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p1, %r5, %r4; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 16, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p2, %r7, %r6; |
| ; CHECK-NEXT: bfe.s32 %r8, %r2, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 8, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p3, %r9, %r8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 0, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p4, %r11, %r10; |
| ; CHECK-NEXT: bfe.s32 %r12, %r3, 0, 8; |
| ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; |
| ; CHECK-NEXT: bfe.s32 %r14, %r3, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; |
| ; CHECK-NEXT: bfi.b32 %r16, %r15, %r13, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r17, %r3, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; |
| ; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r20, %r3, 24, 8; |
| ; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; |
| ; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; |
| ; CHECK-NEXT: ret; |
| %cmp = icmp ne <4 x i8> %a, %b |
| %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_mul( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<19>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_mul_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_mul_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; |
| ; CHECK-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; |
| ; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; |
| ; CHECK-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; |
| ; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; |
| ; CHECK-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; |
| ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; |
| ; CHECK-NEXT: ret; |
| %r = mul <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_or( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<7>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_or_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r4, [test_or_param_0]; |
| ; CHECK-NEXT: or.b32 %r5, %r4, %r3; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; |
| ; CHECK-NEXT: ret; |
| %r = or <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_or_computed(i8 %a) { |
| ; CHECK-LABEL: test_or_computed( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<2>; |
| ; CHECK-NEXT: .reg .b32 %r<9>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_or_computed_param_0]; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; |
| ; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; |
| ; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; |
| ; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; |
| ; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; |
| ; CHECK-NEXT: or.b32 %r8, %r6, %r4; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; |
| ; CHECK-NEXT: ret; |
| %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 |
| %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 |
| %r = or <4 x i8> %ins.1, %ins.0 |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_or_imm_0(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_or_imm_0( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_0_param_0]; |
| ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = or <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_or_imm_1( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_1_param_0]; |
| ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = or <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_xor( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<7>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_xor_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r4, [test_xor_param_0]; |
| ; CHECK-NEXT: xor.b32 %r5, %r4, %r3; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; |
| ; CHECK-NEXT: ret; |
| %r = xor <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_xor_computed(i8 %a) { |
| ; CHECK-LABEL: test_xor_computed( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<2>; |
| ; CHECK-NEXT: .reg .b32 %r<9>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_xor_computed_param_0]; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; |
| ; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; |
| ; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; |
| ; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; |
| ; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; |
| ; CHECK-NEXT: xor.b32 %r8, %r6, %r4; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; |
| ; CHECK-NEXT: ret; |
| %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 |
| %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 |
| %r = xor <4 x i8> %ins.1, %ins.0 |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_xor_imm_0(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_xor_imm_0( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_0_param_0]; |
| ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = xor <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_xor_imm_1( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_1_param_0]; |
| ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = xor <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_and( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<7>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_and_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r4, [test_and_param_0]; |
| ; CHECK-NEXT: and.b32 %r5, %r4, %r3; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; |
| ; CHECK-NEXT: ret; |
| %r = and <4 x i8> %a, %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_and_computed(i8 %a) { |
| ; CHECK-LABEL: test_and_computed( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<2>; |
| ; CHECK-NEXT: .reg .b32 %r<9>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_and_computed_param_0]; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; |
| ; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; |
| ; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; |
| ; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; |
| ; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; |
| ; CHECK-NEXT: and.b32 %r8, %r6, %r4; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; |
| ; CHECK-NEXT: ret; |
| %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 |
| %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 |
| %r = and <4 x i8> %ins.1, %ins.0 |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_and_imm_0(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_and_imm_0( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_0_param_0]; |
| ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = and <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_and_imm_1( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_1_param_0]; |
| ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = and <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4> |
| ret <4 x i8> %r |
| } |
| |
| define void @test_ldst_v2i8(ptr %a, ptr %b) { |
| ; CHECK-LABEL: test_ldst_v2i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<2>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_ldst_v2i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_ldst_v2i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: st.u32 [%rd2], %r1; |
| ; CHECK-NEXT: ret; |
| %t1 = load <4 x i8>, ptr %a |
| store <4 x i8> %t1, ptr %b, align 16 |
| ret void |
| } |
| |
| define void @test_ldst_v3i8(ptr %a, ptr %b) { |
| ; CHECK-LABEL: test_ldst_v3i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_ldst_v3i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_ldst_v3i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: st.u16 [%rd2], %r1; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 16, 8; |
| ; CHECK-NEXT: st.u8 [%rd2+2], %r3; |
| ; CHECK-NEXT: ret; |
| %t1 = load <3 x i8>, ptr %a |
| store <3 x i8> %t1, ptr %b, align 16 |
| ret void |
| } |
| |
| define void @test_ldst_v4i8(ptr %a, ptr %b) { |
| ; CHECK-LABEL: test_ldst_v4i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<2>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_ldst_v4i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_ldst_v4i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: st.u32 [%rd2], %r1; |
| ; CHECK-NEXT: ret; |
| %t1 = load <4 x i8>, ptr %a |
| store <4 x i8> %t1, ptr %b, align 16 |
| ret void |
| } |
| |
| define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) { |
| ; CHECK-LABEL: test_ldst_v4i8_unaligned( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_ldst_v4i8_unaligned_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_ldst_v4i8_unaligned_param_0]; |
| ; CHECK-NEXT: ld.u8 %r1, [%rd1]; |
| ; CHECK-NEXT: ld.u8 %r2, [%rd1+1]; |
| ; CHECK-NEXT: ld.u8 %r3, [%rd1+2]; |
| ; CHECK-NEXT: ld.u8 %r4, [%rd1+3]; |
| ; CHECK-NEXT: st.u8 [%rd2+3], %r4; |
| ; CHECK-NEXT: st.u8 [%rd2+2], %r3; |
| ; CHECK-NEXT: st.u8 [%rd2+1], %r2; |
| ; CHECK-NEXT: st.u8 [%rd2], %r1; |
| ; CHECK-NEXT: ret; |
| %t1 = load <4 x i8>, ptr %a, align 1 |
| store <4 x i8> %t1, ptr %b, align 1 |
| ret void |
| } |
| |
| |
| define void @test_ldst_v8i8(ptr %a, ptr %b) { |
| ; CHECK-LABEL: test_ldst_v8i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-NEXT: .reg .b64 %rd<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_ldst_v8i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_ldst_v8i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: ld.u32 %r2, [%rd1+4]; |
| ; CHECK-NEXT: st.u32 [%rd2+4], %r2; |
| ; CHECK-NEXT: st.u32 [%rd2], %r1; |
| ; CHECK-NEXT: ret; |
| %t1 = load <8 x i8>, ptr %a |
| store <8 x i8> %t1, ptr %b, align 16 |
| ret void |
| } |
| |
| declare <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) #0 |
| |
| define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_call( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_call_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_call_param_0]; |
| ; CHECK-NEXT: { // callseq 0, 0 |
| ; CHECK-NEXT: .reg .b32 temp_param_reg; |
| ; CHECK-NEXT: .param .align 4 .b8 param0[4]; |
| ; CHECK-NEXT: st.param.b32 [param0+0], %r1; |
| ; CHECK-NEXT: .param .align 4 .b8 param1[4]; |
| ; CHECK-NEXT: st.param.b32 [param1+0], %r2; |
| ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; |
| ; CHECK-NEXT: call.uni (retval0), |
| ; CHECK-NEXT: test_callee, |
| ; CHECK-NEXT: ( |
| ; CHECK-NEXT: param0, |
| ; CHECK-NEXT: param1 |
| ; CHECK-NEXT: ); |
| ; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; |
| ; CHECK-NEXT: } // callseq 0 |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_call_flipped( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_call_flipped_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_call_flipped_param_0]; |
| ; CHECK-NEXT: { // callseq 1, 0 |
| ; CHECK-NEXT: .reg .b32 temp_param_reg; |
| ; CHECK-NEXT: .param .align 4 .b8 param0[4]; |
| ; CHECK-NEXT: st.param.b32 [param0+0], %r2; |
| ; CHECK-NEXT: .param .align 4 .b8 param1[4]; |
| ; CHECK-NEXT: st.param.b32 [param1+0], %r1; |
| ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; |
| ; CHECK-NEXT: call.uni (retval0), |
| ; CHECK-NEXT: test_callee, |
| ; CHECK-NEXT: ( |
| ; CHECK-NEXT: param0, |
| ; CHECK-NEXT: param1 |
| ; CHECK-NEXT: ); |
| ; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; |
| ; CHECK-NEXT: } // callseq 1 |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_tailcall_flipped( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_tailcall_flipped_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_tailcall_flipped_param_0]; |
| ; CHECK-NEXT: { // callseq 2, 0 |
| ; CHECK-NEXT: .reg .b32 temp_param_reg; |
| ; CHECK-NEXT: .param .align 4 .b8 param0[4]; |
| ; CHECK-NEXT: st.param.b32 [param0+0], %r2; |
| ; CHECK-NEXT: .param .align 4 .b8 param1[4]; |
| ; CHECK-NEXT: st.param.b32 [param1+0], %r1; |
| ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; |
| ; CHECK-NEXT: call.uni (retval0), |
| ; CHECK-NEXT: test_callee, |
| ; CHECK-NEXT: ( |
| ; CHECK-NEXT: param0, |
| ; CHECK-NEXT: param1 |
| ; CHECK-NEXT: ); |
| ; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; |
| ; CHECK-NEXT: } // callseq 2 |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 { |
| ; CHECK-LABEL: test_select( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<2>; |
| ; CHECK-NEXT: .reg .b16 %rs<3>; |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_select_param_2]; |
| ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; |
| ; CHECK-NEXT: setp.eq.b16 %p1, %rs2, 1; |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_select_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_select_param_0]; |
| ; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %r = select i1 %c, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> %d) #0 { |
| ; CHECK-LABEL: test_select_cc( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<29>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r4, [test_select_cc_param_3]; |
| ; CHECK-NEXT: ld.param.u32 %r3, [test_select_cc_param_2]; |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_select_cc_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_select_cc_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r5, %r4, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r3, 24, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p1, %r6, %r5; |
| ; CHECK-NEXT: bfe.s32 %r7, %r4, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r3, 16, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p2, %r8, %r7; |
| ; CHECK-NEXT: bfe.s32 %r9, %r4, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r3, 8, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p3, %r10, %r9; |
| ; CHECK-NEXT: bfe.s32 %r11, %r4, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r12, %r3, 0, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p4, %r12, %r11; |
| ; CHECK-NEXT: bfe.s32 %r13, %r2, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r1, 0, 8; |
| ; CHECK-NEXT: selp.b32 %r15, %r14, %r13, %p4; |
| ; CHECK-NEXT: bfe.s32 %r16, %r2, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r17, %r1, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3; |
| ; CHECK-NEXT: bfi.b32 %r19, %r18, %r15, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r20, %r2, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r21, %r1, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2; |
| ; CHECK-NEXT: bfi.b32 %r23, %r22, %r19, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r24, %r2, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r25, %r1, 24, 8; |
| ; CHECK-NEXT: selp.b32 %r26, %r25, %r24, %p1; |
| ; CHECK-NEXT: bfi.b32 %r27, %r26, %r23, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r27; |
| ; CHECK-NEXT: ret; |
| %cc = icmp ne <4 x i8> %c, %d |
| %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b, |
| ; CHECK-LABEL: test_select_cc_i32_i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<23>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v4.u32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1]; |
| ; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0]; |
| ; CHECK-NEXT: ld.param.u32 %r10, [test_select_cc_i32_i8_param_3]; |
| ; CHECK-NEXT: ld.param.u32 %r9, [test_select_cc_i32_i8_param_2]; |
| ; CHECK-NEXT: bfe.s32 %r11, %r10, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r12, %r9, 0, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p1, %r12, %r11; |
| ; CHECK-NEXT: bfe.s32 %r13, %r10, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r9, 8, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p2, %r14, %r13; |
| ; CHECK-NEXT: bfe.s32 %r15, %r10, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r16, %r9, 16, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p3, %r16, %r15; |
| ; CHECK-NEXT: bfe.s32 %r17, %r10, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r18, %r9, 24, 8; |
| ; CHECK-NEXT: setp.ne.u32 %p4, %r18, %r17; |
| ; CHECK-NEXT: selp.b32 %r19, %r4, %r8, %p4; |
| ; CHECK-NEXT: selp.b32 %r20, %r3, %r7, %p3; |
| ; CHECK-NEXT: selp.b32 %r21, %r2, %r6, %p2; |
| ; CHECK-NEXT: selp.b32 %r22, %r1, %r5, %p1; |
| ; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r22, %r21, %r20, %r19}; |
| ; CHECK-NEXT: ret; |
| <4 x i8> %c, <4 x i8> %d) #0 { |
| %cc = icmp ne <4 x i8> %c, %d |
| %r = select <4 x i1> %cc, <4 x i32> %a, <4 x i32> %b |
| ret <4 x i32> %r |
| } |
| |
| define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, |
| ; CHECK-LABEL: test_select_cc_i8_i32( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<27>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v4.u32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3]; |
| ; CHECK-NEXT: ld.param.v4.u32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2]; |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_select_cc_i8_i32_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_select_cc_i8_i32_param_0]; |
| ; CHECK-NEXT: setp.ne.s32 %p1, %r6, %r10; |
| ; CHECK-NEXT: setp.ne.s32 %p2, %r5, %r9; |
| ; CHECK-NEXT: setp.ne.s32 %p3, %r4, %r8; |
| ; CHECK-NEXT: setp.ne.s32 %p4, %r3, %r7; |
| ; CHECK-NEXT: bfe.s32 %r11, %r2, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r12, %r1, 0, 8; |
| ; CHECK-NEXT: selp.b32 %r13, %r12, %r11, %p4; |
| ; CHECK-NEXT: bfe.s32 %r14, %r2, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r15, %r1, 8, 8; |
| ; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r18, %r2, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r19, %r1, 16, 8; |
| ; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2; |
| ; CHECK-NEXT: bfi.b32 %r21, %r20, %r17, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r22, %r2, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r23, %r1, 24, 8; |
| ; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p1; |
| ; CHECK-NEXT: bfi.b32 %r25, %r24, %r21, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; |
| ; CHECK-NEXT: ret; |
| <4 x i32> %c, <4 x i32> %d) #0 { |
| %cc = icmp ne <4 x i32> %c, %d |
| %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b |
| ret <4 x i8> %r |
| } |
| |
| |
| define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { |
| ; CHECK-LABEL: test_trunc_2xi32( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<9>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0]; |
| ; CHECK-NEXT: bfi.b32 %r5, %r2, %r1, 8, 8; |
| ; CHECK-NEXT: bfi.b32 %r6, %r3, %r5, 16, 8; |
| ; CHECK-NEXT: bfi.b32 %r7, %r4, %r6, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; |
| ; CHECK-NEXT: ret; |
| %r = trunc <4 x i32> %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { |
| ; CHECK-LABEL: test_trunc_2xi64( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<9>; |
| ; CHECK-NEXT: .reg .b64 %rd<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v2.u64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; |
| ; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; |
| ; CHECK-NEXT: cvt.u32.u64 %r1, %rd1; |
| ; CHECK-NEXT: cvt.u32.u64 %r2, %rd2; |
| ; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u32.u64 %r4, %rd3; |
| ; CHECK-NEXT: bfi.b32 %r5, %r4, %r3, 16, 8; |
| ; CHECK-NEXT: cvt.u32.u64 %r6, %rd4; |
| ; CHECK-NEXT: bfi.b32 %r7, %r6, %r5, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; |
| ; CHECK-NEXT: ret; |
| %r = trunc <4 x i64> %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_zext_2xi32( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<6>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_zext_2xi32_param_0]; |
| ; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; |
| ; CHECK-NEXT: bfe.u32 %r3, %r1, 16, 8; |
| ; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; |
| ; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; |
| ; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r5, %r4, %r3, %r2}; |
| ; CHECK-NEXT: ret; |
| %r = zext <4 x i8> %a to <4 x i32> |
| ret <4 x i32> %r |
| } |
| |
| define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_zext_2xi64( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<6>; |
| ; CHECK-NEXT: .reg .b64 %rd<9>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_zext_2xi64_param_0]; |
| ; CHECK-NEXT: bfe.s32 %r2, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.u64.u32 %rd1, %r2; |
| ; CHECK-NEXT: and.b64 %rd2, %rd1, 255; |
| ; CHECK-NEXT: bfe.s32 %r3, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.u64.u32 %rd3, %r3; |
| ; CHECK-NEXT: and.b64 %rd4, %rd3, 255; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.u64.u32 %rd5, %r4; |
| ; CHECK-NEXT: and.b64 %rd6, %rd5, 255; |
| ; CHECK-NEXT: bfe.s32 %r5, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.u64.u32 %rd7, %r5; |
| ; CHECK-NEXT: and.b64 %rd8, %rd7, 255; |
| ; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd8, %rd6}; |
| ; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd2}; |
| ; CHECK-NEXT: ret; |
| %r = zext <4 x i8> %a to <4 x i64> |
| ret <4 x i64> %r |
| } |
| |
| define <4 x i8> @test_bitcast_i32_to_4xi8(i32 %a) #0 { |
| ; CHECK-LABEL: test_bitcast_i32_to_4xi8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_bitcast_i32_to_4xi8_param_0]; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; |
| ; CHECK-NEXT: ret; |
| %r = bitcast i32 %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_bitcast_float_to_4xi8(float %a) #0 { |
| ; CHECK-LABEL: test_bitcast_float_to_4xi8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<2>; |
| ; CHECK-NEXT: .reg .f32 %f<2>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.f32 %f1, [test_bitcast_float_to_4xi8_param_0]; |
| ; CHECK-NEXT: mov.b32 %r1, %f1; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; |
| ; CHECK-NEXT: ret; |
| %r = bitcast float %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define i32 @test_bitcast_4xi8_to_i32(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_bitcast_4xi8_to_i32( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_i32_param_0]; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %r = bitcast <4 x i8> %a to i32 |
| ret i32 %r |
| } |
| |
| define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_bitcast_4xi8_to_float( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-NEXT: .reg .f32 %f<2>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_float_param_0]; |
| ; CHECK-NEXT: mov.b32 %f1, %r2; |
| ; CHECK-NEXT: st.param.f32 [func_retval0+0], %f1; |
| ; CHECK-NEXT: ret; |
| %r = bitcast <4 x i8> %a to float |
| ret float %r |
| } |
| |
| |
| define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { |
| ; CHECK-LABEL: test_bitcast_4xi8_to_2xhalf( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<2>; |
| ; CHECK-NEXT: .reg .b32 %r<6>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; |
| ; CHECK-NEXT: bfi.b32 %r2, 5, %r1, 8, 8; |
| ; CHECK-NEXT: bfi.b32 %r3, 6, %r2, 16, 8; |
| ; CHECK-NEXT: bfi.b32 %r4, 7, %r3, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; |
| ; CHECK-NEXT: ret; |
| %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0 |
| %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 |
| %ins.2 = insertelement <4 x i8> %ins.1, i8 6, i32 2 |
| %ins.3 = insertelement <4 x i8> %ins.2, i8 7, i32 3 |
| %r = bitcast <4 x i8> %ins.3 to <2 x half> |
| ret <2 x half> %r |
| } |
| |
| |
| define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 { |
| ; CHECK-LABEL: test_shufflevector( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_param_0]; |
| ; CHECK-NEXT: // implicit-def: %r3 |
| ; CHECK-NEXT: prmt.b32 %r2, %r1, %r3, 291; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; |
| ; CHECK-NEXT: ret; |
| %s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> <i32 3, i32 2, i32 1, i32 0> |
| ret <4 x i8> %s |
| } |
| |
| define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 { |
| ; CHECK-LABEL: test_shufflevector_2( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b32 %r<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u32 %r2, [test_shufflevector_2_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_2_param_0]; |
| ; CHECK-NEXT: prmt.b32 %r3, %r1, %r2, 9527; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> <i32 7, i32 3, i32 5, i32 2> |
| ret <4 x i8> %s |
| } |
| |
| |
| define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 { |
| ; CHECK-LABEL: test_insertelement( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<2>; |
| ; CHECK-NEXT: .reg .b32 %r<5>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.u8 %rs1, [test_insertelement_param_1]; |
| ; CHECK-NEXT: ld.param.u32 %r1, [test_insertelement_param_0]; |
| ; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; |
| ; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; |
| ; CHECK-NEXT: ret; |
| %i = insertelement <4 x i8> %a, i8 %x, i64 1 |
| ret <4 x i8> %i |
| } |
| |
| define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { |
| ; CHECK-LABEL: test_fptosi_4xhalf_to_4xi8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<15>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptosi_4xhalf_to_4xi8_param_0]; |
| ; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3; |
| ; CHECK-NEXT: cvt.rzi.s16.f16 %rs3, %rs2; |
| ; CHECK-NEXT: cvt.rzi.s16.f16 %rs4, %rs1; |
| ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; |
| ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; |
| ; CHECK-NEXT: cvt.u32.u16 %r6, %rs5; |
| ; CHECK-NEXT: cvt.u32.u16 %r7, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r8, %r7, %r6, 8, 8; |
| ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; |
| ; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs8; |
| ; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs7; |
| ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; |
| ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; |
| ; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; |
| ; CHECK-NEXT: bfi.b32 %r11, %r10, %r8, 16, 8; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; |
| ; CHECK-NEXT: ret; |
| %r = fptosi <4 x half> %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { |
| ; CHECK-LABEL: test_fptoui_4xhalf_to_4xi8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<15>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptoui_4xhalf_to_4xi8_param_0]; |
| ; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3; |
| ; CHECK-NEXT: cvt.rzi.u16.f16 %rs3, %rs2; |
| ; CHECK-NEXT: cvt.rzi.u16.f16 %rs4, %rs1; |
| ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; |
| ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; |
| ; CHECK-NEXT: cvt.u32.u16 %r6, %rs5; |
| ; CHECK-NEXT: cvt.u32.u16 %r7, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r8, %r7, %r6, 8, 8; |
| ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; |
| ; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs8; |
| ; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs7; |
| ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; |
| ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; |
| ; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; |
| ; CHECK-NEXT: bfi.b32 %r11, %r10, %r8, 16, 8; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 24, 8; |
| ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; |
| ; CHECK-NEXT: ret; |
| %r = fptoui <4 x half> %a to <4 x i8> |
| ret <4 x i8> %r |
| } |
| |
| define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) { |
| ; CHECK-LABEL: test_srem_v4i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<13>; |
| ; CHECK-NEXT: .reg .b32 %r<18>; |
| ; CHECK-NEXT: .reg .b64 %rd<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: // %entry |
| ; CHECK-NEXT: ld.param.u64 %rd3, [test_srem_v4i8_param_2]; |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_srem_v4i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_srem_v4i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: ld.u32 %r2, [%rd2]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs1, %r3; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs2, %r4; |
| ; CHECK-NEXT: rem.s16 %rs3, %rs2, %rs1; |
| ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; |
| ; CHECK-NEXT: bfe.s32 %r6, %r2, 8, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs4, %r6; |
| ; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs5, %r7; |
| ; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4; |
| ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; |
| ; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs7, %r10; |
| ; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs8, %r11; |
| ; CHECK-NEXT: rem.s16 %rs9, %rs8, %rs7; |
| ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs10, %r14; |
| ; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs11, %r15; |
| ; CHECK-NEXT: rem.s16 %rs12, %rs11, %rs10; |
| ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; |
| ; CHECK-NEXT: st.u32 [%rd3], %r17; |
| ; CHECK-NEXT: ret; |
| entry: |
| %t57 = load <4 x i8>, ptr %a, align 4 |
| %t59 = load <4 x i8>, ptr %b, align 4 |
| %x = srem <4 x i8> %t57, %t59 |
| store <4 x i8> %x, ptr %c, align 4 |
| ret void |
| } |
| |
| ;; v3i8 lowering, especially for unaligned loads is terrible. We end up doing |
| ;; tons of pointless scalar_to_vector/bitcast/extract_elt on v2i16/v4i8, which |
| ;; is further complicated by LLVM trying to use i16 as an intermediate type, |
| ;; because we don't have i8 registers. It's a mess. |
| ;; Ideally we want to split it into element-wise ops, but legalizer can't handle |
| ;; odd-sized vectors. TL;DR; don't use odd-sized vectors of v8. |
| define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { |
| ; CHECK-LABEL: test_srem_v3i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<20>; |
| ; CHECK-NEXT: .reg .b32 %r<16>; |
| ; CHECK-NEXT: .reg .b64 %rd<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: // %entry |
| ; CHECK-NEXT: ld.param.u64 %rd3, [test_srem_v3i8_param_2]; |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_srem_v3i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_srem_v3i8_param_0]; |
| ; CHECK-NEXT: ld.u8 %rs1, [%rd1]; |
| ; CHECK-NEXT: ld.u8 %rs2, [%rd1+1]; |
| ; CHECK-NEXT: shl.b16 %rs3, %rs2, 8; |
| ; CHECK-NEXT: or.b16 %rs4, %rs3, %rs1; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs4; |
| ; CHECK-NEXT: ld.s8 %rs5, [%rd1+2]; |
| ; CHECK-NEXT: ld.u8 %rs6, [%rd2]; |
| ; CHECK-NEXT: ld.u8 %rs7, [%rd2+1]; |
| ; CHECK-NEXT: shl.b16 %rs8, %rs7, 8; |
| ; CHECK-NEXT: or.b16 %rs9, %rs8, %rs6; |
| ; CHECK-NEXT: cvt.u32.u16 %r3, %rs9; |
| ; CHECK-NEXT: ld.s8 %rs10, [%rd2+2]; |
| ; CHECK-NEXT: bfe.s32 %r5, %r3, 0, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs11, %r5; |
| ; CHECK-NEXT: bfe.s32 %r6, %r1, 0, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs12, %r6; |
| ; CHECK-NEXT: rem.s16 %rs13, %rs12, %rs11; |
| ; CHECK-NEXT: cvt.u32.u16 %r7, %rs13; |
| ; CHECK-NEXT: bfe.s32 %r8, %r3, 8, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs14, %r8; |
| ; CHECK-NEXT: bfe.s32 %r9, %r1, 8, 8; |
| ; CHECK-NEXT: cvt.s8.s32 %rs15, %r9; |
| ; CHECK-NEXT: rem.s16 %rs16, %rs15, %rs14; |
| ; CHECK-NEXT: cvt.u32.u16 %r10, %rs16; |
| ; CHECK-NEXT: bfi.b32 %r11, %r10, %r7, 8, 8; |
| ; CHECK-NEXT: // implicit-def: %r13 |
| ; CHECK-NEXT: bfi.b32 %r12, %r13, %r11, 16, 8; |
| ; CHECK-NEXT: // implicit-def: %r15 |
| ; CHECK-NEXT: bfi.b32 %r14, %r15, %r12, 24, 8; |
| ; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10; |
| ; CHECK-NEXT: cvt.u16.u32 %rs18, %r14; |
| ; CHECK-NEXT: st.u8 [%rd3], %rs18; |
| ; CHECK-NEXT: shr.u16 %rs19, %rs18, 8; |
| ; CHECK-NEXT: st.u8 [%rd3+1], %rs19; |
| ; CHECK-NEXT: st.u8 [%rd3+2], %rs17; |
| ; CHECK-NEXT: ret; |
| entry: |
| %t57 = load <3 x i8>, ptr %a, align 1 |
| %t59 = load <3 x i8>, ptr %b, align 1 |
| %x = srem <3 x i8> %t57, %t59 |
| store <3 x i8> %x, ptr %c, align 1 |
| ret void |
| } |
| |
| define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) { |
| ; CHECK-LABEL: test_sext_v4i1_to_v4i8( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .pred %p<5>; |
| ; CHECK-NEXT: .reg .b32 %r<18>; |
| ; CHECK-NEXT: .reg .b64 %rd<4>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: // %entry |
| ; CHECK-NEXT: ld.param.u64 %rd3, [test_sext_v4i1_to_v4i8_param_2]; |
| ; CHECK-NEXT: ld.param.u64 %rd2, [test_sext_v4i1_to_v4i8_param_1]; |
| ; CHECK-NEXT: ld.param.u64 %rd1, [test_sext_v4i1_to_v4i8_param_0]; |
| ; CHECK-NEXT: ld.u32 %r1, [%rd1]; |
| ; CHECK-NEXT: ld.u32 %r2, [%rd2]; |
| ; CHECK-NEXT: bfe.s32 %r3, %r2, 24, 8; |
| ; CHECK-NEXT: bfe.s32 %r4, %r1, 24, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p1, %r4, %r3; |
| ; CHECK-NEXT: bfe.s32 %r5, %r2, 16, 8; |
| ; CHECK-NEXT: bfe.s32 %r6, %r1, 16, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p2, %r6, %r5; |
| ; CHECK-NEXT: bfe.s32 %r7, %r2, 8, 8; |
| ; CHECK-NEXT: bfe.s32 %r8, %r1, 8, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p3, %r8, %r7; |
| ; CHECK-NEXT: bfe.s32 %r9, %r2, 0, 8; |
| ; CHECK-NEXT: bfe.s32 %r10, %r1, 0, 8; |
| ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; |
| ; CHECK-NEXT: selp.s32 %r11, -1, 0, %p4; |
| ; CHECK-NEXT: selp.s32 %r12, -1, 0, %p3; |
| ; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; |
| ; CHECK-NEXT: selp.s32 %r14, -1, 0, %p2; |
| ; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; |
| ; CHECK-NEXT: selp.s32 %r16, -1, 0, %p1; |
| ; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; |
| ; CHECK-NEXT: st.u32 [%rd3], %r17; |
| ; CHECK-NEXT: ret; |
| entry: |
| %t1 = load <4 x i8>, ptr %a, align 4 |
| %t2 = load <4 x i8>, ptr %b, align 4 |
| %t5 = icmp ugt <4 x i8> %t1, %t2 |
| %t6 = sext <4 x i1> %t5 to <4 x i8> |
| store <4 x i8> %t6, ptr %c, align 4 |
| ret void |
| } |
| |
| attributes #0 = { nounwind } |