llvm/test/CodeGen/NVPTX/bswap.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}

target triple = "nvptx64-nvidia-cuda"

define i16 @bswap16(i16 %a) {
; CHECK-LABEL: bswap16(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b16 %r1, [bswap16_param_0];
; CHECK-NEXT:    prmt.b32 %r2, %r1, 0, 0x7701U;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
  %b = tail call i16 @llvm.bswap.i16(i16 %a)
  ret i16 %b
}


define i32 @bswap32(i32 %a) {
; CHECK-LABEL: bswap32(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [bswap32_param_0];
; CHECK-NEXT:    prmt.b32 %r2, %r1, 0, 0x123U;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
  %b = tail call i32 @llvm.bswap.i32(i32 %a)
  ret i32 %b
}


define <2 x i16> @bswapv2i16(<2 x i16> %a) #0 {
; CHECK-LABEL: bswapv2i16(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [bswapv2i16_param_0];
; CHECK-NEXT:    prmt.b32 %r2, %r1, 0, 0x2301U;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
  %b = tail call <2 x i16> @llvm.bswap.v2i16(<2 x i16> %a)
  ret <2 x i16> %b
}

define i64 @bswap64(i64 %a) {
; CHECK-LABEL: bswap64(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<5>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b64 %rd1, [bswap64_param_0];
; CHECK-NEXT:    mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT:    prmt.b32 %r3, %r1, 0, 0x123U;
; CHECK-NEXT:    prmt.b32 %r4, %r2, 0, 0x123U;
; CHECK-NEXT:    mov.b64 %rd2, {%r4, %r3};
; CHECK-NEXT:    st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT:    ret;
  %b = tail call i64 @llvm.bswap.i64(i64 %a)
  ret i64 %b
}

define <2 x i32> @bswapv2i32(<2 x i32> %a) {
; CHECK-LABEL: bswapv2i32(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.v2.b32 {%r1, %r2}, [bswapv2i32_param_0];
; CHECK-NEXT:    prmt.b32 %r3, %r2, 0, 0x123U;
; CHECK-NEXT:    prmt.b32 %r4, %r1, 0, 0x123U;
; CHECK-NEXT:    st.param.v2.b32 [func_retval0], {%r4, %r3};
; CHECK-NEXT:    ret;
  %b = tail call <2 x i32> @llvm.bswap.v2i32(<2 x i32> %a)
  ret <2 x i32> %b
}
declare i16 @llvm.bswap.i16(i16)
declare i32 @llvm.bswap.i32(i32)
declare <2 x i16> @llvm.bswap.v2i16(<2 x i16>)
declare i64 @llvm.bswap.i64(i64)
declare <2 x i32> @llvm.bswap.v2i32(<2 x i32>)