43 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 | FileCheck %s3; 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 .b32 %r<5>;12; CHECK-NEXT: .reg .b64 %rd<3>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b64 %rd1, [rotate64_param_0];16; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;17; CHECK-NEXT: shf.l.wrap.b32 %r3, %r1, %r2, 3;18; CHECK-NEXT: shf.l.wrap.b32 %r4, %r2, %r1, 3;19; CHECK-NEXT: mov.b64 %rd2, {%r4, %r3};20; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;21; CHECK-NEXT: ret;22 %val = tail call i64 @llvm.nvvm.rotate.b64(i64 %a, i32 3)23 ret i64 %val24}25 26define i64 @rotateright64(i64 %a, i32 %b) {27; CHECK-LABEL: rotateright64(28; CHECK: {29; CHECK-NEXT: .reg .b32 %r<5>;30; CHECK-NEXT: .reg .b64 %rd<3>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b64 %rd1, [rotateright64_param_0];34; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;35; CHECK-NEXT: shf.r.wrap.b32 %r3, %r2, %r1, 3;36; CHECK-NEXT: shf.r.wrap.b32 %r4, %r1, %r2, 3;37; CHECK-NEXT: mov.b64 %rd2, {%r4, %r3};38; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;39; CHECK-NEXT: ret;40 %val = tail call i64 @llvm.nvvm.rotate.right.b64(i64 %a, i32 3)41 ret i64 %val42}43