Skip to content

Commit

Permalink
address comments - update rotate_64.ll
Browse files Browse the repository at this point in the history
  • Loading branch information
AlexMaclean committed Sep 23, 2024
1 parent b967616 commit 979aa53
Showing 1 changed file with 23 additions and 10 deletions.
33 changes: 23 additions & 10 deletions llvm/test/CodeGen/NVPTX/rotate_64.ll
Original file line number Diff line number Diff line change
@@ -1,25 +1,38 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}

declare i64 @llvm.nvvm.rotate.b64(i64, i32)
declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)

; CHECK: rotate64
define i64 @rotate64(i64 %a, i32 %b) {
; CHECK: shr.u64 [[RHS:%.*]], [[RD1:%.*]], 61;
; CHECK: shl.b64 [[LHS:%.*]], [[RD1]], 3;
; CHECK: or.b64 [[RD2:%.*]], [[LHS]], [[RHS]];
; CHECK: ret
; CHECK-LABEL: rotate64(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [rotate64_param_0];
; CHECK-NEXT: shr.u64 %rd2, %rd1, 61;
; CHECK-NEXT: shl.b64 %rd3, %rd1, 3;
; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4;
; CHECK-NEXT: ret;
%val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 3)
ret i64 %val
}

; CHECK: rotateright64
define i64 @rotateright64(i64 %a, i32 %b) {
; CHECK: shl.b64 [[RHS:%.*]], [[RD1:%.*]], 61;
; CHECK: shr.u64 [[LHS:%.*]], [[RD1]], 3;
; CHECK: or.b64 [[RD2:%.*]], [[LHS]], [[RHS]];
; CHECK: ret
; CHECK-LABEL: rotateright64(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [rotateright64_param_0];
; CHECK-NEXT: shl.b64 %rd2, %rd1, 61;
; CHECK-NEXT: shr.u64 %rd3, %rd1, 3;
; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;
; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd4;
; CHECK-NEXT: ret;
%val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 3)
ret i64 %val
}

0 comments on commit 979aa53

Please sign in to comment.