1274feef7SJustin Fargnoli; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2b279f6b0SFangrui Song; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s 3274feef7SJustin Fargnoli 4274feef7SJustin Fargnolitarget triple = "nvptx64-unknown-unknown" 5aa23e493Smmoadeli 6aa23e493Smmoadelidefine void @kernel_func(ptr %in.vec, ptr %out.vec0) nounwind { 7274feef7SJustin Fargnoli; CHECK-LABEL: kernel_func( 8274feef7SJustin Fargnoli; CHECK: { 9*932d9c13SDrew Kersnar; CHECK-NEXT: .reg .b32 %r<14>; 10274feef7SJustin Fargnoli; CHECK-EMPTY: 11274feef7SJustin Fargnoli; CHECK-NEXT: // %bb.0: 12274feef7SJustin Fargnoli; CHECK-NEXT: ld.param.u32 %r1, [kernel_func_param_0]; 13*932d9c13SDrew Kersnar; CHECK-NEXT: ld.v4.b32 {%r2, %r3, %r4, %r5}, [%r1]; 14*932d9c13SDrew Kersnar; CHECK-NEXT: ld.v4.b32 {%r6, %r7, %r8, %r9}, [%r1+16]; 15*932d9c13SDrew Kersnar; CHECK-NEXT: ld.param.u32 %r10, [kernel_func_param_1]; 16*932d9c13SDrew Kersnar; CHECK-NEXT: prmt.b32 %r11, %r6, %r8, 0x4000U; 17*932d9c13SDrew Kersnar; CHECK-NEXT: prmt.b32 %r12, %r2, %r4, 0x40U; 18*932d9c13SDrew Kersnar; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x7610U; 19*932d9c13SDrew Kersnar; CHECK-NEXT: st.u32 [%r10], %r13; 20274feef7SJustin Fargnoli; CHECK-NEXT: ret; 21aa23e493Smmoadeli %wide.vec = load <32 x i8>, ptr %in.vec, align 64 22aa23e493Smmoadeli %vec0 = shufflevector <32 x i8> %wide.vec, <32 x i8> undef, <4 x i32> <i32 0, i32 8, i32 16, i32 24> 23aa23e493Smmoadeli store <4 x i8> %vec0, ptr %out.vec0, align 64 24aa23e493Smmoadeli ret void 25aa23e493Smmoadeli} 26