Skip to content

Commit 0019ac7

Browse files
AlexMacleanAlexMaclean
committed
pre-commit tests
1 parent 577631f commit 0019ac7

File tree

2 files changed

+314
-0
lines changed

2 files changed

+314
-0
lines changed

‎llvm/test/CodeGen/NVPTX/fast-math.ll‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -131,6 +131,20 @@ define float @fadd_ftz(float %a, float %b) #1 {
131131
declare float @llvm.sin.f32(float)
132132
declare float @llvm.cos.f32(float)
133133

134+
; CHECK-LABEL: fsin_approx_afn
135+
; CHECK: sin.approx.f32
136+
define float @fsin_approx_afn(float %a) {
137+
%r = tail call afn float @llvm.sin.f32(float %a)
138+
ret float %r
139+
}
140+
141+
; CHECK-LABEL: fcos_approx_afn
142+
; CHECK: cos.approx.f32
143+
define float @fcos_approx_afn(float %a) {
144+
%r = tail call afn float @llvm.cos.f32(float %a)
145+
ret float %r
146+
}
147+
134148
; CHECK-LABEL: fsin_approx
135149
; CHECK: sin.approx.f32
136150
define float @fsin_approx(float %a) #0 {

‎llvm/test/CodeGen/NVPTX/frem.ll‎

