diff --git a/clang/lib/DPCT/Diagnostics/Diagnostics.inc b/clang/lib/DPCT/Diagnostics/Diagnostics.inc index 273bf19fc699..c1f4a1effac6 100644 --- a/clang/lib/DPCT/Diagnostics/Diagnostics.inc +++ b/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, currently 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, currently 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 diff --git a/clang/lib/DPCT/RulesLang/APINamesMath.inc b/clang/lib/DPCT/RulesLang/APINamesMath.inc index a8707d92c1db..92a21edc0032 100644 --- a/clang/lib/DPCT/RulesLang/APINamesMath.inc +++ b/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") diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp index a09411383d82..cb3a928fa6d6 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp +++ b/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", diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index 3839ec85263b..8f17aaa8d9a2 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/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") diff --git a/clang/test/dpct/help_option_check/lin/help_all.txt b/clang/test/dpct/help_option_check/lin/help_all.txt index bc535d587eec..0a759cf82505 100644 --- a/clang/test/dpct/help_option_check/lin/help_all.txt +++ b/clang/test/dpct/help_option_check/lin/help_all.txt @@ -128,7 +128,7 @@ All DPCT options --rule-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= - 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= - Specify the extension of migrated source file(s). diff --git a/clang/test/dpct/help_option_check/lin/help_option_check.cpp b/clang/test/dpct/help_option_check/lin/help_option_check.cpp index 7ee2698c29d6..88b4f1c3e6be 100644 --- a/clang/test/dpct/help_option_check/lin/help_option_check.cpp +++ b/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 diff --git a/clang/test/dpct/help_option_check/win/help_all.txt b/clang/test/dpct/help_option_check/win/help_all.txt index 7beb74a521c1..7887c4a5a5dc 100644 --- a/clang/test/dpct/help_option_check/win/help_all.txt +++ b/clang/test/dpct/help_option_check/win/help_all.txt @@ -127,7 +127,7 @@ All DPCT options --rule-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= - 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= - Specify the extension of migrated source file(s). diff --git a/clang/test/dpct/help_option_check/win/help_option_check.cpp b/clang/test/dpct/help_option_check/win/help_option_check.cpp index b4d8e3c60181..e87b7f897d0c 100644 --- a/clang/test/dpct/help_option_check/win/help_option_check.cpp +++ b/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 diff --git a/clang/test/dpct/math/cuda-math-extension.cu b/clang/test/dpct/math/cuda-math-extension.cu index 24b2cec16b1c..80f03703bcd3 100644 --- a/clang/test/dpct/math/cuda-math-extension.cu +++ b/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); diff --git a/clang/test/dpct/math/cuda-math-intrinsics.cu b/clang/test/dpct/math/cuda-math-intrinsics.cu index 7bfcfbddfd47..3ac5d4517eb4 100644 --- a/clang/test/dpct/math/cuda-math-intrinsics.cu +++ b/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: */