1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s 3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %} 4 5declare i64 @llvm.nvvm.rotate.b64(i64, i32) 6declare i64 @llvm.nvvm.rotate.right.b64(i64, i32) 7 8define i64 @rotate64(i64 %a, i32 %b) { 9; CHECK-LABEL: rotate64( 10; CHECK: { 11; CHECK-NEXT: .reg .b64 %rd<5>; 12; CHECK-EMPTY: 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; 19; CHECK-NEXT: ret; 20 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 3) 21 ret i64 %val 22} 23 24define i64 @rotateright64(i64 %a, i32 %b) { 25; CHECK-LABEL: rotateright64( 26; CHECK: { 27; CHECK-NEXT: .reg .b64 %rd<5>; 28; CHECK-EMPTY: 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; 35; CHECK-NEXT: ret; 36 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 3) 37 ret i64 %val 38} 39