1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2 ; RUN: llc < %s -march=nvptx64 | FileCheck %s
3 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
5 declare i64 @llvm.nvvm.rotate.b64(i64, i32)
6 declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
8 define i64 @rotate64(i64 %a, i32 %b) {
9 ; CHECK-LABEL: rotate64(
11 ; CHECK-NEXT: .reg .b64 %rd<5>;
13 ; CHECK-NEXT: // %bb.0:
14 ; CHECK-NEXT: ld.param.u64 %rd1, [rotate64_param_0];
15 ; CHECK-NEXT: shr.u64 %rd2, %rd1, 61;
16 ; CHECK-NEXT: shl.b64 %rd3, %rd1, 3;
17 ; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;
18 ; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;
20 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 3)
24 define i64 @rotateright64(i64 %a, i32 %b) {
25 ; CHECK-LABEL: rotateright64(
27 ; CHECK-NEXT: .reg .b64 %rd<5>;
29 ; CHECK-NEXT: // %bb.0:
30 ; CHECK-NEXT: ld.param.u64 %rd1, [rotateright64_param_0];
31 ; CHECK-NEXT: shl.b64 %rd2, %rd1, 61;
32 ; CHECK-NEXT: shr.u64 %rd3, %rd1, 3;
33 ; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;
34 ; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;
36 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 3)