Lines changed: 300 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,300 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s --enable-unsafe-fp-math | FileCheck %s --check-prefixes=FAST
3+
; RUN: llc < %s | FileCheck %s --check-prefixes=NORMAL
4+
5+
6+
target triple = "nvptx64-unknown-cuda"
7+
8+
define half @frem_f16(half %a, half %b) {
9+
; FAST-LABEL: frem_f16(
10+
; FAST: {
11+
; FAST-NEXT: .reg .b16 %rs<4>;
12+
; FAST-NEXT: .reg .f32 %f<7>;
13+
; FAST-EMPTY:
14+
; FAST-NEXT: // %bb.0:
15+
; FAST-NEXT: ld.param.b16 %rs1, [frem_f16_param_0];
16+
; FAST-NEXT: ld.param.b16 %rs2, [frem_f16_param_1];
17+
; FAST-NEXT: cvt.f32.f16 %f1, %rs2;
18+
; FAST-NEXT: cvt.f32.f16 %f2, %rs1;
19+
; FAST-NEXT: div.rn.f32 %f3, %f2, %f1;
20+
; FAST-NEXT: cvt.rzi.f32.f32 %f4, %f3;
21+
; FAST-NEXT: mul.f32 %f5, %f4, %f1;
22+
; FAST-NEXT: sub.f32 %f6, %f2, %f5;
23+
; FAST-NEXT: cvt.rn.f16.f32 %rs3, %f6;
24+
; FAST-NEXT: st.param.b16 [func_retval0], %rs3;
25+
; FAST-NEXT: ret;
26+
;
27+
; NORMAL-LABEL: frem_f16(
28+
; NORMAL: {
29+
; NORMAL-NEXT: .reg .pred %p<2>;
30+
; NORMAL-NEXT: .reg .b16 %rs<4>;
31+
; NORMAL-NEXT: .reg .f32 %f<8>;
32+
; NORMAL-EMPTY:
33+
; NORMAL-NEXT: // %bb.0:
34+
; NORMAL-NEXT: ld.param.b16 %rs1, [frem_f16_param_0];
35+
; NORMAL-NEXT: ld.param.b16 %rs2, [frem_f16_param_1];
36+
; NORMAL-NEXT: cvt.f32.f16 %f1, %rs2;
37+
; NORMAL-NEXT: cvt.f32.f16 %f2, %rs1;
38+
; NORMAL-NEXT: div.rn.f32 %f3, %f2, %f1;
39+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f4, %f3;
40+
; NORMAL-NEXT: mul.f32 %f5, %f4, %f1;
41+
; NORMAL-NEXT: sub.f32 %f6, %f2, %f5;
42+
; NORMAL-NEXT: testp.infinite.f32 %p1, %f1;
43+
; NORMAL-NEXT: selp.f32 %f7, %f2, %f6, %p1;
44+
; NORMAL-NEXT: cvt.rn.f16.f32 %rs3, %f7;
45+
; NORMAL-NEXT: st.param.b16 [func_retval0], %rs3;
46+
; NORMAL-NEXT: ret;
47+
%r = frem half %a, %b
48+
ret half %r
49+
}
50+
51+
define float @frem_f32(float %a, float %b) {
52+
; FAST-LABEL: frem_f32(
53+
; FAST: {
54+
; FAST-NEXT: .reg .f32 %f<7>;
55+
; FAST-EMPTY:
56+
; FAST-NEXT: // %bb.0:
57+
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_param_0];
58+
; FAST-NEXT: ld.param.f32 %f2, [frem_f32_param_1];
59+
; FAST-NEXT: div.rn.f32 %f3, %f1, %f2;
60+
; FAST-NEXT: cvt.rzi.f32.f32 %f4, %f3;
61+
; FAST-NEXT: mul.f32 %f5, %f4, %f2;
62+
; FAST-NEXT: sub.f32 %f6, %f1, %f5;
63+
; FAST-NEXT: st.param.f32 [func_retval0], %f6;
64+
; FAST-NEXT: ret;
65+
;
66+
; NORMAL-LABEL: frem_f32(
67+
; NORMAL: {
68+
; NORMAL-NEXT: .reg .pred %p<2>;
69+
; NORMAL-NEXT: .reg .f32 %f<8>;
70+
; NORMAL-EMPTY:
71+
; NORMAL-NEXT: // %bb.0:
72+
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_param_0];
73+
; NORMAL-NEXT: ld.param.f32 %f2, [frem_f32_param_1];
74+
; NORMAL-NEXT: div.rn.f32 %f3, %f1, %f2;
75+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f4, %f3;
76+
; NORMAL-NEXT: mul.f32 %f5, %f4, %f2;
77+
; NORMAL-NEXT: sub.f32 %f6, %f1, %f5;
78+
; NORMAL-NEXT: testp.infinite.f32 %p1, %f2;
79+
; NORMAL-NEXT: selp.f32 %f7, %f1, %f6, %p1;
80+
; NORMAL-NEXT: st.param.f32 [func_retval0], %f7;
81+
; NORMAL-NEXT: ret;
82+
%r = frem float %a, %b
83+
ret float %r
84+
}
85+
86+
define double @frem_f64(double %a, double %b) {
87+
; FAST-LABEL: frem_f64(
88+
; FAST: {
89+
; FAST-NEXT: .reg .f64 %fd<7>;
90+
; FAST-EMPTY:
91+
; FAST-NEXT: // %bb.0:
92+
; FAST-NEXT: ld.param.f64 %fd1, [frem_f64_param_0];
93+
; FAST-NEXT: ld.param.f64 %fd2, [frem_f64_param_1];
94+
; FAST-NEXT: div.rn.f64 %fd3, %fd1, %fd2;
95+
; FAST-NEXT: cvt.rzi.f64.f64 %fd4, %fd3;
96+
; FAST-NEXT: mul.f64 %fd5, %fd4, %fd2;
97+
; FAST-NEXT: sub.f64 %fd6, %fd1, %fd5;
98+
; FAST-NEXT: st.param.f64 [func_retval0], %fd6;
99+
; FAST-NEXT: ret;
100+
;
101+
; NORMAL-LABEL: frem_f64(
102+
; NORMAL: {
103+
; NORMAL-NEXT: .reg .pred %p<2>;
104+
; NORMAL-NEXT: .reg .f64 %fd<8>;
105+
; NORMAL-EMPTY:
106+
; NORMAL-NEXT: // %bb.0:
107+
; NORMAL-NEXT: ld.param.f64 %fd1, [frem_f64_param_0];
108+
; NORMAL-NEXT: ld.param.f64 %fd2, [frem_f64_param_1];
109+
; NORMAL-NEXT: div.rn.f64 %fd3, %fd1, %fd2;
110+
; NORMAL-NEXT: cvt.rzi.f64.f64 %fd4, %fd3;
111+
; NORMAL-NEXT: mul.f64 %fd5, %fd4, %fd2;
112+
; NORMAL-NEXT: sub.f64 %fd6, %fd1, %fd5;
113+
; NORMAL-NEXT: testp.infinite.f64 %p1, %fd2;
114+
; NORMAL-NEXT: selp.f64 %fd7, %fd1, %fd6, %p1;
115+
; NORMAL-NEXT: st.param.f64 [func_retval0], %fd7;
116+
; NORMAL-NEXT: ret;
117+
%r = frem double %a, %b
118+
ret double %r
119+
}
120+
121+
define half @frem_f16_ninf(half %a, half %b) {
122+
; FAST-LABEL: frem_f16_ninf(
123+
; FAST: {
124+
; FAST-NEXT: .reg .b16 %rs<4>;
125+
; FAST-NEXT: .reg .f32 %f<7>;
126+
; FAST-EMPTY:
127+
; FAST-NEXT: // %bb.0:
128+
; FAST-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_param_0];
129+
; FAST-NEXT: ld.param.b16 %rs2, [frem_f16_ninf_param_1];
130+
; FAST-NEXT: cvt.f32.f16 %f1, %rs2;
131+
; FAST-NEXT: cvt.f32.f16 %f2, %rs1;
132+
; FAST-NEXT: div.rn.f32 %f3, %f2, %f1;
133+
; FAST-NEXT: cvt.rzi.f32.f32 %f4, %f3;
134+
; FAST-NEXT: mul.f32 %f5, %f4, %f1;
135+
; FAST-NEXT: sub.f32 %f6, %f2, %f5;
136+
; FAST-NEXT: cvt.rn.f16.f32 %rs3, %f6;
137+
; FAST-NEXT: st.param.b16 [func_retval0], %rs3;
138+
; FAST-NEXT: ret;
139+
;
140+
; NORMAL-LABEL: frem_f16_ninf(
141+
; NORMAL: {
142+
; NORMAL-NEXT: .reg .pred %p<2>;
143+
; NORMAL-NEXT: .reg .b16 %rs<4>;
144+
; NORMAL-NEXT: .reg .f32 %f<8>;
145+
; NORMAL-EMPTY:
146+
; NORMAL-NEXT: // %bb.0:
147+
; NORMAL-NEXT: ld.param.b16 %rs1, [frem_f16_ninf_param_0];
148+
; NORMAL-NEXT: ld.param.b16 %rs2, [frem_f16_ninf_param_1];
149+
; NORMAL-NEXT: cvt.f32.f16 %f1, %rs2;
150+
; NORMAL-NEXT: cvt.f32.f16 %f2, %rs1;
151+
; NORMAL-NEXT: div.rn.f32 %f3, %f2, %f1;
152+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f4, %f3;
153+
; NORMAL-NEXT: mul.f32 %f5, %f4, %f1;
154+
; NORMAL-NEXT: sub.f32 %f6, %f2, %f5;
155+
; NORMAL-NEXT: testp.infinite.f32 %p1, %f1;
156+
; NORMAL-NEXT: selp.f32 %f7, %f2, %f6, %p1;
157+
; NORMAL-NEXT: cvt.rn.f16.f32 %rs3, %f7;
158+
; NORMAL-NEXT: st.param.b16 [func_retval0], %rs3;
159+
; NORMAL-NEXT: ret;
160+
%r = frem ninf half %a, %b
161+
ret half %r
162+
}
163+
164+
define float @frem_f32_ninf(float %a, float %b) {
165+
; FAST-LABEL: frem_f32_ninf(
166+
; FAST: {
167+
; FAST-NEXT: .reg .f32 %f<7>;
168+
; FAST-EMPTY:
169+
; FAST-NEXT: // %bb.0:
170+
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_ninf_param_0];
171+
; FAST-NEXT: ld.param.f32 %f2, [frem_f32_ninf_param_1];
172+
; FAST-NEXT: div.rn.f32 %f3, %f1, %f2;
173+
; FAST-NEXT: cvt.rzi.f32.f32 %f4, %f3;
174+
; FAST-NEXT: mul.f32 %f5, %f4, %f2;
175+
; FAST-NEXT: sub.f32 %f6, %f1, %f5;
176+
; FAST-NEXT: st.param.f32 [func_retval0], %f6;
177+
; FAST-NEXT: ret;
178+
;
179+
; NORMAL-LABEL: frem_f32_ninf(
180+
; NORMAL: {
181+
; NORMAL-NEXT: .reg .pred %p<2>;
182+
; NORMAL-NEXT: .reg .f32 %f<8>;
183+
; NORMAL-EMPTY:
184+
; NORMAL-NEXT: // %bb.0:
185+
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_ninf_param_0];
186+
; NORMAL-NEXT: ld.param.f32 %f2, [frem_f32_ninf_param_1];
187+
; NORMAL-NEXT: div.rn.f32 %f3, %f1, %f2;
188+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f4, %f3;
189+
; NORMAL-NEXT: mul.f32 %f5, %f4, %f2;
190+
; NORMAL-NEXT: sub.f32 %f6, %f1, %f5;
191+
; NORMAL-NEXT: testp.infinite.f32 %p1, %f2;
192+
; NORMAL-NEXT: selp.f32 %f7, %f1, %f6, %p1;
193+
; NORMAL-NEXT: st.param.f32 [func_retval0], %f7;
194+
; NORMAL-NEXT: ret;
195+
%r = frem ninf float %a, %b
196+
ret float %r
197+
}
198+
199+
define double @frem_f64_ninf(double %a, double %b) {
200+
; FAST-LABEL: frem_f64_ninf(
201+
; FAST: {
202+
; FAST-NEXT: .reg .f64 %fd<7>;
203+
; FAST-EMPTY:
204+
; FAST-NEXT: // %bb.0:
205+
; FAST-NEXT: ld.param.f64 %fd1, [frem_f64_ninf_param_0];
206+
; FAST-NEXT: ld.param.f64 %fd2, [frem_f64_ninf_param_1];
207+
; FAST-NEXT: div.rn.f64 %fd3, %fd1, %fd2;
208+
; FAST-NEXT: cvt.rzi.f64.f64 %fd4, %fd3;
209+
; FAST-NEXT: mul.f64 %fd5, %fd4, %fd2;
210+
; FAST-NEXT: sub.f64 %fd6, %fd1, %fd5;
211+
; FAST-NEXT: st.param.f64 [func_retval0], %fd6;
212+
; FAST-NEXT: ret;
213+
;
214+
; NORMAL-LABEL: frem_f64_ninf(
215+
; NORMAL: {
216+
; NORMAL-NEXT: .reg .pred %p<2>;
217+
; NORMAL-NEXT: .reg .f64 %fd<8>;
218+
; NORMAL-EMPTY:
219+
; NORMAL-NEXT: // %bb.0:
220+
; NORMAL-NEXT: ld.param.f64 %fd1, [frem_f64_ninf_param_0];
221+
; NORMAL-NEXT: ld.param.f64 %fd2, [frem_f64_ninf_param_1];
222+
; NORMAL-NEXT: div.rn.f64 %fd3, %fd1, %fd2;
223+
; NORMAL-NEXT: cvt.rzi.f64.f64 %fd4, %fd3;
224+
; NORMAL-NEXT: mul.f64 %fd5, %fd4, %fd2;
225+
; NORMAL-NEXT: sub.f64 %fd6, %fd1, %fd5;
226+
; NORMAL-NEXT: testp.infinite.f64 %p1, %fd2;
227+
; NORMAL-NEXT: selp.f64 %fd7, %fd1, %fd6, %p1;
228+
; NORMAL-NEXT: st.param.f64 [func_retval0], %fd7;
229+
; NORMAL-NEXT: ret;
230+
%r = frem ninf double %a, %b
231+
ret double %r
232+
}
233+
234+
define float @frem_f32_imm1(float %a) {
235+
; FAST-LABEL: frem_f32_imm1(
236+
; FAST: {
237+
; FAST-NEXT: .reg .f32 %f<6>;
238+
; FAST-EMPTY:
239+
; FAST-NEXT: // %bb.0:
240+
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_imm1_param_0];
241+
; FAST-NEXT: div.rn.f32 %f2, %f1, 0f40E00000;
242+
; FAST-NEXT: cvt.rzi.f32.f32 %f3, %f2;
243+
; FAST-NEXT: mul.f32 %f4, %f3, 0f40E00000;
244+
; FAST-NEXT: sub.f32 %f5, %f1, %f4;
245+
; FAST-NEXT: st.param.f32 [func_retval0], %f5;
246+
; FAST-NEXT: ret;
247+
;
248+
; NORMAL-LABEL: frem_f32_imm1(
249+
; NORMAL: {
250+
; NORMAL-NEXT: .reg .pred %p<2>;
251+
; NORMAL-NEXT: .reg .f32 %f<7>;
252+
; NORMAL-EMPTY:
253+
; NORMAL-NEXT: // %bb.0:
254+
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_imm1_param_0];
255+
; NORMAL-NEXT: div.rn.f32 %f2, %f1, 0f40E00000;
256+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f3, %f2;
257+
; NORMAL-NEXT: mul.f32 %f4, %f3, 0f40E00000;
258+
; NORMAL-NEXT: sub.f32 %f5, %f1, %f4;
259+
; NORMAL-NEXT: testp.infinite.f32 %p1, 0f40E00000;
260+
; NORMAL-NEXT: selp.f32 %f6, %f1, %f5, %p1;
261+
; NORMAL-NEXT: st.param.f32 [func_retval0], %f6;
262+
; NORMAL-NEXT: ret;
263+
%r = frem float %a, 7.0
264+
ret float %r
265+
}
266+
267+
define float @frem_f32_imm2(float %a) {
268+
; FAST-LABEL: frem_f32_imm2(
269+
; FAST: {
270+
; FAST-NEXT: .reg .f32 %f<7>;
271+
; FAST-EMPTY:
272+
; FAST-NEXT: // %bb.0:
273+
; FAST-NEXT: ld.param.f32 %f1, [frem_f32_imm2_param_0];
274+
; FAST-NEXT: mov.b32 %f2, 0f40E00000;
275+
; FAST-NEXT: div.rn.f32 %f3, %f2, %f1;
276+
; FAST-NEXT: cvt.rzi.f32.f32 %f4, %f3;
277+
; FAST-NEXT: mul.f32 %f5, %f4, %f1;
278+
; FAST-NEXT: sub.f32 %f6, %f2, %f5;
279+
; FAST-NEXT: st.param.f32 [func_retval0], %f6;
280+
; FAST-NEXT: ret;
281+
;
282+
; NORMAL-LABEL: frem_f32_imm2(
283+
; NORMAL: {
284+
; NORMAL-NEXT: .reg .pred %p<2>;
285+
; NORMAL-NEXT: .reg .f32 %f<8>;
286+
; NORMAL-EMPTY:
287+
; NORMAL-NEXT: // %bb.0:
288+
; NORMAL-NEXT: ld.param.f32 %f1, [frem_f32_imm2_param_0];
289+
; NORMAL-NEXT: mov.b32 %f2, 0f40E00000;
290+
; NORMAL-NEXT: div.rn.f32 %f3, %f2, %f1;
291+
; NORMAL-NEXT: cvt.rzi.f32.f32 %f4, %f3;
292+
; NORMAL-NEXT: mul.f32 %f5, %f4, %f1;
293+
; NORMAL-NEXT: sub.f32 %f6, %f2, %f5;
294+
; NORMAL-NEXT: testp.infinite.f32 %p1, %f1;
295+
; NORMAL-NEXT: selp.f32 %f7, %f2, %f6, %p1;
296+
; NORMAL-NEXT: st.param.f32 [func_retval0], %f7;
297+
; NORMAL-NEXT: ret;
298+
%r = frem float 7.0, %a
299+
ret float %r
300+
}

0 commit comments

Comments
 (0)