[SYCLomatic] Migrate __fmaf_ieee_r* to sycl::ext::intel::math::fmaf_r*

clang/lib/DPCT/Diagnostics/Diagnostics.inc

@@ -300,6 +300,8 @@ DEF_WARNING(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Imag
 DEF_COMMENT(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource using NT handle on Windows. If assert({0}.get_win32_handle()) fails, you may need to adjust the code to use ({0}.get_win32_handle()).")
 DEF_WARNING(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"%0\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
 DEF_COMMENT(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"{0}\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
+DEF_WARNING(FTZ_BEHAVIOR, 1138, MEDIUM_LEVEL, "The API %0 ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.")
+DEF_COMMENT(FTZ_BEHAVIOR, 1138, MEDIUM_LEVEL, "The API {0} ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.")
 
 // clang-format on
 

clang/lib/DPCT/RulesLang/APINamesMath.inc

@@ -116,6 +116,10 @@ ENTRY_REWRITE("__fmaf_rd")
 ENTRY_REWRITE("__fmaf_rn")
 ENTRY_REWRITE("__fmaf_ru")
 ENTRY_REWRITE("__fmaf_rz")
+ENTRY_REWRITE("__fmaf_ieee_rd")
+ENTRY_REWRITE("__fmaf_ieee_rn")
+ENTRY_REWRITE("__fmaf_ieee_ru")
+ENTRY_REWRITE("__fmaf_ieee_rz")
 ENTRY_RENAMED_SINGLE("__frcp_rd", MapNames::getClNamespace(false, true) + "native::recip")
 ENTRY_RENAMED_SINGLE("__frcp_rn", MapNames::getClNamespace(false, true) + "native::recip")
 ENTRY_RENAMED_SINGLE("__frcp_ru", MapNames::getClNamespace(false, true) + "native::recip")

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

@@ -170,6 +170,98 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
                            CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
                            CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
                   Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
+      // __fmaf_ieee_rd
+      MATH_API_REWRITERS_V2(
+          "__fmaf_ieee_rd",
+          MATH_API_REWRITER_PAIR(
+              math::Tag::math_libdevice,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rd",
+                  CALL_FACTORY_ENTRY("__fmaf_ieee_rd",
+                                     CALL(MapNames::getClNamespace() +
+                                              "ext::intel::math::fmaf_rd",
+                                          ARG(0), ARG(1), ARG(2))),
+                  Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rd"))),
+          MATH_API_REWRITER_PAIR(
+              math::Tag::emulation,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rd",
+                  CALL_FACTORY_ENTRY(
+                      "__fmaf_ieee_rd",
+                      CALL(MapNames::getClNamespace(false, true) + "fma",
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
+                  Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
+      // __fmaf_ieee_rn
+      MATH_API_REWRITERS_V2(
+          "__fmaf_ieee_rn",
+          MATH_API_REWRITER_PAIR(
+              math::Tag::math_libdevice,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rn",
+                  CALL_FACTORY_ENTRY("__fmaf_ieee_rn",
+                                     CALL(MapNames::getClNamespace() +
+                                              "ext::intel::math::fmaf_rn",
+                                          ARG(0), ARG(1), ARG(2))),
+                  Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rn"))),
+          MATH_API_REWRITER_PAIR(
+              math::Tag::emulation,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rn",
+                  CALL_FACTORY_ENTRY(
+                      "__fmaf_ieee_rn",
+                      CALL(MapNames::getClNamespace(false, true) + "fma",
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
+                  Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
+      // __fmaf_ieee_ru
+      MATH_API_REWRITERS_V2(
+          "__fmaf_ieee_ru",
+          MATH_API_REWRITER_PAIR(
+              math::Tag::math_libdevice,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_ru",
+                  CALL_FACTORY_ENTRY("__fmaf_ieee_ru",
+                                     CALL(MapNames::getClNamespace() +
+                                              "ext::intel::math::fmaf_ru",
+                                          ARG(0), ARG(1), ARG(2))),
+                  Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_ru"))),
+          MATH_API_REWRITER_PAIR(
+              math::Tag::emulation,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_ru",
+                  CALL_FACTORY_ENTRY(
+                      "__fmaf_ieee_ru",
+                      CALL(MapNames::getClNamespace(false, true) + "fma",
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
+                  Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
+      // __fmaf_ieee_rz
+      MATH_API_REWRITERS_V2(
+          "__fmaf_ieee_rz",
+          MATH_API_REWRITER_PAIR(
+              math::Tag::math_libdevice,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rz",
+                  CALL_FACTORY_ENTRY("__fmaf_ieee_rz",
+                                     CALL(MapNames::getClNamespace() +
+                                              "ext::intel::math::fmaf_rz",
+                                          ARG(0), ARG(1), ARG(2))),
+                  Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rz"))),
+          MATH_API_REWRITER_PAIR(
+              math::Tag::emulation,
+              WARNING_FACTORY_ENTRY(
+                  "__fmaf_ieee_rz",
+                  CALL_FACTORY_ENTRY(
+                      "__fmaf_ieee_rz",
+                      CALL(MapNames::getClNamespace(false, true) + "fma",
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(0)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(1)),
+                           CAST_IF_NOT_SAME(makeLiteral("float"), ARG(2)))),
+                  Diagnostics::ROUNDING_MODE_UNSUPPORTED)))
       // __fmul_rd
       MATH_API_REWRITER_DEVICE(
           "__fmul_rd",

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -1097,10 +1097,10 @@ ENTRY(__fdiv_rn, __fdiv_rn, true, NO_FLAG, P4, "Successful:  DPCT1013")
 ENTRY(__fdiv_ru, __fdiv_ru, true, NO_FLAG, P0, "Successful:  DPCT1013")
 ENTRY(__fdiv_rz, __fdiv_rz, true, NO_FLAG, P4, "Successful:  DPCT1013")
 ENTRY(__fdividef, __fdividef, true, NO_FLAG, P0, "Successful")
-ENTRY(__fmaf_ieee_rd, __fmaf_ieee_rd, false, NO_FLAG, P4, "comment")
-ENTRY(__fmaf_ieee_rn, __fmaf_ieee_rn, false, NO_FLAG, P4, "comment")
-ENTRY(__fmaf_ieee_ru, __fmaf_ieee_ru, false, NO_FLAG, P4, "comment")
-ENTRY(__fmaf_ieee_rz, __fmaf_ieee_rz, false, NO_FLAG, P4, "comment")
+ENTRY(__fmaf_ieee_rd, __fmaf_ieee_rd, true, NO_FLAG, P4, "Successful:  DPCT1013")
+ENTRY(__fmaf_ieee_rn, __fmaf_ieee_rn, true, NO_FLAG, P4, "Successful:  DPCT1013")
+ENTRY(__fmaf_ieee_ru, __fmaf_ieee_ru, true, NO_FLAG, P4, "Successful:  DPCT1013")
+ENTRY(__fmaf_ieee_rz, __fmaf_ieee_rz, true, NO_FLAG, P4, "Successful:  DPCT1013")
 ENTRY(__fmaf_rd, __fmaf_rd, true, NO_FLAG, P4, "Successful:  DPCT1013")
 ENTRY(__fmaf_rn, __fmaf_rn, true, NO_FLAG, P0, "Successful:  DPCT1013")
 ENTRY(__fmaf_ru, __fmaf_ru, true, NO_FLAG, P4, "Successful:  DPCT1013")

clang/test/dpct/help_option_check/lin/help_all.txt

@@ -128,7 +128,7 @@ All DPCT options
   --rule-file=<file>                      - Specify the rule file for migration. Also, reference the predefined rules in the "extensions" directory in the root folder of the tool.
   --stop-on-parse-err                     - Stop migration and generation of reports if parsing errors happened. Default: off. 
   --suppress-warnings=<value>             - A comma separated list of migration warnings to suppress. Valid warning IDs range
-                                            from 1000 to 1137. Hyphen separated ranges are also allowed. For example:
+                                            from 1000 to 1138. Hyphen separated ranges are also allowed. For example:
                                             --suppress-warnings=1000-1010,1011.
   --suppress-warnings-all                 - Suppress all migration warnings. Default: off.
   --sycl-file-extension=<value>           - Specify the extension of migrated source file(s).

clang/test/dpct/help_option_check/lin/help_option_check.cpp

@@ -4,8 +4,8 @@
 // RUN: cd %T/help_option_check
 
 // RUN: dpct --help > output.txt
-// RUN: diff --strip-trailing-cr %S/help_all.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_all.txt %T/help_option_check/output.txt
 // RUN: dpct --help=basic > output.txt
-// RUN: diff --strip-trailing-cr %S/help_basic.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_basic.txt %T/help_option_check/output.txt
 // RUN: dpct --help=advanced > output.txt
-// RUN: diff --strip-trailing-cr %S/help_advanced.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_advanced.txt %T/help_option_check/output.txt

clang/test/dpct/help_option_check/win/help_all.txt

@@ -127,7 +127,7 @@ All DPCT options
   --rule-file=<file>                      - Specify the rule file for migration. Also, reference the predefined rules in the "extensions" directory in the root folder of the tool.
   --stop-on-parse-err                     - Stop migration and generation of reports if parsing errors happened. Default: off. 
   --suppress-warnings=<value>             - A comma separated list of migration warnings to suppress. Valid warning IDs range
-                                            from 1000 to 1137. Hyphen separated ranges are also allowed. For example:
+                                            from 1000 to 1138. Hyphen separated ranges are also allowed. For example:
                                             --suppress-warnings=1000-1010,1011.
   --suppress-warnings-all                 - Suppress all migration warnings. Default: off.
   --sycl-file-extension=<value>           - Specify the extension of migrated source file(s).

clang/test/dpct/help_option_check/win/help_option_check.cpp

@@ -4,8 +4,8 @@
 // RUN: cd %T/help_option_check
 
 // RUN: dpct --help > output.txt
-// RUN: diff --strip-trailing-cr %S/help_all.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_all.txt %T/help_option_check/output.txt
 // RUN: dpct --help=basic > output.txt
-// RUN: diff --strip-trailing-cr %S/help_basic.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_basic.txt %T/help_option_check/output.txt
 // RUN: dpct --help=advanced > output.txt
-// RUN: diff --strip-trailing-cr %S/help_advanced.txt %T/help_option_check/output.txt >> %T/diff.txt
+// RUN: diff --strip-trailing-cr %S/help_advanced.txt %T/help_option_check/output.txt

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

@@ -181,6 +181,25 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
   f2 = __fmaf_ru(f0, f1, f2);
   // CHECK: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
   f2 = __fmaf_rz(f0, f1, f2);
+  // CHECK: DPCT1138:{{[0-9]+}}: The API __fmaf_ieee_rd ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rd(f0, f1, f2);
+  f2 = __fmaf_ieee_rd(f0, f1, f2);
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The API __fmaf_ieee_rn ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2);
+  f2 = __fmaf_ieee_rn(f0, f1, f2);
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The API __fmaf_ieee_ru ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2);
+  f2 = __fmaf_ieee_ru(f0, f1, f2);
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The API __fmaf_ieee_rz ignores the flush-to-zero compiler option when handling denormalized data, while it impacts the migrated SYCL API. Please verify the correctness of the migrated code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
+  f2 = __fmaf_ieee_rz(f0, f1, f2);
   // CHECK: f2 = sycl::ext::intel::math::fmul_rd(f0, f1);
   f2 = __fmul_rd(f0, f1);
   // CHECK: f2 = sycl::ext::intel::math::fmul_rn(f0, f1);

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

@@ -1339,6 +1339,27 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
   // CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
   f2 = __fmaf_rz(f0, f1, f2);
 
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
+  f2 = __fmaf_ieee_rd(f0, f1, f2);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
+  f2 = __fmaf_ieee_rn(f0, f1, f2);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
+  f2 = __fmaf_ieee_ru(f0, f1, f2);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma(f0, f1, f2);
+  f2 = __fmaf_ieee_rz(f0, f1, f2);
+
   // CHECK: /*
   // 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.
   // CHECK-NEXT: */
@@ -1360,6 +1381,27 @@ __global__ void kernelFuncFloat(float *deviceArrayFloat) {
   // CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
   f2 = __fmaf_rz(i, i, i);
 
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
+  f2 = __fmaf_ieee_rd(i, i, i);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
+  f2 = __fmaf_ieee_rn(i, i, i);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
+  f2 = __fmaf_ieee_ru(i, i, i);
+  // CHECK: /*
+  // 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.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: f2 = sycl::fma((float)i, (float)i, (float)i);
+  f2 = __fmaf_ieee_rz(i, i, i);
+
   // CHECK: /*
   // 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.
   // CHECK-NEXT: */