1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2 // Test with OCML_BASIC_ROUNDED_OPERATIONS
3 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
4 // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
5 // RUN: -internal-isystem %S/Inputs/include \
6 // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
7 // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
8 // RUN: -D__HIPCC_RTC__ -DOCML_BASIC_ROUNDED_OPERATIONS | FileCheck %s
10 // CHECK-LABEL: @test___fadd_rd(
12 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3:[0-9]+]]
13 // CHECK-NEXT: ret float [[CALL_I]]
15 extern "C" __device__ float test___fadd_rd(float x, float y) {
16 return __fadd_rd(x, y);
19 // CHECK-LABEL: @test___fadd_rn(
21 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
22 // CHECK-NEXT: ret float [[CALL_I]]
24 extern "C" __device__ float test___fadd_rn(float x, float y) {
25 return __fadd_rn(x, y);
28 // CHECK-LABEL: @test___fadd_ru(
30 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
31 // CHECK-NEXT: ret float [[CALL_I]]
33 extern "C" __device__ float test___fadd_ru(float x, float y) {
34 return __fadd_ru(x, y);
37 // CHECK-LABEL: @test___fadd_rz(
39 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
40 // CHECK-NEXT: ret float [[CALL_I]]
42 extern "C" __device__ float test___fadd_rz(float x, float y) {
43 return __fadd_rz(x, y);
46 // CHECK-LABEL: @test__fmaf_rd(
48 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
49 // CHECK-NEXT: ret float [[CALL_I]]
51 extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {
52 return __fmaf_rd(x, y, z);
55 // CHECK-LABEL: @test__fmaf_rn(
57 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
58 // CHECK-NEXT: ret float [[CALL_I]]
60 extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
61 return __fmaf_rn(x, y, z);
64 // CHECK-LABEL: @test__fmaf_ru(
66 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
67 // CHECK-NEXT: ret float [[CALL_I]]
69 extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {
70 return __fmaf_ru(x, y, z);
73 // CHECK-LABEL: @test__fmaf_rz(
75 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR3]]
76 // CHECK-NEXT: ret float [[CALL_I]]
78 extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {
79 return __fmaf_rz(x, y, z);
82 // CHECK-LABEL: @test___fmul_rd(
84 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
85 // CHECK-NEXT: ret float [[CALL_I]]
87 extern "C" __device__ float test___fmul_rd(float x, float y) {
88 return __fmul_rd(x, y);
91 // CHECK-LABEL: @test___fmul_rn(
93 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
94 // CHECK-NEXT: ret float [[CALL_I]]
96 extern "C" __device__ float test___fmul_rn(float x, float y) {
97 return __fmul_rn(x, y);
100 // CHECK-LABEL: @test___fmul_ru(
101 // CHECK-NEXT: entry:
102 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
103 // CHECK-NEXT: ret float [[CALL_I]]
105 extern "C" __device__ float test___fmul_ru(float x, float y) {
106 return __fmul_ru(x, y);
109 // CHECK-LABEL: @test___fmul_rz(
110 // CHECK-NEXT: entry:
111 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
112 // CHECK-NEXT: ret float [[CALL_I]]
114 extern "C" __device__ float test___fmul_rz(float x, float y) {
115 return __fmul_rz(x, y);
118 // CHECK-LABEL: @test___frcp_rd(
119 // CHECK-NEXT: entry:
120 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
121 // CHECK-NEXT: ret float [[CALL_I]]
123 extern "C" __device__ float test___frcp_rd(float x) {
127 // CHECK-LABEL: @test___frcp_rn(
128 // CHECK-NEXT: entry:
129 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
130 // CHECK-NEXT: ret float [[CALL_I]]
132 extern "C" __device__ float test___frcp_rn(float x) {
136 // CHECK-LABEL: @test___frcp_ru(
137 // CHECK-NEXT: entry:
138 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
139 // CHECK-NEXT: ret float [[CALL_I]]
141 extern "C" __device__ float test___frcp_ru(float x) {
145 // CHECK-LABEL: @test___frcp_rz(
146 // CHECK-NEXT: entry:
147 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR3]]
148 // CHECK-NEXT: ret float [[CALL_I]]
150 extern "C" __device__ float test___frcp_rz(float x) {
154 // CHECK-LABEL: @test___fsqrt_rd(
155 // CHECK-NEXT: entry:
156 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR3]]
157 // CHECK-NEXT: ret float [[CALL_I]]
159 extern "C" __device__ float test___fsqrt_rd(float x) {
160 return __fsqrt_rd(x);
163 // CHECK-LABEL: @test___fsqrt_rn(
164 // CHECK-NEXT: entry:
165 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR3]]
166 // CHECK-NEXT: ret float [[CALL_I]]
168 extern "C" __device__ float test___fsqrt_rn(float x) {
169 return __fsqrt_rn(x);
172 // CHECK-LABEL: @test___fsqrt_ru(
173 // CHECK-NEXT: entry:
174 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR3]]
175 // CHECK-NEXT: ret float [[CALL_I]]
177 extern "C" __device__ float test___fsqrt_ru(float x) {
178 return __fsqrt_ru(x);
181 // CHECK-LABEL: @test___fsqrt_rz(
182 // CHECK-NEXT: entry:
183 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR3]]
184 // CHECK-NEXT: ret float [[CALL_I]]
186 extern "C" __device__ float test___fsqrt_rz(float x) {
187 return __fsqrt_rz(x);
190 // CHECK-LABEL: @test___fsub_rd(
191 // CHECK-NEXT: entry:
192 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
193 // CHECK-NEXT: ret float [[CALL_I]]
195 extern "C" __device__ float test___fsub_rd(float x, float y) {
196 return __fsub_rd(x, y);
199 // CHECK-LABEL: @test___fsub_rn(
200 // CHECK-NEXT: entry:
201 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
202 // CHECK-NEXT: ret float [[CALL_I]]
204 extern "C" __device__ float test___fsub_rn(float x, float y) {
205 return __fsub_rn(x, y);
208 // CHECK-LABEL: @test___fsub_ru(
209 // CHECK-NEXT: entry:
210 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
211 // CHECK-NEXT: ret float [[CALL_I]]
213 extern "C" __device__ float test___fsub_ru(float x, float y) {
214 return __fsub_ru(x, y);
217 // CHECK-LABEL: @test___fsub_rz(
218 // CHECK-NEXT: entry:
219 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR3]]
220 // CHECK-NEXT: ret float [[CALL_I]]
222 extern "C" __device__ float test___fsub_rz(float x, float y) {
223 return __fsub_rz(x, y);
226 // CHECK-LABEL: @test___dadd_rd(
227 // CHECK-NEXT: entry:
228 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
229 // CHECK-NEXT: ret double [[CALL_I]]
231 extern "C" __device__ double test___dadd_rd(double x, double y) {
232 return __dadd_rd(x, y);
235 // CHECK-LABEL: @test___dadd_rn(
236 // CHECK-NEXT: entry:
237 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
238 // CHECK-NEXT: ret double [[CALL_I]]
240 extern "C" __device__ double test___dadd_rn(double x, double y) {
241 return __dadd_rn(x, y);
244 // CHECK-LABEL: @test___dadd_ru(
245 // CHECK-NEXT: entry:
246 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
247 // CHECK-NEXT: ret double [[CALL_I]]
249 extern "C" __device__ double test___dadd_ru(double x, double y) {
250 return __dadd_ru(x, y);
253 // CHECK-LABEL: @test___dadd_rz(
254 // CHECK-NEXT: entry:
255 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
256 // CHECK-NEXT: ret double [[CALL_I]]
258 extern "C" __device__ double test___dadd_rz(double x, double y) {
259 return __dadd_rz(x, y);
262 // CHECK-LABEL: @test___dmul_rd(
263 // CHECK-NEXT: entry:
264 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
265 // CHECK-NEXT: ret double [[CALL_I]]
267 extern "C" __device__ double test___dmul_rd(double x, double y) {
268 return __dmul_rd(x, y);
271 // CHECK-LABEL: @test___dmul_rn(
272 // CHECK-NEXT: entry:
273 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
274 // CHECK-NEXT: ret double [[CALL_I]]
276 extern "C" __device__ double test___dmul_rn(double x, double y) {
277 return __dmul_rn(x, y);
280 // CHECK-LABEL: @test___dmul_ru(
281 // CHECK-NEXT: entry:
282 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
283 // CHECK-NEXT: ret double [[CALL_I]]
285 extern "C" __device__ double test___dmul_ru(double x, double y) {
286 return __dmul_ru(x, y);
289 // CHECK-LABEL: @test___dmul_rz(
290 // CHECK-NEXT: entry:
291 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR3]]
292 // CHECK-NEXT: ret double [[CALL_I]]
294 extern "C" __device__ double test___dmul_rz(double x, double y) {
295 return __dmul_rz(x, y);
298 // CHECK-LABEL: @test___drcp_rd(
299 // CHECK-NEXT: entry:
300 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
301 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
302 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
303 // CHECK-NEXT: ret float [[CONV1]]
305 extern "C" __device__ float test___drcp_rd(float x) {
309 // CHECK-LABEL: @test___drcp_rn(
310 // CHECK-NEXT: entry:
311 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
312 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
313 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
314 // CHECK-NEXT: ret float [[CONV1]]
316 extern "C" __device__ float test___drcp_rn(float x) {
320 // CHECK-LABEL: @test___drcp_ru(
321 // CHECK-NEXT: entry:
322 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
323 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
324 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
325 // CHECK-NEXT: ret float [[CONV1]]
327 extern "C" __device__ float test___drcp_ru(float x) {
331 // CHECK-LABEL: @test___drcp_rz(
332 // CHECK-NEXT: entry:
333 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
334 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR3]]
335 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
336 // CHECK-NEXT: ret float [[CONV1]]
338 extern "C" __device__ float test___drcp_rz(float x) {
342 // CHECK-LABEL: @test___dsqrt_rd(
343 // CHECK-NEXT: entry:
344 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
345 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR3]]
346 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
347 // CHECK-NEXT: ret float [[CONV1]]
349 extern "C" __device__ float test___dsqrt_rd(float x) {
350 return __dsqrt_rd(x);
353 // CHECK-LABEL: @test___dsqrt_rn(
354 // CHECK-NEXT: entry:
355 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
356 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR3]]
357 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
358 // CHECK-NEXT: ret float [[CONV1]]
360 extern "C" __device__ float test___dsqrt_rn(float x) {
361 return __dsqrt_rn(x);
364 // CHECK-LABEL: @test___dsqrt_ru(
365 // CHECK-NEXT: entry:
366 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
367 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR3]]
368 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
369 // CHECK-NEXT: ret float [[CONV1]]
371 extern "C" __device__ float test___dsqrt_ru(float x) {
372 return __dsqrt_ru(x);
375 // CHECK-LABEL: @test___dsqrt_rz(
376 // CHECK-NEXT: entry:
377 // CHECK-NEXT: [[CONV:%.*]] = fpext float [[X:%.*]] to double
378 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR3]]
379 // CHECK-NEXT: [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
380 // CHECK-NEXT: ret float [[CONV1]]
382 extern "C" __device__ float test___dsqrt_rz(float x) {
383 return __dsqrt_rz(x);
386 // CHECK-LABEL: @test__fma_rd(
387 // CHECK-NEXT: entry:
388 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
389 // CHECK-NEXT: ret double [[CALL_I]]
391 extern "C" __device__ double test__fma_rd(double x, double y, double z) {
392 return __fma_rd(x, y, z);
395 // CHECK-LABEL: @test__fma_rn(
396 // CHECK-NEXT: entry:
397 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
398 // CHECK-NEXT: ret double [[CALL_I]]
400 extern "C" __device__ double test__fma_rn(double x, double y, double z) {
401 return __fma_rn(x, y, z);
404 // CHECK-LABEL: @test__fma_ru(
405 // CHECK-NEXT: entry:
406 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
407 // CHECK-NEXT: ret double [[CALL_I]]
409 extern "C" __device__ double test__fma_ru(double x, double y, double z) {
410 return __fma_ru(x, y, z);
413 // CHECK-LABEL: @test__fma_rz(
414 // CHECK-NEXT: entry:
415 // CHECK-NEXT: [[CALL_I:%.*]] = tail call contract noundef double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR3]]
416 // CHECK-NEXT: ret double [[CALL_I]]
418 extern "C" __device__ double test__fma_rz(double x, double y, double z) {
419 return __fma_rz(x, y, z);