1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
2 ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s
3 ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}
5 define i16 @test_mulhi_i16(i16 %x, i16 %y) {
6 ; CHECK-LABEL: test_mulhi_i16(
8 ; CHECK-NEXT: .reg .b16 %rs<4>;
9 ; CHECK-NEXT: .reg .b32 %r<2>;
11 ; CHECK-NEXT: // %bb.0:
12 ; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_i16_param_0];
13 ; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_i16_param_1];
14 ; CHECK-NEXT: mul.hi.s16 %rs3, %rs1, %rs2;
15 ; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
16 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
18 %1 = call i16 @llvm.nvvm.mulhi.s(i16 %x, i16 %y)
22 define i16 @test_mulhi_u16(i16 %x, i16 %y) {
23 ; CHECK-LABEL: test_mulhi_u16(
25 ; CHECK-NEXT: .reg .b16 %rs<4>;
26 ; CHECK-NEXT: .reg .b32 %r<2>;
28 ; CHECK-NEXT: // %bb.0:
29 ; CHECK-NEXT: ld.param.u16 %rs1, [test_mulhi_u16_param_0];
30 ; CHECK-NEXT: ld.param.u16 %rs2, [test_mulhi_u16_param_1];
31 ; CHECK-NEXT: mul.hi.u16 %rs3, %rs1, %rs2;
32 ; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
33 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
35 %1 = call i16 @llvm.nvvm.mulhi.us(i16 %x, i16 %y)
39 define i32 @test_mulhi_i32(i32 %x, i32 %y) {
40 ; CHECK-LABEL: test_mulhi_i32(
42 ; CHECK-NEXT: .reg .b32 %r<4>;
44 ; CHECK-NEXT: // %bb.0:
45 ; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_i32_param_0];
46 ; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_i32_param_1];
47 ; CHECK-NEXT: mul.hi.s32 %r3, %r1, %r2;
48 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
50 %1 = call i32 @llvm.nvvm.mulhi.i(i32 %x, i32 %y)
54 define i32 @test_mulhi_u32(i32 %x, i32 %y) {
55 ; CHECK-LABEL: test_mulhi_u32(
57 ; CHECK-NEXT: .reg .b32 %r<4>;
59 ; CHECK-NEXT: // %bb.0:
60 ; CHECK-NEXT: ld.param.u32 %r1, [test_mulhi_u32_param_0];
61 ; CHECK-NEXT: ld.param.u32 %r2, [test_mulhi_u32_param_1];
62 ; CHECK-NEXT: mul.hi.u32 %r3, %r1, %r2;
63 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3;
65 %1 = call i32 @llvm.nvvm.mulhi.ui(i32 %x, i32 %y)
69 define i64 @test_mulhi_i64(i64 %x, i64 %y) {
70 ; CHECK-LABEL: test_mulhi_i64(
72 ; CHECK-NEXT: .reg .b64 %rd<4>;
74 ; CHECK-NEXT: // %bb.0:
75 ; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_i64_param_0];
76 ; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_i64_param_1];
77 ; CHECK-NEXT: mul.hi.s64 %rd3, %rd1, %rd2;
78 ; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
80 %1 = call i64 @llvm.nvvm.mulhi.ll(i64 %x, i64 %y)
84 define i64 @test_mulhi_u64(i64 %x, i64 %y) {
85 ; CHECK-LABEL: test_mulhi_u64(
87 ; CHECK-NEXT: .reg .b64 %rd<4>;
89 ; CHECK-NEXT: // %bb.0:
90 ; CHECK-NEXT: ld.param.u64 %rd1, [test_mulhi_u64_param_0];
91 ; CHECK-NEXT: ld.param.u64 %rd2, [test_mulhi_u64_param_1];
92 ; CHECK-NEXT: mul.hi.u64 %rd3, %rd1, %rd2;
93 ; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
95 %1 = call i64 @llvm.nvvm.mulhi.ull(i64 %x, i64 %y)
99 declare i16 @llvm.nvvm.mulhi.s(i16, i16)
100 declare i16 @llvm.nvvm.mulhi.us(i16, i16)
101 declare i32 @llvm.nvvm.mulhi.i(i32, i32)
102 declare i32 @llvm.nvvm.mulhi.ui(i32, i32)
103 declare i64 @llvm.nvvm.mulhi.ll(i64, i64)
104 declare i64 @llvm.nvvm.mulhi.ull(i64, i64)