; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
target triple = "nvptx-nvidia-cuda"
define <6 x half> @half6() {
; CHECK-LABEL: half6(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b16 %rs1, 0x0000;
; CHECK-NEXT: st.param.v4.b16 [func_retval0+0], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v2.b16 [func_retval0+8], {%rs1, %rs1};
; CHECK-NEXT: ret;
ret <6 x half> zeroinitializer
}
define <10 x half> @half10() {
; CHECK-LABEL: half10(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b16 %rs1, 0x0000;
; CHECK-NEXT: st.param.v4.b16 [func_retval0+0], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b16 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v2.b16 [func_retval0+16], {%rs1, %rs1};
; CHECK-NEXT: ret;
ret <10 x half> zeroinitializer
}
define <12 x i8> @byte12() {
; CHECK-LABEL: byte12(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.u16 %rs1, 0;
; CHECK-NEXT: st.param.v4.b8 [func_retval0+0], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: ret;
ret <12 x i8> zeroinitializer
}
define <20 x i8> @byte20() {
; CHECK-LABEL: byte20(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.u16 %rs1, 0;
; CHECK-NEXT: st.param.v4.b8 [func_retval0+0], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: st.param.v4.b8 [func_retval0+16], {%rs1, %rs1, %rs1, %rs1};
; CHECK-NEXT: ret;
ret <20 x i8> zeroinitializer
}