1cdaebf6fSJustin Fargnoli; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2b279f6b0SFangrui Song; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s 3cdaebf6fSJustin Fargnoli 4cdaebf6fSJustin Fargnolitarget triple = "nvptx-nvidia-cuda" 5cdaebf6fSJustin Fargnoli 6cdaebf6fSJustin Fargnolidefine <6 x half> @half6() { 7cdaebf6fSJustin Fargnoli; CHECK-LABEL: half6( 8cdaebf6fSJustin Fargnoli; CHECK: { 9cdaebf6fSJustin Fargnoli; CHECK-NEXT: .reg .b16 %rs<2>; 10cdaebf6fSJustin Fargnoli; CHECK-EMPTY: 11cdaebf6fSJustin Fargnoli; CHECK-NEXT: // %bb.0: 12cdaebf6fSJustin Fargnoli; CHECK-NEXT: mov.b16 %rs1, 0x0000; 130f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 14cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v2.b16 [func_retval0+8], {%rs1, %rs1}; 15cdaebf6fSJustin Fargnoli; CHECK-NEXT: ret; 16cdaebf6fSJustin Fargnoli ret <6 x half> zeroinitializer 17cdaebf6fSJustin Fargnoli} 18cdaebf6fSJustin Fargnoli 19cdaebf6fSJustin Fargnolidefine <10 x half> @half10() { 20cdaebf6fSJustin Fargnoli; CHECK-LABEL: half10( 21cdaebf6fSJustin Fargnoli; CHECK: { 22cdaebf6fSJustin Fargnoli; CHECK-NEXT: .reg .b16 %rs<2>; 23cdaebf6fSJustin Fargnoli; CHECK-EMPTY: 24cdaebf6fSJustin Fargnoli; CHECK-NEXT: // %bb.0: 25cdaebf6fSJustin Fargnoli; CHECK-NEXT: mov.b16 %rs1, 0x0000; 260f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 27cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b16 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 28cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v2.b16 [func_retval0+16], {%rs1, %rs1}; 29cdaebf6fSJustin Fargnoli; CHECK-NEXT: ret; 30cdaebf6fSJustin Fargnoli ret <10 x half> zeroinitializer 31cdaebf6fSJustin Fargnoli} 32cdaebf6fSJustin Fargnoli 33cdaebf6fSJustin Fargnolidefine <12 x i8> @byte12() { 34cdaebf6fSJustin Fargnoli; CHECK-LABEL: byte12( 35cdaebf6fSJustin Fargnoli; CHECK: { 36cdaebf6fSJustin Fargnoli; CHECK-NEXT: .reg .b16 %rs<2>; 37cdaebf6fSJustin Fargnoli; CHECK-EMPTY: 38cdaebf6fSJustin Fargnoli; CHECK-NEXT: // %bb.0: 39*310e7987SAlex MacLean; CHECK-NEXT: mov.b16 %rs1, 0; 400f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 41cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 42cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 43cdaebf6fSJustin Fargnoli; CHECK-NEXT: ret; 44cdaebf6fSJustin Fargnoli ret <12 x i8> zeroinitializer 45cdaebf6fSJustin Fargnoli} 46cdaebf6fSJustin Fargnoli 47cdaebf6fSJustin Fargnolidefine <20 x i8> @byte20() { 48cdaebf6fSJustin Fargnoli; CHECK-LABEL: byte20( 49cdaebf6fSJustin Fargnoli; CHECK: { 50cdaebf6fSJustin Fargnoli; CHECK-NEXT: .reg .b16 %rs<2>; 51cdaebf6fSJustin Fargnoli; CHECK-EMPTY: 52cdaebf6fSJustin Fargnoli; CHECK-NEXT: // %bb.0: 53*310e7987SAlex MacLean; CHECK-NEXT: mov.b16 %rs1, 0; 540f0a96b8SYoungsuk Kim; CHECK-NEXT: st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1}; 55cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1}; 56cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1}; 57cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1}; 58cdaebf6fSJustin Fargnoli; CHECK-NEXT: st.param.v4.b8 [func_retval0+16], {%rs1, %rs1, %rs1, %rs1}; 59cdaebf6fSJustin Fargnoli; CHECK-NEXT: ret; 60cdaebf6fSJustin Fargnoli ret <20 x i8> zeroinitializer 61cdaebf6fSJustin Fargnoli} 62