xref: /llvm-project/llvm/test/CodeGen/NVPTX/funnel-shift-clamp.ll (revision b279f6b098d3849f7f1c1f539b108307d5f8ae2d) 
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2; RUN: llc < %s -mtriple=nvptx -mcpu=sm_61 | FileCheck %s
3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_61 | FileCheck %s
4
5target triple = "nvptx-nvidia-cuda"
6
7declare i32 @llvm.nvvm.fshr.clamp.i32(i32, i32, i32)
8declare i32 @llvm.nvvm.fshl.clamp.i32(i32, i32, i32)
9
10define i32 @fshr_clamp_r(i32 %hi, i32 %lo, i32 %n) {
11; CHECK-LABEL: fshr_clamp_r(
12; CHECK:       {
13; CHECK-NEXT:    .reg .b32 %r<5>;
14; CHECK-EMPTY:
15; CHECK-NEXT:  // %bb.0:
16; CHECK-NEXT:    ld.param.u32 %r1, [fshr_clamp_r_param_0];
17; CHECK-NEXT:    ld.param.u32 %r2, [fshr_clamp_r_param_1];
18; CHECK-NEXT:    ld.param.u32 %r3, [fshr_clamp_r_param_2];
19; CHECK-NEXT:    shf.r.clamp.b32 %r4, %r2, %r1, %r3;
20; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
21; CHECK-NEXT:    ret;
22  %call = call i32 @llvm.nvvm.fshr.clamp.i32(i32 %hi, i32 %lo, i32 %n)
23  ret i32 %call
24}
25
26define i32 @fshl_clamp_r(i32 %hi, i32 %lo, i32 %n) {
27; CHECK-LABEL: fshl_clamp_r(
28; CHECK:       {
29; CHECK-NEXT:    .reg .b32 %r<5>;
30; CHECK-EMPTY:
31; CHECK-NEXT:  // %bb.0:
32; CHECK-NEXT:    ld.param.u32 %r1, [fshl_clamp_r_param_0];
33; CHECK-NEXT:    ld.param.u32 %r2, [fshl_clamp_r_param_1];
34; CHECK-NEXT:    ld.param.u32 %r3, [fshl_clamp_r_param_2];
35; CHECK-NEXT:    shf.l.clamp.b32 %r4, %r2, %r1, %r3;
36; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
37; CHECK-NEXT:    ret;
38  %call = call i32 @llvm.nvvm.fshl.clamp.i32(i32 %hi, i32 %lo, i32 %n)
39  ret i32 %call
40}
41
42define i32 @fshr_clamp_i(i32 %hi, i32 %lo) {
43; CHECK-LABEL: fshr_clamp_i(
44; CHECK:       {
45; CHECK-NEXT:    .reg .b32 %r<4>;
46; CHECK-EMPTY:
47; CHECK-NEXT:  // %bb.0:
48; CHECK-NEXT:    ld.param.u32 %r1, [fshr_clamp_i_param_0];
49; CHECK-NEXT:    ld.param.u32 %r2, [fshr_clamp_i_param_1];
50; CHECK-NEXT:    shf.r.clamp.b32 %r3, %r2, %r1, 3;
51; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
52; CHECK-NEXT:    ret;
53  %call = call i32 @llvm.nvvm.fshr.clamp.i32(i32 %hi, i32 %lo, i32 3)
54  ret i32 %call
55}
56
57define i32 @fshl_clamp_i(i32 %hi, i32 %lo) {
58; CHECK-LABEL: fshl_clamp_i(
59; CHECK:       {
60; CHECK-NEXT:    .reg .b32 %r<4>;
61; CHECK-EMPTY:
62; CHECK-NEXT:  // %bb.0:
63; CHECK-NEXT:    ld.param.u32 %r1, [fshl_clamp_i_param_0];
64; CHECK-NEXT:    ld.param.u32 %r2, [fshl_clamp_i_param_1];
65; CHECK-NEXT:    shf.l.clamp.b32 %r3, %r2, %r1, 3;
66; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
67; CHECK-NEXT:    ret;
68  %call = call i32 @llvm.nvvm.fshl.clamp.i32(i32 %hi, i32 %lo, i32 3)
69  ret i32 %call
70}
71