Skip to content

Commit 399475c

Browse files
zhiweij1zhiweij1
committed
[SYCLomatic] Migrate __fmaf_ieee_r* to sycl::ext::intel::math::fmaf_r*
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent 8513308 commit 399475c

5 files changed

Lines changed: 138 additions & 4 deletions

File tree

‎clang/lib/DPCT/RulesLang/APINamesMath.inc‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -116,6 +116,10 @@ ENTRY_REWRITE("__fmaf_rd")
116116
ENTRY_REWRITE("__fmaf_rn")
117117
ENTRY_REWRITE("__fmaf_ru")
118118
ENTRY_REWRITE("__fmaf_rz")
119+
ENTRY_REWRITE("__fmaf_ieee_rd")
120+
ENTRY_REWRITE("__fmaf_ieee_rn")
121+
ENTRY_REWRITE("__fmaf_ieee_ru")
122+
ENTRY_REWRITE("__fmaf_ieee_rz")
119123
ENTRY_RENAMED_SINGLE("__frcp_rd", MapNames::getClNamespace(false, true) + "native::recip")
120124
ENTRY_RENAMED_SINGLE("__frcp_rn", MapNames::getClNamespace(false, true) + "native::recip")
121125
ENTRY_RENAMED_SINGLE("__frcp_ru", MapNames::getClNamespace(false, true) + "native::recip")

