llvm/test/CodeGen/NVPTX/mulwide.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | FileCheck %s --check-prefixes=CHECK,OPT
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefixes=CHECK,NOOPT
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O3 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}

define i32 @mulwide16(i16 %a, i16 %b) {
; OPT-LABEL: mulwide16(
; OPT:       {
; OPT-NEXT:    .reg .b16 %rs<3>;
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b16 %rs1, [mulwide16_param_0];
; OPT-NEXT:    ld.param.b16 %rs2, [mulwide16_param_1];
; OPT-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; OPT-NEXT:    st.param.b32 [func_retval0], %r1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide16(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s16 %r1, [mulwide16_param_0];
; NOOPT-NEXT:    ld.param.s16 %r2, [mulwide16_param_1];
; NOOPT-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r3;
; NOOPT-NEXT:    ret;
  %val0 = sext i16 %a to i32
  %val1 = sext i16 %b to i32
  %val2 = mul i32 %val0, %val1
  ret i32 %val2
}

define i32 @mulwideu16(i16 %a, i16 %b) {
; OPT-LABEL: mulwideu16(
; OPT:       {
; OPT-NEXT:    .reg .b16 %rs<3>;
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b16 %rs1, [mulwideu16_param_0];
; OPT-NEXT:    ld.param.b16 %rs2, [mulwideu16_param_1];
; OPT-NEXT:    mul.wide.u16 %r1, %rs1, %rs2;
; OPT-NEXT:    st.param.b32 [func_retval0], %r1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwideu16(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b16 %r1, [mulwideu16_param_0];
; NOOPT-NEXT:    ld.param.b16 %r2, [mulwideu16_param_1];
; NOOPT-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r3;
; NOOPT-NEXT:    ret;
  %val0 = zext i16 %a to i32
  %val1 = zext i16 %b to i32
  %val2 = mul i32 %val0, %val1
  ret i32 %val2
}

define i32 @mulwide8(i8 %a, i8 %b) {
; OPT-LABEL: mulwide8(
; OPT:       {
; OPT-NEXT:    .reg .b16 %rs<3>;
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.s8 %rs1, [mulwide8_param_0];
; OPT-NEXT:    ld.param.s8 %rs2, [mulwide8_param_1];
; OPT-NEXT:    mul.wide.s16 %r1, %rs1, %rs2;
; OPT-NEXT:    st.param.b32 [func_retval0], %r1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide8(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s8 %r1, [mulwide8_param_0];
; NOOPT-NEXT:    ld.param.s8 %r2, [mulwide8_param_1];
; NOOPT-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r3;
; NOOPT-NEXT:    ret;
  %val0 = sext i8 %a to i32
  %val1 = sext i8 %b to i32
  %val2 = mul i32 %val0, %val1
  ret i32 %val2
}

define i32 @mulwideu8(i8 %a, i8 %b) {
; OPT-LABEL: mulwideu8(
; OPT:       {
; OPT-NEXT:    .reg .b16 %rs<3>;
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b8 %rs1, [mulwideu8_param_0];
; OPT-NEXT:    ld.param.b8 %rs2, [mulwideu8_param_1];
; OPT-NEXT:    mul.wide.u16 %r1, %rs1, %rs2;
; OPT-NEXT:    st.param.b32 [func_retval0], %r1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwideu8(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %r1, [mulwideu8_param_0];
; NOOPT-NEXT:    ld.param.b8 %r2, [mulwideu8_param_1];
; NOOPT-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r3;
; NOOPT-NEXT:    ret;
  %val0 = zext i8 %a to i32
  %val1 = zext i8 %b to i32
  %val2 = mul i32 %val0, %val1
  ret i32 %val2
}

define i64 @mulwide32(i32 %a, i32 %b) {
; OPT-LABEL: mulwide32(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide32_param_0];
; OPT-NEXT:    ld.param.b32 %r2, [mulwide32_param_1];
; OPT-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide32(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s32 %rd1, [mulwide32_param_0];
; NOOPT-NEXT:    ld.param.s32 %rd2, [mulwide32_param_1];
; NOOPT-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT:    ret;
  %val0 = sext i32 %a to i64
  %val1 = sext i32 %b to i64
  %val2 = mul i64 %val0, %val1
  ret i64 %val2
}

define i64 @mulwideu32(i32 %a, i32 %b) {
; OPT-LABEL: mulwideu32(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwideu32_param_0];
; OPT-NEXT:    ld.param.b32 %r2, [mulwideu32_param_1];
; OPT-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwideu32(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %rd1, [mulwideu32_param_0];
; NOOPT-NEXT:    ld.param.b32 %rd2, [mulwideu32_param_1];
; NOOPT-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT:    ret;
  %val0 = zext i32 %a to i64
  %val1 = zext i32 %b to i64
  %val2 = mul i64 %val0, %val1
  ret i64 %val2
}

define i64 @mulwideu7(i7 %a, i7 %b) {
; OPT-LABEL: mulwideu7(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<5>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b8 %r1, [mulwideu7_param_1];
; OPT-NEXT:    and.b32 %r2, %r1, 127;
; OPT-NEXT:    ld.param.b8 %r3, [mulwideu7_param_0];
; OPT-NEXT:    and.b32 %r4, %r3, 127;
; OPT-NEXT:    mul.wide.u32 %rd1, %r4, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwideu7(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<6>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rd1, [mulwideu7_param_0];
; NOOPT-NEXT:    and.b64 %rd2, %rd1, 127;
; NOOPT-NEXT:    ld.param.b8 %rd3, [mulwideu7_param_1];
; NOOPT-NEXT:    and.b64 %rd4, %rd3, 127;
; NOOPT-NEXT:    mul.lo.s64 %rd5, %rd2, %rd4;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd5;
; NOOPT-NEXT:    ret;
  %val0 = zext i7 %a to i64
  %val1 = zext i7 %b to i64
  %val2 = mul i64 %val0, %val1
  ret i64 %val2
}

define i64 @mulwides7(i7 %a, i7 %b) {
; OPT-LABEL: mulwides7(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<5>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b8 %r1, [mulwides7_param_0];
; OPT-NEXT:    ld.param.b8 %r2, [mulwides7_param_1];
; OPT-NEXT:    bfe.s32 %r3, %r2, 0, 7;
; OPT-NEXT:    bfe.s32 %r4, %r1, 0, 7;
; OPT-NEXT:    mul.wide.s32 %rd1, %r4, %r3;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwides7(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<8>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rd1, [mulwides7_param_1];
; NOOPT-NEXT:    ld.param.b8 %rd2, [mulwides7_param_0];
; NOOPT-NEXT:    shl.b64 %rd3, %rd2, 57;
; NOOPT-NEXT:    shr.s64 %rd4, %rd3, 57;
; NOOPT-NEXT:    shl.b64 %rd5, %rd1, 57;
; NOOPT-NEXT:    shr.s64 %rd6, %rd5, 57;
; NOOPT-NEXT:    mul.lo.s64 %rd7, %rd4, %rd6;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd7;
; NOOPT-NEXT:    ret;
  %val0 = sext i7 %a to i64
  %val1 = sext i7 %b to i64
  %val2 = mul i64 %val0, %val1
  ret i64 %val2
}

define i64 @shl30(i32 %a) {
; OPT-LABEL: shl30(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [shl30_param_0];
; OPT-NEXT:    mul.wide.s32 %rd1, %r1, 1073741824;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: shl30(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s32 %rd1, [shl30_param_0];
; NOOPT-NEXT:    shl.b64 %rd2, %rd1, 30;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT:    ret;
  %conv = sext i32 %a to i64
  %shl = shl i64 %conv, 30
  ret i64 %shl
}

define i64 @shl31(i32 %a) {
; CHECK-LABEL: shl31(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.s32 %rd1, [shl31_param_0];
; CHECK-NEXT:    shl.b64 %rd2, %rd1, 31;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT:    ret;
  %conv = sext i32 %a to i64
  %shl = shl i64 %conv, 31
  ret i64 %shl
}


define i32 @mulwide_sext_shl_topbit_i16(i16 %x) {
; CHECK-LABEL: mulwide_sext_shl_topbit_i16(
; 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, [mulwide_sext_shl_topbit_i16_param_0];
; CHECK-NEXT:    shl.b16 %rs2, %rs1, 15;
; CHECK-NEXT:    cvt.s32.s16 %r1, %rs2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
  %s = shl nsw i16 %x, 15
  %e = sext i16 %s to i32
  ret i32 %e
}

define i64 @mulwide_sext_shl_topbit_i32(i32 %x) {
; CHECK-LABEL: mulwide_sext_shl_topbit_i32(
; 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, [mulwide_sext_shl_topbit_i32_param_0];
; CHECK-NEXT:    shl.b32 %r2, %r1, 31;
; CHECK-NEXT:    cvt.s64.s32 %rd1, %r2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT:    ret;
  %s = shl nsw i32 %x, 31
  %e = sext i32 %s to i64
  ret i64 %e
}

; A folded trunc/zext still provides a 32-bit unsigned multiply operand.
define i64 @mulwide_u32_mask(i64 %a, i32 %b) {
; OPT-LABEL: mulwide_u32_mask(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_u32_mask_param_0];
; OPT-NEXT:    ld.param.b32 %r2, [mulwide_u32_mask_param_1];
; OPT-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_u32_mask(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %rd1, [mulwide_u32_mask_param_0];
; NOOPT-NEXT:    ld.param.b32 %rd2, [mulwide_u32_mask_param_1];
; NOOPT-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT:    ret;
  %x = and i64 %a, 4294967295
  %y = zext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

; This is the high-half extraction and multiply used by Philox.
define i64 @mulwide_u32_lshr(i64 %a) {
; OPT-LABEL: mulwide_u32_lshr(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_u32_lshr_param_0+4];
; OPT-NEXT:    mul.wide.u32 %rd1, %r1, -766435501;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_u32_lshr(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %rd1, [mulwide_u32_lshr_param_0+4];
; NOOPT-NEXT:    mul.lo.s64 %rd2, %rd1, 3528531795;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT:    ret;
  %x = lshr i64 %a, 32
  %p = mul i64 %x, 3528531795
  ret i64 %p
}

define i64 @mulwide_u32_xor(i64 %a, i64 %b) {
; OPT-LABEL: mulwide_u32_xor(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-NEXT:    .reg .b64 %rd<5>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %rd1, [mulwide_u32_xor_param_0+4];
; OPT-NEXT:    ld.param.b64 %rd2, [mulwide_u32_xor_param_1];
; OPT-NEXT:    xor.b64 %rd3, %rd1, %rd2;
; OPT-NEXT:    cvt.u32.u64 %r1, %rd3;
; OPT-NEXT:    mul.wide.u32 %rd4, %r1, -845247145;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd4;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_u32_xor(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<5>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %rd1, [mulwide_u32_xor_param_0+4];
; NOOPT-NEXT:    ld.param.b32 %rd2, [mulwide_u32_xor_param_1];
; NOOPT-NEXT:    xor.b64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    mul.lo.s64 %rd4, %rd3, 3449720151;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd4;
; NOOPT-NEXT:    ret;
  %x = lshr i64 %a, 32
  %y = and i64 %b, 4294967295
  %z = xor i64 %x, %y
  %p = mul i64 %z, 3449720151
  ret i64 %p
}

define i64 @mulwide_s32_ashr(i64 %a, i32 %b) {
; OPT-LABEL: mulwide_s32_ashr(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_s32_ashr_param_0+4];
; OPT-NEXT:    ld.param.b32 %r2, [mulwide_s32_ashr_param_1];
; OPT-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_s32_ashr(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s32 %rd1, [mulwide_s32_ashr_param_0+4];
; NOOPT-NEXT:    ld.param.s32 %rd2, [mulwide_s32_ashr_param_1];
; NOOPT-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT:    ret;
  %x = ashr i64 %a, 32
  %y = sext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

define i64 @mulwide_s32_sext_inreg(i64 %a, i32 %b) {
; OPT-LABEL: mulwide_s32_sext_inreg(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_s32_sext_inreg_param_0];
; OPT-NEXT:    ld.param.b32 %r2, [mulwide_s32_sext_inreg_param_1];
; OPT-NEXT:    mul.wide.s32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_s32_sext_inreg(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s32 %rd1, [mulwide_s32_sext_inreg_param_0];
; NOOPT-NEXT:    ld.param.s32 %rd2, [mulwide_s32_sext_inreg_param_1];
; NOOPT-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd3;
; NOOPT-NEXT:    ret;
  %hi = shl i64 %a, 32
  %x = ashr i64 %hi, 32
  %y = sext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

; The nonnegative operand fits either signedness; use signed for the other.
define i64 @mulwide_s32_nonnegative(i32 %a, i32 %b) {
; OPT-LABEL: mulwide_s32_nonnegative(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<4>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_s32_nonnegative_param_0];
; OPT-NEXT:    and.b32 %r2, %r1, 2147483647;
; OPT-NEXT:    ld.param.b32 %r3, [mulwide_s32_nonnegative_param_1];
; OPT-NEXT:    mul.wide.s32 %rd1, %r2, %r3;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_s32_nonnegative(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<5>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %rd1, [mulwide_s32_nonnegative_param_0];
; NOOPT-NEXT:    ld.param.s32 %rd2, [mulwide_s32_nonnegative_param_1];
; NOOPT-NEXT:    and.b64 %rd3, %rd1, 2147483647;
; NOOPT-NEXT:    mul.lo.s64 %rd4, %rd3, %rd2;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd4;
; NOOPT-NEXT:    ret;
  %nonnegative = and i32 %a, 2147483647
  %x = zext i32 %nonnegative to i64
  %y = sext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

define i64 @mulwide_s32_negative_constant(i64 %a) {
; OPT-LABEL: mulwide_s32_negative_constant(
; OPT:       {
; OPT-NEXT:    .reg .b32 %r<2>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b32 %r1, [mulwide_s32_negative_constant_param_0+4];
; OPT-NEXT:    mul.wide.s32 %rd1, %r1, -766435501;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwide_s32_negative_constant(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.s32 %rd1, [mulwide_s32_negative_constant_param_0+4];
; NOOPT-NEXT:    mul.lo.s64 %rd2, %rd1, -766435501;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT:    ret;
  %x = ashr i64 %a, 32
  %p = mul i64 %x, -766435501
  ret i64 %p
}

; A plain shift of a small value should not acquire 16-bit conversions.
define i32 @shl_narrow_mask(i32 %a) {
; CHECK-LABEL: shl_narrow_mask(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [shl_narrow_mask_param_0];
; CHECK-NEXT:    and.b32 %r2, %r1, 3;
; CHECK-NEXT:    shl.b32 %r3, %r2, 3;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %x = and i32 %a, 3
  %p = shl i32 %x, 3
  ret i32 %p
}

; Implicit narrowing of i32 operands should keep ordinary 32-bit multiplies.
define i32 @mulwide_u16_mask(i32 %a, i16 %b) {
; CHECK-LABEL: mulwide_u16_mask(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %r1, [mulwide_u16_mask_param_0];
; CHECK-NEXT:    ld.param.b16 %r2, [mulwide_u16_mask_param_1];
; CHECK-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %x = and i32 %a, 65535
  %y = zext i16 %b to i32
  %p = mul i32 %x, %y
  ret i32 %p
}

define i32 @mulwide_s16_ashr(i32 %a, i16 %b) {
; CHECK-LABEL: mulwide_s16_ashr(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.s16 %r1, [mulwide_s16_ashr_param_0+2];
; CHECK-NEXT:    ld.param.s16 %r2, [mulwide_s16_ashr_param_1];
; CHECK-NEXT:    mul.lo.s32 %r3, %r1, %r2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %x = ashr i32 %a, 16
  %y = sext i16 %b to i32
  %p = mul i32 %x, %y
  ret i32 %p
}

; A full unsigned 32-bit value and a potentially negative signed one do not
; share a 32-bit interpretation.
define i64 @mulwide_reject_mixed_signedness(i32 %a, i32 %b) {
; CHECK-LABEL: mulwide_reject_mixed_signedness(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %rd1, [mulwide_reject_mixed_signedness_param_0];
; CHECK-NEXT:    ld.param.s32 %rd2, [mulwide_reject_mixed_signedness_param_1];
; CHECK-NEXT:    mul.lo.s64 %rd3, %rd1, %rd2;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT:    ret;
  %x = zext i32 %a to i64
  %y = sext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

define i64 @mulwide_reject_u33_mask(i64 %a, i32 %b) {
; CHECK-LABEL: mulwide_reject_u33_mask(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [mulwide_reject_u33_mask_param_0];
; CHECK-NEXT:    and.b64 %rd2, %rd1, 8589934591;
; CHECK-NEXT:    ld.param.b32 %rd3, [mulwide_reject_u33_mask_param_1];
; CHECK-NEXT:    mul.lo.s64 %rd4, %rd2, %rd3;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd4;
; CHECK-NEXT:    ret;
  %x = and i64 %a, 8589934591
  %y = zext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

define i64 @mulwide_reject_s33_ashr(i64 %a, i32 %b) {
; CHECK-LABEL: mulwide_reject_s33_ashr(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [mulwide_reject_s33_ashr_param_0];
; CHECK-NEXT:    shr.s64 %rd2, %rd1, 31;
; CHECK-NEXT:    ld.param.s32 %rd3, [mulwide_reject_s33_ashr_param_1];
; CHECK-NEXT:    mul.lo.s64 %rd4, %rd2, %rd3;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd4;
; CHECK-NEXT:    ret;
  %x = ashr i64 %a, 31
  %y = sext i32 %b to i64
  %p = mul i64 %x, %y
  ret i64 %p
}

define i64 @mulwide_reject_u33_constant(i64 %a) {
; CHECK-LABEL: mulwide_reject_u33_constant(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %rd1, [mulwide_reject_u33_constant_param_0];
; CHECK-NEXT:    mul.lo.s64 %rd2, %rd1, 5000000001;
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT:    ret;
  %x = and i64 %a, 4294967295
  %p = mul i64 %x, 5000000001
  ret i64 %p
}

define i32 @mulwide_reject_u17_mask(i32 %a, i16 %b) {
; CHECK-LABEL: mulwide_reject_u17_mask(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [mulwide_reject_u17_mask_param_0];
; CHECK-NEXT:    and.b32 %r2, %r1, 131071;
; CHECK-NEXT:    ld.param.b16 %r3, [mulwide_reject_u17_mask_param_1];
; CHECK-NEXT:    mul.lo.s32 %r4, %r2, %r3;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
; CHECK-NEXT:    ret;
  %x = and i32 %a, 131071
  %y = zext i16 %b to i32
  %p = mul i32 %x, %y
  ret i32 %p
}

define i32 @mulwide_reject_s17_ashr(i32 %a, i16 %b) {
; CHECK-LABEL: mulwide_reject_s17_ashr(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [mulwide_reject_s17_ashr_param_0];
; CHECK-NEXT:    shr.s32 %r2, %r1, 15;
; CHECK-NEXT:    ld.param.s16 %r3, [mulwide_reject_s17_ashr_param_1];
; CHECK-NEXT:    mul.lo.s32 %r4, %r2, %r3;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
; CHECK-NEXT:    ret;
  %x = ashr i32 %a, 15
  %y = sext i16 %b to i32
  %p = mul i32 %x, %y
  ret i32 %p
}

; Preserve paired low/high products for ptxas to share partial products.
; The low half fits a signed wide multiply, but the full unsigned product
; also needs the high half when %b is negative.
define void @mulwide_keep_unsigned_lohi(ptr %out, i32 %a, i32 %b) {
; CHECK-LABEL: mulwide_keep_unsigned_lohi(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<7>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [mulwide_keep_unsigned_lohi_param_0];
; CHECK-NEXT:    ld.param.b32 %rd2, [mulwide_keep_unsigned_lohi_param_1];
; CHECK-NEXT:    ld.param.s32 %rd3, [mulwide_keep_unsigned_lohi_param_2];
; CHECK-NEXT:    and.b64 %rd4, %rd2, 2147483647;
; CHECK-NEXT:    mul.hi.u64 %rd5, %rd4, %rd3;
; CHECK-NEXT:    mul.lo.s64 %rd6, %rd4, %rd3;
; CHECK-NEXT:    st.b64 [%rd1], %rd6;
; CHECK-NEXT:    st.b64 [%rd1+8], %rd5;
; CHECK-NEXT:    ret;
  %nonnegative = and i32 %a, 2147483647
  %x64 = zext i32 %nonnegative to i64
  %y64 = sext i32 %b to i64
  %x = zext i64 %x64 to i128
  %y = zext i64 %y64 to i128
  %p = mul i128 %x, %y
  %lo = trunc i128 %p to i64
  %shifted = lshr i128 %p, 64
  %hi = trunc i128 %shifted to i64
  store i64 %lo, ptr %out
  %out_hi = getelementptr i64, ptr %out, i64 1
  store i64 %hi, ptr %out_hi
  ret void
}

; Commuting the source product must preserve the same pair.
define void @mulwide_keep_unsigned_lohi_commuted(ptr %out, i32 %a, i32 %b) {
; CHECK-LABEL: mulwide_keep_unsigned_lohi_commuted(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<7>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [mulwide_keep_unsigned_lohi_commuted_param_0];
; CHECK-NEXT:    ld.param.b32 %rd2, [mulwide_keep_unsigned_lohi_commuted_param_1];
; CHECK-NEXT:    ld.param.s32 %rd3, [mulwide_keep_unsigned_lohi_commuted_param_2];
; CHECK-NEXT:    and.b64 %rd4, %rd2, 2147483647;
; CHECK-NEXT:    mul.hi.u64 %rd5, %rd3, %rd4;
; CHECK-NEXT:    mul.lo.s64 %rd6, %rd3, %rd4;
; CHECK-NEXT:    st.b64 [%rd1], %rd6;
; CHECK-NEXT:    st.b64 [%rd1+8], %rd5;
; CHECK-NEXT:    ret;
  %nonnegative = and i32 %a, 2147483647
  %x64 = zext i32 %nonnegative to i64
  %y64 = sext i32 %b to i64
  %x = zext i64 %x64 to i128
  %y = zext i64 %y64 to i128
  %p = mul i128 %y, %x
  %lo = trunc i128 %p to i64
  %shifted = lshr i128 %p, 64
  %hi = trunc i128 %shifted to i64
  store i64 %lo, ptr %out
  %out_hi = getelementptr i64, ptr %out, i64 1
  store i64 %hi, ptr %out_hi
  ret void
}