1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 | FileCheck %s 3; RUN: %if ptxas-11.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %} 4 5define <2 x i16> @sext_setcc_v2i1_to_v2i16(ptr %p) { 6; CHECK-LABEL: sext_setcc_v2i1_to_v2i16( 7; CHECK: { 8; CHECK-NEXT: .reg .pred %p<3>; 9; CHECK-NEXT: .reg .b16 %rs<5>; 10; CHECK-NEXT: .reg .b32 %r<3>; 11; CHECK-NEXT: .reg .b64 %rd<2>; 12; CHECK-EMPTY: 13; CHECK-NEXT: // %bb.0: // %entry 14; CHECK-NEXT: ld.param.u64 %rd1, [sext_setcc_v2i1_to_v2i16_param_0]; 15; CHECK-NEXT: ld.u32 %r1, [%rd1]; 16; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1; 17; CHECK-NEXT: setp.eq.s16 %p1, %rs1, 0; 18; CHECK-NEXT: setp.eq.s16 %p2, %rs2, 0; 19; CHECK-NEXT: selp.s16 %rs3, -1, 0, %p2; 20; CHECK-NEXT: selp.s16 %rs4, -1, 0, %p1; 21; CHECK-NEXT: mov.b32 %r2, {%rs4, %rs3}; 22; CHECK-NEXT: st.param.b32 [func_retval0], %r2; 23; CHECK-NEXT: ret; 24entry: 25 %v = load <2 x i16>, ptr %p, align 4 26 %cmp = icmp eq <2 x i16> %v, zeroinitializer 27 %sext = sext <2 x i1> %cmp to <2 x i16> 28 ret <2 x i16> %sext 29} 30 31define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { 32; CHECK-LABEL: sext_setcc_v4i1_to_v4i8( 33; CHECK: { 34; CHECK-NEXT: .reg .pred %p<5>; 35; CHECK-NEXT: .reg .b16 %rs<9>; 36; CHECK-NEXT: .reg .b32 %r<13>; 37; CHECK-NEXT: .reg .b64 %rd<2>; 38; CHECK-EMPTY: 39; CHECK-NEXT: // %bb.0: // %entry 40; CHECK-NEXT: ld.param.u64 %rd1, [sext_setcc_v4i1_to_v4i8_param_0]; 41; CHECK-NEXT: ld.u32 %r1, [%rd1]; 42; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; 43; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; 44; CHECK-NEXT: and.b16 %rs2, %rs1, 255; 45; CHECK-NEXT: setp.eq.s16 %p1, %rs2, 0; 46; CHECK-NEXT: bfe.u32 %r3, %r1, 8, 8; 47; CHECK-NEXT: cvt.u16.u32 %rs3, %r3; 48; CHECK-NEXT: and.b16 %rs4, %rs3, 255; 49; CHECK-NEXT: setp.eq.s16 %p2, %rs4, 0; 50; CHECK-NEXT: bfe.u32 %r4, %r1, 16, 8; 51; CHECK-NEXT: cvt.u16.u32 %rs5, %r4; 52; CHECK-NEXT: and.b16 %rs6, %rs5, 255; 53; CHECK-NEXT: setp.eq.s16 %p3, %rs6, 0; 54; CHECK-NEXT: bfe.u32 %r5, %r1, 24, 8; 55; CHECK-NEXT: cvt.u16.u32 %rs7, %r5; 56; CHECK-NEXT: and.b16 %rs8, %rs7, 255; 57; CHECK-NEXT: setp.eq.s16 %p4, %rs8, 0; 58; CHECK-NEXT: selp.s32 %r6, -1, 0, %p4; 59; CHECK-NEXT: selp.s32 %r7, -1, 0, %p3; 60; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U; 61; CHECK-NEXT: selp.s32 %r9, -1, 0, %p2; 62; CHECK-NEXT: selp.s32 %r10, -1, 0, %p1; 63; CHECK-NEXT: prmt.b32 %r11, %r10, %r9, 0x3340U; 64; CHECK-NEXT: prmt.b32 %r12, %r11, %r8, 0x5410U; 65; CHECK-NEXT: st.param.b32 [func_retval0], %r12; 66; CHECK-NEXT: ret; 67entry: 68 %v = load <4 x i8>, ptr %p, align 4 69 %cmp = icmp eq <4 x i8> %v, zeroinitializer 70 %sext = sext <4 x i1> %cmp to <4 x i8> 71 ret <4 x i8> %sext 72} 73