88 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7define i16 @bswap16(i16 %a) {8; CHECK-LABEL: bswap16(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b16 %r1, [bswap16_param_0];14; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7701U;15; CHECK-NEXT: st.param.b32 [func_retval0], %r2;16; CHECK-NEXT: ret;17 %b = tail call i16 @llvm.bswap.i16(i16 %a)18 ret i16 %b19}20 21 22define i32 @bswap32(i32 %a) {23; CHECK-LABEL: bswap32(24; CHECK: {25; CHECK-NEXT: .reg .b32 %r<3>;26; CHECK-EMPTY:27; CHECK-NEXT: // %bb.0:28; CHECK-NEXT: ld.param.b32 %r1, [bswap32_param_0];29; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x123U;30; CHECK-NEXT: st.param.b32 [func_retval0], %r2;31; CHECK-NEXT: ret;32 %b = tail call i32 @llvm.bswap.i32(i32 %a)33 ret i32 %b34}35 36 37define <2 x i16> @bswapv2i16(<2 x i16> %a) #0 {38; CHECK-LABEL: bswapv2i16(39; CHECK: {40; CHECK-NEXT: .reg .b32 %r<3>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: ld.param.b32 %r1, [bswapv2i16_param_0];44; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x2301U;45; CHECK-NEXT: st.param.b32 [func_retval0], %r2;46; CHECK-NEXT: ret;47 %b = tail call <2 x i16> @llvm.bswap.v2i16(<2 x i16> %a)48 ret <2 x i16> %b49}50 51define i64 @bswap64(i64 %a) {52; CHECK-LABEL: bswap64(53; CHECK: {54; CHECK-NEXT: .reg .b32 %r<5>;55; CHECK-NEXT: .reg .b64 %rd<3>;56; CHECK-EMPTY:57; CHECK-NEXT: // %bb.0:58; CHECK-NEXT: ld.param.b64 %rd1, [bswap64_param_0];59; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;60; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x123U;61; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x123U;62; CHECK-NEXT: mov.b64 %rd2, {%r4, %r3};63; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;64; CHECK-NEXT: ret;65 %b = tail call i64 @llvm.bswap.i64(i64 %a)66 ret i64 %b67}68 69define <2 x i32> @bswapv2i32(<2 x i32> %a) {70; CHECK-LABEL: bswapv2i32(71; CHECK: {72; CHECK-NEXT: .reg .b32 %r<5>;73; CHECK-EMPTY:74; CHECK-NEXT: // %bb.0:75; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [bswapv2i32_param_0];76; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x123U;77; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x123U;78; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};79; CHECK-NEXT: ret;80 %b = tail call <2 x i32> @llvm.bswap.v2i32(<2 x i32> %a)81 ret <2 x i32> %b82}83declare i16 @llvm.bswap.i16(i16)84declare i32 @llvm.bswap.i32(i32)85declare <2 x i16> @llvm.bswap.v2i16(<2 x i16>)86declare i64 @llvm.bswap.i64(i64)87declare <2 x i32> @llvm.bswap.v2i32(<2 x i32>)88