1 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
3 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
5 target triple = "nvptx64-nvidia-cuda"
7 define i16 @bswap16(i16 %a) {
8 ; CHECK-LABEL: bswap16(
10 ; CHECK-NEXT: .reg .b16 %rs<5>;
11 ; CHECK-NEXT: .reg .b32 %r<2>;
13 ; CHECK-NEXT: // %bb.0:
14 ; CHECK-NEXT: ld.param.u16 %rs1, [bswap16_param_0];
15 ; CHECK-NEXT: shr.u16 %rs2, %rs1, 8;
16 ; CHECK-NEXT: shl.b16 %rs3, %rs1, 8;
17 ; CHECK-NEXT: or.b16 %rs4, %rs3, %rs2;
18 ; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
19 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1;
21 %b = tail call i16 @llvm.bswap.i16(i16 %a)
26 define i32 @bswap32(i32 %a) {
27 ; CHECK-LABEL: bswap32(
29 ; CHECK-NEXT: .reg .b32 %r<3>;
31 ; CHECK-NEXT: // %bb.0:
32 ; CHECK-NEXT: ld.param.u32 %r1, [bswap32_param_0];
33 ; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 291;
34 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2;
36 %b = tail call i32 @llvm.bswap.i32(i32 %a)
41 define <2 x i16> @bswapv2i16(<2 x i16> %a) #0 {
42 ; CHECK-LABEL: bswapv2i16(
44 ; CHECK-NEXT: .reg .b32 %r<3>;
46 ; CHECK-NEXT: // %bb.0:
47 ; CHECK-NEXT: ld.param.u32 %r1, [bswapv2i16_param_0];
48 ; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 8961;
49 ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2;
51 %b = tail call <2 x i16> @llvm.bswap.v2i16(<2 x i16> %a)
55 define i64 @bswap64(i64 %a) {
56 ; CHECK-LABEL: bswap64(
58 ; CHECK-NEXT: .reg .b32 %r<5>;
59 ; CHECK-NEXT: .reg .b64 %rd<3>;
61 ; CHECK-NEXT: // %bb.0:
62 ; CHECK-NEXT: ld.param.u64 %rd1, [bswap64_param_0];
63 ; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {%r1, tmp}, %rd1; }
64 ; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 291;
65 ; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r3}, %rd1; }
66 ; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 291;
67 ; CHECK-NEXT: mov.b64 %rd2, {%r4, %r2};
68 ; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd2;
70 %b = tail call i64 @llvm.bswap.i64(i64 %a)
74 declare i16 @llvm.bswap.i16(i16)
75 declare i32 @llvm.bswap.i32(i32)
76 declare <2 x i16> @llvm.bswap.v2i16(<2 x i16>)
77 declare i64 @llvm.bswap.i64(i64)