42e0d30268
Implements `(sign_extend|zero_extend (mul|shl) x, y) -> (mul.wide x, y)` as a DAG combine. Implements `(add (mul.wide a, b), c) -> (mad.wide a, b, c)` in instruction selection.
1340 lines
35 KiB
LLVM
1340 lines
35 KiB
LLVM
; 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<3>;
|
|
; 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: ld.param.b64 %rd1, [t1_param_2];
|
|
; O1-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; 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<3>;
|
|
; 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: ld.param.b64 %rd1, [t2_param_2];
|
|
; O1-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; 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<2>;
|
|
; 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: mad.wide.s32 %rd1, %r1, %r2, 1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd1;
|
|
; 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<3>;
|
|
; 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: mad.wide.s32 %rd2, %r1, 3, %rd1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; 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<2>;
|
|
; O1-EMPTY:
|
|
; O1-NEXT: // %bb.0:
|
|
; O1-NEXT: ld.param.b32 %r1, [t4_1_param_0];
|
|
; O1-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd1;
|
|
; 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<3>;
|
|
; 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: ld.param.b64 %rd1, [t5_param_2];
|
|
; O1-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; 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<3>;
|
|
; 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: ld.param.b64 %rd1, [t6_param_2];
|
|
; O1-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
|
|
; O1-NEXT: st.param.b64 [func_retval0], %rd2;
|
|
; 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<3>;
|
|
; 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: ld.param.b32 %r1, [t32_param_2];
|
|
; O1-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; 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<3>;
|
|
; 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: ld.param.b32 %r1, [t33_param_2];
|
|
; O1-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; 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<2>;
|
|
; 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: mad.wide.s16 %r1, %rs1, %rs2, 1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r1;
|
|
; 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<3>;
|
|
; 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: mad.wide.s16 %r2, %rs1, 3, %r1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; 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<2>;
|
|
; O1-EMPTY:
|
|
; O1-NEXT: // %bb.0:
|
|
; O1-NEXT: ld.param.b16 %rs1, [t36_param_0];
|
|
; O1-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r1;
|
|
; 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<3>;
|
|
; 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: ld.param.b32 %r1, [t37_param_2];
|
|
; O1-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; 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<3>;
|
|
; 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: ld.param.b32 %r1, [t38_param_2];
|
|
; O1-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
|
|
; O1-NEXT: st.param.b32 [func_retval0], %r2;
|
|
; 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
|
|
}
|