213 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}4target triple = "nvptx64-nvidia-cuda"5 6 7define i16 @test_v2i8(i16 %a) {8; CHECK-LABEL: test_v2i8(9; CHECK: {10; CHECK-NEXT: .reg .b16 %rs<5>;11; CHECK-NEXT: .reg .b32 %r<2>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b16 %rs1, [test_v2i8_param_0];15; CHECK-NEXT: cvt.s16.s8 %rs2, %rs1;16; CHECK-NEXT: shr.s16 %rs3, %rs1, 8;17; CHECK-NEXT: add.s16 %rs4, %rs2, %rs3;18; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;19; CHECK-NEXT: st.param.b32 [func_retval0], %r1;20; CHECK-NEXT: ret;21 %v = bitcast i16 %a to <2 x i8>22 %r0 = extractelement <2 x i8> %v, i64 023 %r1 = extractelement <2 x i8> %v, i64 124 %r0i = sext i8 %r0 to i1625 %r1i = sext i8 %r1 to i1626 %r01 = add i16 %r0i, %r1i27 ret i16 %r0128}29 30define i1 @test_v2i8_load(ptr %a) {31; CHECK-LABEL: test_v2i8_load(32; CHECK: {33; CHECK-NEXT: .reg .pred %p<2>;34; CHECK-NEXT: .reg .b16 %rs<5>;35; CHECK-NEXT: .reg .b32 %r<2>;36; CHECK-NEXT: .reg .b64 %rd<2>;37; CHECK-EMPTY:38; CHECK-NEXT: // %bb.0:39; CHECK-NEXT: ld.param.b64 %rd1, [test_v2i8_load_param_0];40; CHECK-NEXT: ld.v2.b8 {%rs1, %rs2}, [%rd1];41; CHECK-NEXT: or.b16 %rs3, %rs1, %rs2;42; CHECK-NEXT: and.b16 %rs4, %rs3, 255;43; CHECK-NEXT: setp.eq.b16 %p1, %rs4, 0;44; CHECK-NEXT: selp.b32 %r1, -1, 0, %p1;45; CHECK-NEXT: st.param.b32 [func_retval0], %r1;46; CHECK-NEXT: ret;47 %v = load <2 x i8>, ptr %a, align 448 %r0 = extractelement <2 x i8> %v, i64 049 %r1 = extractelement <2 x i8> %v, i64 150 %icmp = icmp eq i8 %r0, 051 %icmp3 = icmp eq i8 %r1, 052 %select = select i1 %icmp, i1 %icmp3, i1 false53 ret i1 %select54}55define i16 @test_v4i8(i32 %a) {56; CHECK-LABEL: test_v4i8(57; CHECK: {58; CHECK-NEXT: .reg .b16 %rs<8>;59; CHECK-NEXT: .reg .b32 %r<6>;60; CHECK-EMPTY:61; CHECK-NEXT: // %bb.0:62; CHECK-NEXT: ld.param.b32 %r1, [test_v4i8_param_0];63; CHECK-NEXT: cvt.s8.s32 %rs1, %r1;64; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x9991U;65; CHECK-NEXT: cvt.u16.u32 %rs2, %r2;66; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0xaaa2U;67; CHECK-NEXT: cvt.u16.u32 %rs3, %r3;68; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;69; CHECK-NEXT: cvt.u16.u32 %rs4, %r4;70; CHECK-NEXT: add.s16 %rs5, %rs1, %rs2;71; CHECK-NEXT: add.s16 %rs6, %rs3, %rs4;72; CHECK-NEXT: add.s16 %rs7, %rs5, %rs6;73; CHECK-NEXT: cvt.u32.u16 %r5, %rs7;74; CHECK-NEXT: st.param.b32 [func_retval0], %r5;75; CHECK-NEXT: ret;76 %v = bitcast i32 %a to <4 x i8>77 %r0 = extractelement <4 x i8> %v, i64 078 %r1 = extractelement <4 x i8> %v, i64 179 %r2 = extractelement <4 x i8> %v, i64 280 %r3 = extractelement <4 x i8> %v, i64 381 %r0i = sext i8 %r0 to i1682 %r1i = sext i8 %r1 to i1683 %r2i = sext i8 %r2 to i1684 %r3i = sext i8 %r3 to i1685 %r01 = add i16 %r0i, %r1i86 %r23 = add i16 %r2i, %r3i87 %r = add i16 %r01, %r2388 ret i16 %r89}90 91define i32 @test_v4i8_s32(i32 %a) {92; CHECK-LABEL: test_v4i8_s32(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<9>;95; CHECK-EMPTY:96; CHECK-NEXT: // %bb.0:97; CHECK-NEXT: ld.param.b32 %r1, [test_v4i8_s32_param_0];98; CHECK-NEXT: cvt.s32.s8 %r2, %r1;99; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x9991U;100; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xaaa2U;101; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0xbbb3U;102; CHECK-NEXT: add.s32 %r6, %r2, %r3;103; CHECK-NEXT: add.s32 %r7, %r4, %r5;104; CHECK-NEXT: add.s32 %r8, %r6, %r7;105; CHECK-NEXT: st.param.b32 [func_retval0], %r8;106; CHECK-NEXT: ret;107 %v = bitcast i32 %a to <4 x i8>108 %r0 = extractelement <4 x i8> %v, i64 0109 %r1 = extractelement <4 x i8> %v, i64 1110 %r2 = extractelement <4 x i8> %v, i64 2111 %r3 = extractelement <4 x i8> %v, i64 3112 %r0i = sext i8 %r0 to i32113 %r1i = sext i8 %r1 to i32114 %r2i = sext i8 %r2 to i32115 %r3i = sext i8 %r3 to i32116 %r01 = add i32 %r0i, %r1i117 %r23 = add i32 %r2i, %r3i118 %r = add i32 %r01, %r23119 ret i32 %r120}121 122define i32 @test_v4i8_u32(i32 %a) {123; CHECK-LABEL: test_v4i8_u32(124; CHECK: {125; CHECK-NEXT: .reg .b32 %r<9>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0:128; CHECK-NEXT: ld.param.b32 %r1, [test_v4i8_u32_param_0];129; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7771U;130; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U;131; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;132; CHECK-NEXT: and.b32 %r5, %r1, 255;133; CHECK-NEXT: add.s32 %r6, %r5, %r2;134; CHECK-NEXT: add.s32 %r7, %r3, %r4;135; CHECK-NEXT: add.s32 %r8, %r6, %r7;136; CHECK-NEXT: st.param.b32 [func_retval0], %r8;137; CHECK-NEXT: ret;138 %v = bitcast i32 %a to <4 x i8>139 %r0 = extractelement <4 x i8> %v, i64 0140 %r1 = extractelement <4 x i8> %v, i64 1141 %r2 = extractelement <4 x i8> %v, i64 2142 %r3 = extractelement <4 x i8> %v, i64 3143 %r0i = zext i8 %r0 to i32144 %r1i = zext i8 %r1 to i32145 %r2i = zext i8 %r2 to i32146 %r3i = zext i8 %r3 to i32147 %r01 = add i32 %r0i, %r1i148 %r23 = add i32 %r2i, %r3i149 %r = add i32 %r01, %r23150 ret i32 %r151}152 153 154 155define i16 @test_v8i8(i64 %a) {156; CHECK-LABEL: test_v8i8(157; CHECK: {158; CHECK-NEXT: .reg .b16 %rs<16>;159; CHECK-NEXT: .reg .b32 %r<10>;160; CHECK-EMPTY:161; CHECK-NEXT: // %bb.0:162; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_v8i8_param_0];163; CHECK-NEXT: cvt.s8.s32 %rs1, %r1;164; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x9991U;165; CHECK-NEXT: cvt.u16.u32 %rs2, %r3;166; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xaaa2U;167; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;168; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0xbbb3U;169; CHECK-NEXT: cvt.u16.u32 %rs4, %r5;170; CHECK-NEXT: cvt.s8.s32 %rs5, %r2;171; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x9991U;172; CHECK-NEXT: cvt.u16.u32 %rs6, %r6;173; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0xaaa2U;174; CHECK-NEXT: cvt.u16.u32 %rs7, %r7;175; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0xbbb3U;176; CHECK-NEXT: cvt.u16.u32 %rs8, %r8;177; CHECK-NEXT: add.s16 %rs9, %rs1, %rs2;178; CHECK-NEXT: add.s16 %rs10, %rs3, %rs4;179; CHECK-NEXT: add.s16 %rs11, %rs5, %rs6;180; CHECK-NEXT: add.s16 %rs12, %rs7, %rs8;181; CHECK-NEXT: add.s16 %rs13, %rs9, %rs10;182; CHECK-NEXT: add.s16 %rs14, %rs11, %rs12;183; CHECK-NEXT: add.s16 %rs15, %rs13, %rs14;184; CHECK-NEXT: cvt.u32.u16 %r9, %rs15;185; CHECK-NEXT: st.param.b32 [func_retval0], %r9;186; CHECK-NEXT: ret;187 %v = bitcast i64 %a to <8 x i8>188 %r0 = extractelement <8 x i8> %v, i64 0189 %r1 = extractelement <8 x i8> %v, i64 1190 %r2 = extractelement <8 x i8> %v, i64 2191 %r3 = extractelement <8 x i8> %v, i64 3192 %r4 = extractelement <8 x i8> %v, i64 4193 %r5 = extractelement <8 x i8> %v, i64 5194 %r6 = extractelement <8 x i8> %v, i64 6195 %r7 = extractelement <8 x i8> %v, i64 7196 %r0i = sext i8 %r0 to i16197 %r1i = sext i8 %r1 to i16198 %r2i = sext i8 %r2 to i16199 %r3i = sext i8 %r3 to i16200 %r4i = sext i8 %r4 to i16201 %r5i = sext i8 %r5 to i16202 %r6i = sext i8 %r6 to i16203 %r7i = sext i8 %r7 to i16204 %r01 = add i16 %r0i, %r1i205 %r23 = add i16 %r2i, %r3i206 %r45 = add i16 %r4i, %r5i207 %r67 = add i16 %r6i, %r7i208 %r0123 = add i16 %r01, %r23209 %r4567 = add i16 %r45, %r67210 %r = add i16 %r0123, %r4567211 ret i16 %r212}213