‎clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp‎

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -170,6 +170,86 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
170170
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
171171
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
172172
Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
173+
// __fmaf_ieee_rd
174+
MATH_API_REWRITERS_V2(
175+
"__fmaf_ieee_rd",
176+
MATH_API_REWRITER_PAIR(
177+
math::Tag::math_libdevice,
178+
CALL_FACTORY_ENTRY(
179+
"__fmaf_ieee_rd",
180+
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd",
181+
ARG(0), ARG(1), ARG(2)))),
182+
MATH_API_REWRITER_PAIR(
183+
math::Tag::emulation,
184+
WARNING_FACTORY_ENTRY(
185+
"__fmaf_ieee_rd",
186+
CALL_FACTORY_ENTRY(
187+
"__fmaf_ieee_rd",
188+
CALL(MapNames::getClNamespace(false, true) + "fma",
189+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
190+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
191+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
192+
Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
193+
// __fmaf_ieee_rn
194+
MATH_API_REWRITERS_V2(
195+
"__fmaf_ieee_rn",
196+
MATH_API_REWRITER_PAIR(
197+
math::Tag::math_libdevice,
198+
CALL_FACTORY_ENTRY(
199+
"__fmaf_ieee_rn",
200+
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn",
201+
ARG(0), ARG(1), ARG(2)))),
202+
MATH_API_REWRITER_PAIR(
203+
math::Tag::emulation,
204+
WARNING_FACTORY_ENTRY(
205+
"__fmaf_ieee_rn",
206+
CALL_FACTORY_ENTRY(
207+
"__fmaf_ieee_rn",
208+
CALL(MapNames::getClNamespace(false, true) + "fma",
209+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
210+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
211+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
212+
Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
213+
// __fmaf_ieee_ru
214+
MATH_API_REWRITERS_V2(
215+
"__fmaf_ieee_ru",
216+
MATH_API_REWRITER_PAIR(
217+
math::Tag::math_libdevice,
218+
CALL_FACTORY_ENTRY(
219+
"__fmaf_ieee_ru",
220+
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru",
221+
ARG(0), ARG(1), ARG(2)))),
222+
MATH_API_REWRITER_PAIR(
223+
math::Tag::emulation,
224+
WARNING_FACTORY_ENTRY(
225+
"__fmaf_ieee_ru",
226+
CALL_FACTORY_ENTRY(
227+
"__fmaf_ieee_ru",
228+
CALL(MapNames::getClNamespace(false, true) + "fma",
229+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
230+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
231+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
232+
Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
233+
// __fmaf_ieee_rz
234+
MATH_API_REWRITERS_V2(
235+
"__fmaf_ieee_rz",
236+
MATH_API_REWRITER_PAIR(
237+
math::Tag::math_libdevice,
238+
CALL_FACTORY_ENTRY(
239+
"__fmaf_ieee_rz",
240+
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz",
241+
ARG(0), ARG(1), ARG(2)))),
242+
MATH_API_REWRITER_PAIR(
243+
math::Tag::emulation,
244+
WARNING_FACTORY_ENTRY(
245+
"__fmaf_ieee_rz",
246+
CALL_FACTORY_ENTRY(
247+
"__fmaf_ieee_rz",
248+
CALL(MapNames::getClNamespace(false, true) + "fma",
249+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
250+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
251+
CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
252+
Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
173253
// __fmul_rd
174254
MATH_API_REWRITER_DEVICE(
175255
"__fmul_rd",

‎clang/lib/DPCT/SrcAPI/APINames.inc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1097,10 +1097,10 @@ ENTRY(__fdiv_rn, __fdiv_rn, true, NO_FLAG, P4, "Successful: DPCT1013")
10971097
ENTRY(__fdiv_ru, __fdiv_ru, true, NO_FLAG, P0, "Successful: DPCT1013")
10981098
ENTRY(__fdiv_rz, __fdiv_rz, true, NO_FLAG, P4, "Successful: DPCT1013")
10991099
ENTRY(__fdividef, __fdividef, true, NO_FLAG, P0, "Successful")
1100-
ENTRY(__fmaf_ieee_rd, __fmaf_ieee_rd, false, NO_FLAG, P4, "comment")
1101-
ENTRY(__fmaf_ieee_rn, __fmaf_ieee_rn, false, NO_FLAG, P4, "comment")
1102-
ENTRY(__fmaf_ieee_ru, __fmaf_ieee_ru, false, NO_FLAG, P4, "comment")
1103-
ENTRY(__fmaf_ieee_rz, __fmaf_ieee_rz, false, NO_FLAG, P4, "comment")
1100+
ENTRY(__fmaf_ieee_rd, __fmaf_ieee_rd, true, NO_FLAG, P4, "Successful: DPCT1013")
1101+
ENTRY(__fmaf_ieee_rn, __fmaf_ieee_rn, true, NO_FLAG, P4, "Successful: DPCT1013")
1102+
ENTRY(__fmaf_ieee_ru, __fmaf_ieee_ru, true, NO_FLAG, P4, "Successful: DPCT1013")
1103+
ENTRY(__fmaf_ieee_rz, __fmaf_ieee_rz, true, NO_FLAG, P4, "Successful: DPCT1013")
11041104
ENTRY(__fmaf_rd, __fmaf_rd, true, NO_FLAG, P4, "Successful: DPCT1013")
11051105
ENTRY(__fmaf_rn, __fmaf_rn, true, NO_FLAG, P0, "Successful: DPCT1013")
11061106
ENTRY(__fmaf_ru, __fmaf_ru, true, NO_FLAG, P4, "Successful: DPCT1013")

‎clang/test/dpct/math/cuda-math-extension.cu‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -181,6 +181,14 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
181181
f2 = __fmaf_ru(f0, f1, f2);
182182
// CHECK: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
183183
f2 = __fmaf_rz(f0, f1, f2);
184+
// CHECK: f2 = sycl::ext::intel::math::fmaf_rd(f0, f1, f2);
185+
f2 = __fmaf_ieee_rd(f0, f1, f2);
186+
// CHECK: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2);
187+
f2 = __fmaf_ieee_rn(f0, f1, f2);
188+
// CHECK: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2);
189+
f2 = __fmaf_ieee_ru(f0, f1, f2);
190+
// CHECK: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
191+
f2 = __fmaf_ieee_rz(f0, f1, f2);
184192
// CHECK: f2 = sycl::ext::intel::math::fmul_rd(f0, f1);
185193
f2 = __fmul_rd(f0, f1);
186194
// CHECK: f2 = sycl::ext::intel::math::fmul_rn(f0, f1);

‎clang/test/dpct/math/cuda-math-intrinsics.cu‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1339,6 +1339,27 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
13391339
// CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
13401340
f2 = __fmaf_rz(f0, f1, f2);
13411341

1342+
// CHECK: /*
1343+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1344+
// CHECK-NEXT: */
1345+
// CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
1346+
f2 = __fmaf_ieee_rd(f0, f1, f2);
1347+
// CHECK: /*
1348+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1349+
// CHECK-NEXT: */
1350+
// CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
1351+
f2 = __fmaf_ieee_rn(f0, f1, f2);
1352+
// CHECK: /*
1353+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1354+
// CHECK-NEXT: */
1355+
// CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
1356+
f2 = __fmaf_ieee_ru(f0, f1, f2);
1357+
// CHECK: /*
1358+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1359+
// CHECK-NEXT: */
1360+
// CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
1361+
f2 = __fmaf_ieee_rz(f0, f1, f2);
1362+
13421363
// CHECK: /*
13431364
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
13441365
// CHECK-NEXT: */
@@ -1360,6 +1381,27 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
13601381
// CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
13611382
f2 = __fmaf_rz(i, i, i);
13621383

1384+
// CHECK: /*
1385+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1386+
// CHECK-NEXT: */
1387+
// CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
1388+
f2 = __fmaf_ieee_rd(i, i, i);
1389+
// CHECK: /*
1390+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1391+
// CHECK-NEXT: */
1392+
// CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
1393+
f2 = __fmaf_ieee_rn(i, i, i);
1394+
// CHECK: /*
1395+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1396+
// CHECK-NEXT: */
1397+
// CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
1398+
f2 = __fmaf_ieee_ru(i, i, i);
1399+
// CHECK: /*
1400+
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1401+
// CHECK-NEXT: */
1402+
// CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
1403+
f2 = __fmaf_ieee_rz(i, i, i);
1404+
13631405
// CHECK: /*
13641406
// CHECK-NEXT: DPCT1013:{{[0-9]+}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
13651407
// CHECK-NEXT: */

0 commit comments

Comments
 (0)