llvm/test/CodeGen/NVPTX/combine-wide.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -O1 | FileCheck %s --check-prefixes=CHECK,O1
; RUN: llc < %s -O0 | FileCheck %s --check-prefixes=CHECK,O0

target triple = "nvptx64-nvidia-cuda"

define i64 @t1(i32 %a, i32 %b, i64 %c) {
;
; O1-LABEL: t1(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t1_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t1_param_1];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; O1-NEXT:    ld.param.b64 %rd2, [t1_param_2];
; O1-NEXT:    add.s64 %rd3, %rd2, %rd1;
; O1-NEXT:    st.param.b64 [func_retval0], %rd3;
; O1-NEXT:    ret;
;
; O0-LABEL: t1(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t1_param_2];
; O0-NEXT:    ld.param.b32 %r2, [t1_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t1_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd2, %r3;
; O0-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O0-NEXT:    st.param.b64 [func_retval0], %rd3;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, %b
  %sext = sext i32 %mul to i64
  %add = add i64 %c, %sext
  ret i64 %add
}

define i64 @t2(i32 %a, i32 %b, i64 %c) {
;
; O1-LABEL: t2(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t2_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t2_param_1];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; O1-NEXT:    ld.param.b64 %rd2, [t2_param_2];
; O1-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O1-NEXT:    st.param.b64 [func_retval0], %rd3;
; O1-NEXT:    ret;
;
; O0-LABEL: t2(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t2_param_2];
; O0-NEXT:    ld.param.b32 %r2, [t2_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t2_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd2, %r3;
; O0-NEXT:    add.s64 %rd3, %rd2, %rd1;
; O0-NEXT:    st.param.b64 [func_retval0], %rd3;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, %b
  %sext = sext i32 %mul to i64
  %add = add i64 %sext, %c
  ret i64 %add
}

define i64 @t3(i32 %a, i32 %b) {
;
; O1-LABEL: t3(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<3>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t3_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t3_param_1];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; O1-NEXT:    add.s64 %rd2, %rd1, 1;
; O1-NEXT:    st.param.b64 [func_retval0], %rd2;
; O1-NEXT:    ret;
;
; O0-LABEL: t3(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<3>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t3_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t3_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd1, %r3;
; O0-NEXT:    add.s64 %rd2, %rd1, 1;
; O0-NEXT:    st.param.b64 [func_retval0], %rd2;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, %b
  %sext = sext i32 %mul to i64
  %add = add i64 1, %sext
  ret i64 %add
}

define i64 @t4(i32 %a, i64 %c) {
;
; O1-LABEL: t4(
; O1:       {
; O1-NEXT:    .reg .b32 %r<2>;
; O1-NEXT:    .reg .b64 %rd<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t4_param_0];
; O1-NEXT:    ld.param.b64 %rd1, [t4_param_1];
; O1-NEXT:    mul.wide.s32 %rd2, %r1, 3;
; O1-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O1-NEXT:    st.param.b64 [func_retval0], %rd3;
; O1-NEXT:    ret;
;
; O0-LABEL: t4(
; O0:       {
; O0-NEXT:    .reg .b32 %r<3>;
; O0-NEXT:    .reg .b64 %rd<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t4_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t4_param_0];
; O0-NEXT:    mul.lo.s32 %r2, %r1, 3;
; O0-NEXT:    cvt.s64.s32 %rd2, %r2;
; O0-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O0-NEXT:    st.param.b64 [func_retval0], %rd3;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, 3
  %sext = sext i32 %mul to i64
  %add = add i64 %c, %sext
  ret i64 %add
}

define i64 @t4_1(i32 %a, i64 %c) {
;
; O1-LABEL: t4_1(
; O1:       {
; O1-NEXT:    .reg .b32 %r<2>;
; O1-NEXT:    .reg .b64 %rd<3>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t4_1_param_0];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, 3;
; O1-NEXT:    add.s64 %rd2, %rd1, 5;
; O1-NEXT:    st.param.b64 [func_retval0], %rd2;
; O1-NEXT:    ret;
;
; O0-LABEL: t4_1(
; O0:       {
; O0-NEXT:    .reg .b32 %r<3>;
; O0-NEXT:    .reg .b64 %rd<3>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t4_1_param_0];
; O0-NEXT:    mul.lo.s32 %r2, %r1, 3;
; O0-NEXT:    cvt.s64.s32 %rd1, %r2;
; O0-NEXT:    add.s64 %rd2, %rd1, 5;
; O0-NEXT:    st.param.b64 [func_retval0], %rd2;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, 3
  %sext = sext i32 %mul to i64
  %add = add i64 5, %sext
  ret i64 %add
}

define i64 @t5(i32 %a, i32 %b, i64 %c) {
;
; O1-LABEL: t5(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t5_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t5_param_1];
; O1-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; O1-NEXT:    ld.param.b64 %rd2, [t5_param_2];
; O1-NEXT:    add.s64 %rd3, %rd2, %rd1;
; O1-NEXT:    st.param.b64 [func_retval0], %rd3;
; O1-NEXT:    ret;
;
; O0-LABEL: t5(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t5_param_2];
; O0-NEXT:    ld.param.b32 %r2, [t5_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t5_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.u64.u32 %rd2, %r3;
; O0-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O0-NEXT:    st.param.b64 [func_retval0], %rd3;
; O0-NEXT:    ret;
  %mul = mul nuw i32 %a, %b
  %zext = zext i32 %mul to i64
  %add = add i64 %c, %zext
  ret i64 %add
}

define i64 @t6(i32 %a, i32 %b, i64 %c) {
;
; O1-LABEL: t6(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t6_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t6_param_1];
; O1-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; O1-NEXT:    ld.param.b64 %rd2, [t6_param_2];
; O1-NEXT:    add.s64 %rd3, %rd1, %rd2;
; O1-NEXT:    st.param.b64 [func_retval0], %rd3;
; O1-NEXT:    ret;
;
; O0-LABEL: t6(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t6_param_2];
; O0-NEXT:    ld.param.b32 %r2, [t6_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t6_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.u64.u32 %rd2, %r3;
; O0-NEXT:    add.s64 %rd3, %rd2, %rd1;
; O0-NEXT:    st.param.b64 [func_retval0], %rd3;
; O0-NEXT:    ret;
  %mul = mul nuw i32 %a, %b
  %zext = zext i32 %mul to i64
  %add = add i64 %zext, %c
  ret i64 %add
}

define i32 @t7(i16 %a, i16 %b) {
;
; O1-LABEL: t7(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t7_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t7_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.u32.u16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t7(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t7_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t7_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul i16 %a, %b
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t8(i16 %a, i16 %b) {
;
; O1-LABEL: t8(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t8_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t8_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.s32.s16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t8(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t8_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t8_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul i16 %a, %b
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t9(i32 %a, i32 %b) {
;
; O1-LABEL: t9(
; O1:       {
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t9_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t9_param_1];
; O1-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O1-NEXT:    cvt.u64.u32 %rd1, %r3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t9(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t9_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t9_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.u64.u32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul i32 %a, %b
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t10(i32 %a, i32 %b) {
;
; O1-LABEL: t10(
; O1:       {
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t10_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t10_param_1];
; O1-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O1-NEXT:    cvt.s64.s32 %rd1, %r3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t10(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t10_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t10_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul i32 %a, %b
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t11(i16 %a, i16 %b) {
;
; O1-LABEL: t11(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t11_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t11_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.u32.u16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t11(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t11_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t11_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t12(i16 %a, i16 %b) {
;
; O1-LABEL: t12(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t12_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t12_param_1];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t12(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t12_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t12_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t13(i32 %a, i32 %b) {
;
; O1-LABEL: t13(
; O1:       {
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t13_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t13_param_1];
; O1-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O1-NEXT:    cvt.u64.u32 %rd1, %r3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t13(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t13_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t13_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.u64.u32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, %b
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t14(i32 %a, i32 %b) {
;
; O1-LABEL: t14(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t14_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t14_param_1];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t14(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t14_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t14_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nsw i32 %a, %b
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t15(i16 %a, i16 %b) {
;
; O1-LABEL: t15(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t15_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t15_param_1];
; O1-NEXT:    mul.wide.u16 %r1, %rs1, %rs2;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t15(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t15_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t15_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t16(i16 %a, i16 %b) {
;
; O1-LABEL: t16(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t16_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t16_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.s32.s16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t16(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t16_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t16_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t17(i32 %a, i32 %b) {
;
; O1-LABEL: t17(
; O1:       {
; O1-NEXT:    .reg .b32 %r<3>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t17_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t17_param_1];
; O1-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t17(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t17_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t17_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.u64.u32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nuw i32 %a, %b
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t18(i32 %a, i32 %b) {
;
; O1-LABEL: t18(
; O1:       {
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t18_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t18_param_1];
; O1-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O1-NEXT:    cvt.s64.s32 %rd1, %r3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t18(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t18_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t18_param_0];
; O0-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nuw i32 %a, %b
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t19(i16 %a, i16 %b) {
;
; O1-LABEL: t19(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t19_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t19_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.u32.u16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t19(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t19_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t19_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul i16 %a, %b
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t20(i16 %a) {
;
; CHECK-LABEL: t20(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [t20_param_0];
; CHECK-NEXT:    shl.b16 %rs2, %rs1, 4;
; CHECK-NEXT:    cvt.s32.s16 %r1, %rs2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %mul = shl i16 %a, 4
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t21(i32 %a) {
;
; CHECK-LABEL: t21(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [t21_param_0];
; CHECK-NEXT:    shl.b32 %r2, %r1, 4;
; CHECK-NEXT:    cvt.u64.u32 %rd1, %r2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT:    ret;
  %mul = shl i32 %a, 4
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t22(i32 %a) {
;
; CHECK-LABEL: t22(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [t22_param_0];
; CHECK-NEXT:    shl.b32 %r2, %r1, 4;
; CHECK-NEXT:    cvt.s64.s32 %rd1, %r2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT:    ret;
  %mul = shl i32 %a, 4
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t23(i16 %a, i16 %b) {
;
; CHECK-LABEL: t23(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [t23_param_0];
; CHECK-NEXT:    shl.b16 %rs2, %rs1, 4;
; CHECK-NEXT:    cvt.u32.u16 %r1, %rs2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %mul = shl nsw i16 %a, 4
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t24(i16 %a, i16 %b) {
;
; O1-LABEL: t24(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<2>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t24_param_0];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, 16;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t24(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<3>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs1, [t24_param_0];
; O0-NEXT:    shl.b16 %rs2, %rs1, 4;
; O0-NEXT:    cvt.s32.s16 %r1, %rs2;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = shl nsw i16 %a, 4
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t25(i32 %a) {
;
; CHECK-LABEL: t25(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [t25_param_0];
; CHECK-NEXT:    shl.b32 %r2, %r1, 4;
; CHECK-NEXT:    cvt.u64.u32 %rd1, %r2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT:    ret;
  %mul = shl nsw i32 %a, 4
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t26(i32 %a) {
;
; O1-LABEL: t26(
; O1:       {
; O1-NEXT:    .reg .b32 %r<2>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t26_param_0];
; O1-NEXT:    mul.wide.s32 %rd1, %r1, 16;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t26(
; O0:       {
; O0-NEXT:    .reg .b32 %r<3>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t26_param_0];
; O0-NEXT:    shl.b32 %r2, %r1, 4;
; O0-NEXT:    cvt.s64.s32 %rd1, %r2;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = shl nsw i32 %a, 4
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t27(i16 %a, i16 %b) {
;
; O1-LABEL: t27(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<2>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t27_param_0];
; O1-NEXT:    mul.wide.u16 %r1, %rs1, 16;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t27(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<3>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs1, [t27_param_0];
; O0-NEXT:    shl.b16 %rs2, %rs1, 4;
; O0-NEXT:    cvt.u32.u16 %r1, %rs2;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = shl nuw i16 %a, 4
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t28(i16 %a, i16 %b) {
;
; CHECK-LABEL: t28(
; CHECK:       {
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %rs1, [t28_param_0];
; CHECK-NEXT:    shl.b16 %rs2, %rs1, 4;
; CHECK-NEXT:    cvt.s32.s16 %r1, %rs2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %mul = shl nuw i16 %a, 4
  %sext = sext i16 %mul to i32
  ret i32 %sext
}

define i64 @t29(i32 %a) {
;
; O1-LABEL: t29(
; O1:       {
; O1-NEXT:    .reg .b32 %r<2>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t29_param_0];
; O1-NEXT:    mul.wide.u32 %rd1, %r1, 16;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t29(
; O0:       {
; O0-NEXT:    .reg .b32 %r<3>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t29_param_0];
; O0-NEXT:    shl.b32 %r2, %r1, 4;
; O0-NEXT:    cvt.u64.u32 %rd1, %r2;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = shl nuw i32 %a, 4
  %zext = zext i32 %mul to i64
  ret i64 %zext
}

define i64 @t30(i32 %a) {
;
; CHECK-LABEL: t30(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [t30_param_0];
; CHECK-NEXT:    shl.b32 %r2, %r1, 4;
; CHECK-NEXT:    cvt.s64.s32 %rd1, %r2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT:    ret;
  %mul = shl nuw i32 %a, 4
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i64 @t31(i32 %a, i32 %b) {
;
; O1-LABEL: t31(
; O1:       {
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b32 %r1, [t31_param_0];
; O1-NEXT:    ld.param.b32 %r2, [t31_param_1];
; O1-NEXT:    shl.b32 %r3, %r1, %r2;
; O1-NEXT:    cvt.s64.s32 %rd1, %r3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t31(
; O0:       {
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r2, [t31_param_1];
; O0-NEXT:    ld.param.b32 %r1, [t31_param_0];
; O0-NEXT:    shl.b32 %r3, %r1, %r2;
; O0-NEXT:    cvt.s64.s32 %rd1, %r3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = shl nuw i32 %a, %b
  %sext = sext i32 %mul to i64
  ret i64 %sext
}

define i32 @t32(i16 %a, i16 %b, i32 %c) {
;
; O1-LABEL: t32(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t32_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t32_param_1];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; O1-NEXT:    ld.param.b32 %r2, [t32_param_2];
; O1-NEXT:    add.s32 %r3, %r2, %r1;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t32(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t32_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t32_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t32_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r2, %rs3;
; O0-NEXT:    add.s32 %r3, %r1, %r2;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %sext = sext i16 %mul to i32
  %add = add i32 %c, %sext
  ret i32 %add
}

define i32 @t33(i16 %a, i16 %b, i32 %c) {
;
; O1-LABEL: t33(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t33_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t33_param_1];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; O1-NEXT:    ld.param.b32 %r2, [t33_param_2];
; O1-NEXT:    add.s32 %r3, %r2, %r1;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t33(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t33_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t33_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t33_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r2, %rs3;
; O0-NEXT:    add.s32 %r3, %r1, %r2;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %sext = sext i16 %mul to i32
  %add = add i32 %c, %sext
  ret i32 %add
}

define i32 @t34(i16 %a, i16 %b) {
;
; O1-LABEL: t34(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<3>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t34_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t34_param_1];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; O1-NEXT:    add.s32 %r2, %r1, 1;
; O1-NEXT:    st.param.b32 [func_retval0], %r2;
; O1-NEXT:    ret;
;
; O0-LABEL: t34(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<3>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t34_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t34_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s32.s16 %r1, %rs3;
; O0-NEXT:    add.s32 %r2, %r1, 1;
; O0-NEXT:    st.param.b32 [func_retval0], %r2;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %sext = sext i16 %mul to i32
  %add = add i32 1, %sext
  ret i32 %add
}

define i32 @t35(i16 %a, i32 %c) {
;
; O1-LABEL: t35(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<2>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t35_param_0];
; O1-NEXT:    ld.param.b32 %r1, [t35_param_1];
; O1-NEXT:    mul.wide.s16 %r2, %rs1, 3;
; O1-NEXT:    add.s32 %r3, %r1, %r2;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t35(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<3>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t35_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t35_param_0];
; O0-NEXT:    mul.lo.s16 %rs2, %rs1, 3;
; O0-NEXT:    cvt.s32.s16 %r2, %rs2;
; O0-NEXT:    add.s32 %r3, %r1, %r2;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, 3
  %sext = sext i16 %mul to i32
  %add = add i32 %c, %sext
  ret i32 %add
}

define i32 @t36(i16 %a, i32 %c) {
;
; O1-LABEL: t36(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<2>;
; O1-NEXT:    .reg .b32 %r<3>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t36_param_0];
; O1-NEXT:    mul.wide.s16 %r1, %rs1, 3;
; O1-NEXT:    add.s32 %r2, %r1, 5;
; O1-NEXT:    st.param.b32 [func_retval0], %r2;
; O1-NEXT:    ret;
;
; O0-LABEL: t36(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<3>;
; O0-NEXT:    .reg .b32 %r<3>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs1, [t36_param_0];
; O0-NEXT:    mul.lo.s16 %rs2, %rs1, 3;
; O0-NEXT:    cvt.s32.s16 %r1, %rs2;
; O0-NEXT:    add.s32 %r2, %r1, 5;
; O0-NEXT:    st.param.b32 [func_retval0], %r2;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, 3
  %sext = sext i16 %mul to i32
  %add = add i32 5, %sext
  ret i32 %add
}

define i32 @t37(i16 %a, i16 %b, i32 %c) {
;
; O1-LABEL: t37(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t37_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t37_param_1];
; O1-NEXT:    mul.wide.u16 %r1, %rs1, %rs2;
; O1-NEXT:    ld.param.b32 %r2, [t37_param_2];
; O1-NEXT:    add.s32 %r3, %r2, %r1;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t37(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t37_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t37_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t37_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r2, %rs3;
; O0-NEXT:    add.s32 %r3, %r1, %r2;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  %zext = zext i16 %mul to i32
  %add = add i32 %c, %zext
  ret i32 %add
}

define i32 @t38(i16 %a, i16 %b, i32 %c) {
;
; O1-LABEL: t38(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<3>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t38_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t38_param_1];
; O1-NEXT:    mul.wide.u16 %r1, %rs1, %rs2;
; O1-NEXT:    ld.param.b32 %r2, [t38_param_2];
; O1-NEXT:    add.s32 %r3, %r1, %r2;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t38(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b32 %r1, [t38_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t38_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t38_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u32.u16 %r2, %rs3;
; O0-NEXT:    add.s32 %r3, %r2, %r1;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  %zext = zext i16 %mul to i32
  %add = add i32 %zext, %c
  ret i32 %add
}

define i64 @t39(i16 %a, i16 %b) {
; O1-LABEL: t39(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t39_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t39_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.u64.u16 %rd1, %rs3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t39(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t39_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t39_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u64.u16 %rd1, %rs3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul i16 %a, %b
  %zext = zext i16 %mul to i64
  ret i64 %zext
}

define i64 @t40(i16 %a, i16 %b) {
; O1-LABEL: t40(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t40_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t40_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.u64.u16 %rd1, %rs3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t40(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t40_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t40_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.u64.u16 %rd1, %rs3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  %zext = zext i16 %mul to i64
  ret i64 %zext
}

define i64 @t41(i16 %a, i16 %b) {
; O1-LABEL: t41(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t41_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t41_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    cvt.s64.s16 %rd1, %rs3;
; O1-NEXT:    st.param.b64 [func_retval0], %rd1;
; O1-NEXT:    ret;
;
; O0-LABEL: t41(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b16 %rs2, [t41_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t41_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    cvt.s64.s16 %rd1, %rs3;
; O0-NEXT:    st.param.b64 [func_retval0], %rd1;
; O0-NEXT:    ret;
  %mul = mul nsw i16 %a, %b
  %sext = sext i16 %mul to i64
  ret i64 %sext
}

define i32 @t42(i16 %a, i16 %b, ptr %ptr) {
; O1-LABEL: t42(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<2>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t42_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t42_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    ld.param.b64 %rd1, [t42_param_2];
; O1-NEXT:    st.b16 [%rd1], %rs3;
; O1-NEXT:    cvt.u32.u16 %r1, %rs3;
; O1-NEXT:    st.param.b32 [func_retval0], %r1;
; O1-NEXT:    ret;
;
; O0-LABEL: t42(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<2>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t42_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t42_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t42_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    st.b16 [%rd1], %rs3;
; O0-NEXT:    cvt.u32.u16 %r1, %rs3;
; O0-NEXT:    st.param.b32 [func_retval0], %r1;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  store i16 %mul, ptr %ptr
  %zext = zext i16 %mul to i32
  ret i32 %zext
}

define i32 @t43(i16 %a, i16 %b, i32 %c, ptr %ptr) {
; O1-LABEL: t43(
; O1:       {
; O1-NEXT:    .reg .b16 %rs<4>;
; O1-NEXT:    .reg .b32 %r<4>;
; O1-NEXT:    .reg .b64 %rd<2>;
; O1-EMPTY:
; O1-NEXT:  // %bb.0:
; O1-NEXT:    ld.param.b16 %rs1, [t43_param_0];
; O1-NEXT:    ld.param.b16 %rs2, [t43_param_1];
; O1-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O1-NEXT:    ld.param.b64 %rd1, [t43_param_3];
; O1-NEXT:    st.b16 [%rd1], %rs3;
; O1-NEXT:    ld.param.b32 %r1, [t43_param_2];
; O1-NEXT:    cvt.u32.u16 %r2, %rs3;
; O1-NEXT:    add.s32 %r3, %r2, %r1;
; O1-NEXT:    st.param.b32 [func_retval0], %r3;
; O1-NEXT:    ret;
;
; O0-LABEL: t43(
; O0:       {
; O0-NEXT:    .reg .b16 %rs<4>;
; O0-NEXT:    .reg .b32 %r<4>;
; O0-NEXT:    .reg .b64 %rd<2>;
; O0-EMPTY:
; O0-NEXT:  // %bb.0:
; O0-NEXT:    ld.param.b64 %rd1, [t43_param_3];
; O0-NEXT:    ld.param.b32 %r1, [t43_param_2];
; O0-NEXT:    ld.param.b16 %rs2, [t43_param_1];
; O0-NEXT:    ld.param.b16 %rs1, [t43_param_0];
; O0-NEXT:    mul.lo.s16 %rs3, %rs1, %rs2;
; O0-NEXT:    st.b16 [%rd1], %rs3;
; O0-NEXT:    cvt.u32.u16 %r2, %rs3;
; O0-NEXT:    add.s32 %r3, %r2, %r1;
; O0-NEXT:    st.param.b32 [func_retval0], %r3;
; O0-NEXT:    ret;
  %mul = mul nuw i16 %a, %b
  store i16 %mul, ptr %ptr
  %zext = zext i16 %mul to i32
  %add = add i32 %zext, %c
  ret i32 %add
}