<html>
<head>
<base href="https://bugs.llvm.org/">
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW - [NVPTX] Incorrect lowering for llvm.fshl intrinsic"
href="https://bugs.llvm.org/show_bug.cgi?id=41347">41347</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>[NVPTX] Incorrect lowering for llvm.fshl intrinsic
</td>
</tr>
<tr>
<th>Product</th>
<td>libraries
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>All
</td>
</tr>
<tr>
<th>OS</th>
<td>All
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>normal
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>Backend: PTX
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>J.Price@bristol.ac.uk
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>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.</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>