[llvm] 49dc163 - [NVPTX] mulwide.ll - regenerate test checks
Simon Pilgrim via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 7 05:05:37 PDT 2025
Author: Simon Pilgrim
Date: 2025-07-07T13:04:33+01:00
New Revision: 49dc163d15ca6ade4778971560dcf3700a11687f
URL: https://github.com/llvm/llvm-project/commit/49dc163d15ca6ade4778971560dcf3700a11687f
DIFF: https://github.com/llvm/llvm-project/commit/49dc163d15ca6ade4778971560dcf3700a11687f.diff
LOG: [NVPTX] mulwide.ll - regenerate test checks
Added:
Modified:
llvm/test/CodeGen/NVPTX/mulwide.ll
Removed:
################################################################################
diff --git a/llvm/test/CodeGen/NVPTX/mulwide.ll b/llvm/test/CodeGen/NVPTX/mulwide.ll
index 7db7dec561992..666c7a160e1f0 100644
--- a/llvm/test/CodeGen/NVPTX/mulwide.ll
+++ b/llvm/test/CodeGen/NVPTX/mulwide.ll
@@ -1,111 +1,325 @@
+; 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 %}
-; OPT-LABEL: @mulwide16
-; NOOPT-LABEL: @mulwide16
define i32 @mulwide16(i16 %a, i16 %b) {
-; OPT: mul.wide.s16
-; NOOPT: mul.lo.s32
+; 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
}
-; OPT-LABEL: @mulwideu16
-; NOOPT-LABEL: @mulwideu16
define i32 @mulwideu16(i16 %a, i16 %b) {
-; OPT: mul.wide.u16
-; NOOPT: mul.lo.s32
+; 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
}
-; OPT-LABEL: @mulwide8
-; NOOPT-LABEL: @mulwide8
define i32 @mulwide8(i8 %a, i8 %b) {
-; OPT: mul.wide.s16
-; NOOPT: mul.lo.s32
+; 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
}
-; OPT-LABEL: @mulwideu8
-; NOOPT-LABEL: @mulwideu8
define i32 @mulwideu8(i8 %a, i8 %b) {
-; OPT: mul.wide.u16
-; NOOPT: mul.lo.s32
+; 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
}
-; OPT-LABEL: @mulwide32
-; NOOPT-LABEL: @mulwide32
define i64 @mulwide32(i32 %a, i32 %b) {
-; OPT: mul.wide.s32
-; NOOPT: mul.lo.s64
+; 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
}
-; OPT-LABEL: @mulwideu32
-; NOOPT-LABEL: @mulwideu32
define i64 @mulwideu32(i32 %a, i32 %b) {
-; OPT: mul.wide.u32
-; NOOPT: mul.lo.s64
+; 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
}
-; OPT-LABEL: @mulwideu7
-; NOOPT-LABEL: @mulwideu7
define i64 @mulwideu7(i7 %a, i7 %b) {
-; OPT: mul.wide.u32
-; NOOPT: mul.lo.s64
+; 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
}
-; OPT-LABEL: @mulwides7
-; NOOPT-LABEL: @mulwides7
define i64 @mulwides7(i7 %a, i7 %b) {
-; OPT: mul.wide.s32
-; NOOPT: mul.lo.s64
+; 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
}
-; OPT-LABEL: @shl30
-; NOOPT-LABEL: @shl30
define i64 @shl30(i32 %a) {
-; OPT: mul.wide
-; NOOPT: shl.b64
+; 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
}
-; OPT-LABEL: @shl31
-; NOOPT-LABEL: @shl31
define i64 @shl31(i32 %a) {
-; OPT-NOT: mul.wide
-; NOOPT-NOT: mul.wide
+; 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
More information about the llvm-commits
mailing list