[llvm-bugs] [Bug 41347] New: [NVPTX] Incorrect lowering for llvm.fshl intrinsic

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=41347

            Bug ID: 41347
           Summary: [NVPTX] Incorrect lowering for llvm.fshl intrinsic
           Product: libraries
           Version: trunk
          Hardware: All
                OS: All
            Status: NEW
          Severity: normal
          Priority: P
         Component: Backend: PTX
          Assignee: unassignedbugs at nondot.org
          Reporter: J.Price at bristol.ac.uk
                CC: llvm-bugs at lists.llvm.org

The llvm.fshl intrinsic is supposed to behave like a rotate left operation when
the first two operands are the same. The PTX backend seems to generate code
that doesn't account for the shift amount being larger than the element size
(explicitly called out by the LLVM reference manual).

Example LLVM snippet:

```
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64"

; Function Attrs: noinline nounwind
define i64 @test_rotate(i64 %y) {
entry:
  %x = tail call i64 @llvm.fshl.i64(i64 72340172838076673, i64
72340172838076673, i64 %y)
  ret i64 %x
}

; Function Attrs: nounwind readnone speculatable
declare i64 @llvm.fshl.i64(i64, i64, i64)
```

This produces the following PTX:


```
//
// Generated by LLVM NVPTX Back-End
//

.version 3.2
.target sm_20
.address_size 64

        // .globl       test_rotate     // -- Begin function test_rotate
                                        // @test_rotate
.visible .func  (.param .b64 func_retval0) test_rotate(
        .param .b64 test_rotate_param_0
)
{
        .reg .b32       %r<2>;
        .reg .b64       %rd<3>;

// %bb.0:                               // %entry
        ld.param.u32    %r1, [test_rotate_param_0];
        mov.u64         %rd1, 72340172838076673;
        {
        .reg .b64 %lhs;
        .reg .b64 %rhs;
        .reg .u32 %amt2;
        shl.b64         %lhs, %rd1, %r1;
        sub.u32         %amt2, 64, %r1;
        shr.b64         %rhs, %rd1, %amt2;
        add.u64         %rd2, %lhs, %rhs;
        }
        st.param.b64    [func_retval0+0], %rd2;
        ret;
                                        // -- End function
}
```

Running this with an argument >64 just produces a result of zero.

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20190402/493b9b56/attachment-0001.html>