Skip to content

Commit 605ac34

Browse files
AlexMacleanAlexMaclean
committed
pre-commit tests
1 parent fc3ec13 commit 605ac34

File tree

1 file changed

+172
-0
lines changed

1 file changed

+172
-0
lines changed

‎llvm/test/CodeGen/NVPTX/add-rotate.ll

Lines changed: 172 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,172 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_50 | FileCheck %s
3+
4+
target triple = "nvptx64-nvidia-cuda"
5+
6+
define i32 @test_rotl(i32 %x) {
7+
; CHECK-LABEL: test_rotl(
8+
; CHECK: {
9+
; CHECK-NEXT: .reg .b32 %r<3>;
10+
; CHECK-EMPTY:
11+
; CHECK-NEXT: // %bb.0:
12+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotl_param_0];
13+
; CHECK-NEXT: shf.l.wrap.b32 %r2, %r1, %r1, 7;
14+
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
15+
; CHECK-NEXT: ret;
16+
%shl = shl i32 %x, 7
17+
%shr = lshr i32 %x, 25
18+
%add = add i32 %shl, %shr
19+
ret i32 %add
20+
}
21+
22+
define i32 @test_rotr(i32 %x) {
23+
; CHECK-LABEL: test_rotr(
24+
; CHECK: {
25+
; CHECK-NEXT: .reg .b32 %r<3>;
26+
; CHECK-EMPTY:
27+
; CHECK-NEXT: // %bb.0:
28+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotr_param_0];
29+
; CHECK-NEXT: shf.l.wrap.b32 %r2, %r1, %r1, 25;
30+
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
31+
; CHECK-NEXT: ret;
32+
%shr = lshr i32 %x, 7
33+
%shl = shl i32 %x, 25
34+
%add = add i32 %shr, %shl
35+
ret i32 %add
36+
}
37+
38+
define i32 @test_rotl_var(i32 %x, i32 %y) {
39+
; CHECK-LABEL: test_rotl_var(
40+
; CHECK: {
41+
; CHECK-NEXT: .reg .b32 %r<7>;
42+
; CHECK-EMPTY:
43+
; CHECK-NEXT: // %bb.0:
44+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotl_var_param_0];
45+
; CHECK-NEXT: ld.param.u32 %r2, [test_rotl_var_param_1];
46+
; CHECK-NEXT: shl.b32 %r3, %r1, %r2;
47+
; CHECK-NEXT: sub.s32 %r4, 32, %r2;
48+
; CHECK-NEXT: shr.u32 %r5, %r1, %r4;
49+
; CHECK-NEXT: add.s32 %r6, %r3, %r5;
50+
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
51+
; CHECK-NEXT: ret;
52+
%shl = shl i32 %x, %y
53+
%sub = sub i32 32, %y
54+
%shr = lshr i32 %x, %sub
55+
%add = add i32 %shl, %shr
56+
ret i32 %add
57+
}
58+
59+
define i32 @test_rotr_var(i32 %x, i32 %y) {
60+
; CHECK-LABEL: test_rotr_var(
61+
; CHECK: {
62+
; CHECK-NEXT: .reg .b32 %r<7>;
63+
; CHECK-EMPTY:
64+
; CHECK-NEXT: // %bb.0:
65+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotr_var_param_0];
66+
; CHECK-NEXT: ld.param.u32 %r2, [test_rotr_var_param_1];
67+
; CHECK-NEXT: shr.u32 %r3, %r1, %r2;
68+
; CHECK-NEXT: sub.s32 %r4, 32, %r2;
69+
; CHECK-NEXT: shl.b32 %r5, %r1, %r4;
70+
; CHECK-NEXT: add.s32 %r6, %r3, %r5;
71+
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
72+
; CHECK-NEXT: ret;
73+
%shr = lshr i32 %x, %y
74+
%sub = sub i32 32, %y
75+
%shl = shl i32 %x, %sub
76+
%add = add i32 %shr, %shl
77+
ret i32 %add
78+
}
79+
80+
define i32 @test_rotl_var_and(i32 %x, i32 %y) {
81+
; CHECK-LABEL: test_rotl_var_and(
82+
; CHECK: {
83+
; CHECK-NEXT: .reg .b32 %r<8>;
84+
; CHECK-EMPTY:
85+
; CHECK-NEXT: // %bb.0:
86+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotl_var_and_param_0];
87+
; CHECK-NEXT: ld.param.u32 %r2, [test_rotl_var_and_param_1];
88+
; CHECK-NEXT: shl.b32 %r3, %r1, %r2;
89+
; CHECK-NEXT: neg.s32 %r4, %r2;
90+
; CHECK-NEXT: and.b32 %r5, %r4, 31;
91+
; CHECK-NEXT: shr.u32 %r6, %r1, %r5;
92+
; CHECK-NEXT: add.s32 %r7, %r6, %r3;
93+
; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
94+
; CHECK-NEXT: ret;
95+
%shr = shl i32 %x, %y
96+
%sub = sub nsw i32 0, %y
97+
%and = and i32 %sub, 31
98+
%shl = lshr i32 %x, %and
99+
%add = add i32 %shl, %shr
100+
ret i32 %add
101+
}
102+
103+
define i32 @test_rotr_var_and(i32 %x, i32 %y) {
104+
; CHECK-LABEL: test_rotr_var_and(
105+
; CHECK: {
106+
; CHECK-NEXT: .reg .b32 %r<8>;
107+
; CHECK-EMPTY:
108+
; CHECK-NEXT: // %bb.0:
109+
; CHECK-NEXT: ld.param.u32 %r1, [test_rotr_var_and_param_0];
110+
; CHECK-NEXT: ld.param.u32 %r2, [test_rotr_var_and_param_1];
111+
; CHECK-NEXT: shr.u32 %r3, %r1, %r2;
112+
; CHECK-NEXT: neg.s32 %r4, %r2;
113+
; CHECK-NEXT: and.b32 %r5, %r4, 31;
114+
; CHECK-NEXT: shl.b32 %r6, %r1, %r5;
115+
; CHECK-NEXT: add.s32 %r7, %r3, %r6;
116+
; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
117+
; CHECK-NEXT: ret;
118+
%shr = lshr i32 %x, %y
119+
%sub = sub nsw i32 0, %y
120+
%and = and i32 %sub, 31
121+
%shl = shl i32 %x, %and
122+
%add = add i32 %shr, %shl
123+
ret i32 %add
124+
}
125+
126+
define i32 @test_fshl_special_case(i32 %x0, i32 %x1, i32 %y) {
127+
; CHECK-LABEL: test_fshl_special_case(
128+
; CHECK: {
129+
; CHECK-NEXT: .reg .b32 %r<9>;
130+
; CHECK-EMPTY:
131+
; CHECK-NEXT: // %bb.0:
132+
; CHECK-NEXT: ld.param.u32 %r1, [test_fshl_special_case_param_0];
133+
; CHECK-NEXT: ld.param.u32 %r2, [test_fshl_special_case_param_2];
134+
; CHECK-NEXT: shl.b32 %r3, %r1, %r2;
135+
; CHECK-NEXT: ld.param.u32 %r4, [test_fshl_special_case_param_1];
136+
; CHECK-NEXT: shr.u32 %r5, %r4, 1;
137+
; CHECK-NEXT: xor.b32 %r6, %r2, 31;
138+
; CHECK-NEXT: shr.u32 %r7, %r5, %r6;
139+
; CHECK-NEXT: add.s32 %r8, %r3, %r7;
140+
; CHECK-NEXT: st.param.b32 [func_retval0], %r8;
141+
; CHECK-NEXT: ret;
142+
%shl = shl i32 %x0, %y
143+
%srli = lshr i32 %x1, 1
144+
%x = xor i32 %y, 31
145+
%srlo = lshr i32 %srli, %x
146+
%o = add i32 %shl, %srlo
147+
ret i32 %o
148+
}
149+
150+
define i32 @test_fshr_special_case(i32 %x0, i32 %x1, i32 %y) {
151+
; CHECK-LABEL: test_fshr_special_case(
152+
; CHECK: {
153+
; CHECK-NEXT: .reg .b32 %r<9>;
154+
; CHECK-EMPTY:
155+
; CHECK-NEXT: // %bb.0:
156+
; CHECK-NEXT: ld.param.u32 %r1, [test_fshr_special_case_param_0];
157+
; CHECK-NEXT: ld.param.u32 %r2, [test_fshr_special_case_param_1];
158+
; CHECK-NEXT: ld.param.u32 %r3, [test_fshr_special_case_param_2];
159+
; CHECK-NEXT: shr.u32 %r4, %r2, %r3;
160+
; CHECK-NEXT: shl.b32 %r5, %r1, 1;
161+
; CHECK-NEXT: xor.b32 %r6, %r3, 31;
162+
; CHECK-NEXT: shl.b32 %r7, %r5, %r6;
163+
; CHECK-NEXT: add.s32 %r8, %r4, %r7;
164+
; CHECK-NEXT: st.param.b32 [func_retval0], %r8;
165+
; CHECK-NEXT: ret;
166+
%shl = lshr i32 %x1, %y
167+
%srli = shl i32 %x0, 1
168+
%x = xor i32 %y, 31
169+
%srlo = shl i32 %srli, %x
170+
%o = add i32 %shl, %srlo
171+
ret i32 %o
172+
}

0 commit comments

Comments
 (0)