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-prefix=OPT
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=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 .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b16 %rs2, [mulwide16_param_1];
; NOOPT-NEXT:    ld.param.b16 %rs1, [mulwide16_param_0];
; NOOPT-NEXT:    cvt.s32.s16 %r1, %rs1;
; NOOPT-NEXT:    cvt.s32.s16 %r2, %rs2;
; 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 .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b32 %r<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b16 %rs2, [mulwideu16_param_1];
; NOOPT-NEXT:    ld.param.b16 %rs1, [mulwideu16_param_0];
; NOOPT-NEXT:    cvt.u32.u16 %r1, %rs1;
; NOOPT-NEXT:    cvt.u32.u16 %r2, %rs2;
; 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 .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b32 %r<6>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rs2, [mulwide8_param_1];
; NOOPT-NEXT:    ld.param.b8 %rs1, [mulwide8_param_0];
; NOOPT-NEXT:    cvt.u32.u16 %r1, %rs1;
; NOOPT-NEXT:    cvt.s32.s8 %r2, %r1;
; NOOPT-NEXT:    cvt.u32.u16 %r3, %rs2;
; NOOPT-NEXT:    cvt.s32.s8 %r4, %r3;
; NOOPT-NEXT:    mul.lo.s32 %r5, %r2, %r4;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r5;
; 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 .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b32 %r<6>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rs2, [mulwideu8_param_1];
; NOOPT-NEXT:    ld.param.b8 %rs1, [mulwideu8_param_0];
; NOOPT-NEXT:    cvt.u32.u16 %r1, %rs1;
; NOOPT-NEXT:    and.b32 %r2, %r1, 255;
; NOOPT-NEXT:    cvt.u32.u16 %r3, %rs2;
; NOOPT-NEXT:    and.b32 %r4, %r3, 255;
; NOOPT-NEXT:    mul.lo.s32 %r5, %r2, %r4;
; NOOPT-NEXT:    st.param.b32 [func_retval0], %r5;
; 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 .b32 %r<3>;
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %r2, [mulwide32_param_1];
; NOOPT-NEXT:    ld.param.b32 %r1, [mulwide32_param_0];
; NOOPT-NEXT:    cvt.s64.s32 %rd1, %r1;
; NOOPT-NEXT:    cvt.s64.s32 %rd2, %r2;
; 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 .b32 %r<3>;
; NOOPT-NEXT:    .reg .b64 %rd<4>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %r2, [mulwideu32_param_1];
; NOOPT-NEXT:    ld.param.b32 %r1, [mulwideu32_param_0];
; NOOPT-NEXT:    cvt.u64.u32 %rd1, %r1;
; NOOPT-NEXT:    cvt.u64.u32 %rd2, %r2;
; 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<3>;
; OPT-NEXT:    .reg .b64 %rd<2>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.b8 %r1, [mulwideu7_param_0];
; OPT-NEXT:    ld.param.b8 %r2, [mulwideu7_param_1];
; OPT-NEXT:    mul.wide.u32 %rd1, %r1, %r2;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwideu7(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b64 %rd<6>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rs2, [mulwideu7_param_1];
; NOOPT-NEXT:    ld.param.b8 %rs1, [mulwideu7_param_0];
; NOOPT-NEXT:    cvt.u64.u16 %rd1, %rs1;
; NOOPT-NEXT:    and.b64 %rd2, %rd1, 127;
; NOOPT-NEXT:    cvt.u64.u16 %rd3, %rs2;
; 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:    bfe.s32 %r2, %r1, 0, 7;
; OPT-NEXT:    ld.param.b8 %r3, [mulwides7_param_1];
; OPT-NEXT:    bfe.s32 %r4, %r3, 0, 7;
; OPT-NEXT:    mul.wide.s32 %rd1, %r2, %r4;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd1;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: mulwides7(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b16 %rs<3>;
; NOOPT-NEXT:    .reg .b64 %rd<6>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b8 %rs2, [mulwides7_param_1];
; NOOPT-NEXT:    ld.param.b8 %rs1, [mulwides7_param_0];
; NOOPT-NEXT:    cvt.u64.u16 %rd1, %rs1;
; NOOPT-NEXT:    bfe.s64 %rd2, %rd1, 0, 7;
; NOOPT-NEXT:    cvt.u64.u16 %rd3, %rs2;
; NOOPT-NEXT:    bfe.s64 %rd4, %rd3, 0, 7;
; NOOPT-NEXT:    mul.lo.s64 %rd5, %rd2, %rd4;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd5;
; 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 .b32 %r<2>;
; NOOPT-NEXT:    .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %r1, [shl30_param_0];
; NOOPT-NEXT:    cvt.s64.s32 %rd1, %r1;
; 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) {
; OPT-LABEL: shl31(
; OPT:       {
; OPT-NEXT:    .reg .b64 %rd<3>;
; OPT-EMPTY:
; OPT-NEXT:  // %bb.0:
; OPT-NEXT:    ld.param.s32 %rd1, [shl31_param_0];
; OPT-NEXT:    shl.b64 %rd2, %rd1, 31;
; OPT-NEXT:    st.param.b64 [func_retval0], %rd2;
; OPT-NEXT:    ret;
;
; NOOPT-LABEL: shl31(
; NOOPT:       {
; NOOPT-NEXT:    .reg .b32 %r<2>;
; NOOPT-NEXT:    .reg .b64 %rd<3>;
; NOOPT-EMPTY:
; NOOPT-NEXT:  // %bb.0:
; NOOPT-NEXT:    ld.param.b32 %r1, [shl31_param_0];
; NOOPT-NEXT:    cvt.s64.s32 %rd1, %r1;
; NOOPT-NEXT:    shl.b64 %rd2, %rd1, 31;
; NOOPT-NEXT:    st.param.b64 [func_retval0], %rd2;
; NOOPT-NEXT:    ret;
  %conv = sext i32 %a to i64
  %shl = shl i64 %conv, 31
  ret i64 %shl
}