From 399475c07e82a5dc291e0d6d670398b29a4ed2b1 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Fri, 25 Apr 2025 13:53:01 +0800 Subject: [PATCH 1/6] [SYCLomatic] Migrate __fmaf_ieee_r* to sycl::ext::intel::math::fmaf_r* Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RulesLang/APINamesMath.inc | 4 + .../RewriterSinglePrecisionIntrinsics.cpp | 80 +++++++++++++++++++ clang/lib/DPCT/SrcAPI/APINames.inc | 8 +- clang/test/dpct/math/cuda-math-extension.cu | 8 ++ clang/test/dpct/math/cuda-math-intrinsics.cu | 42 ++++++++++ 5 files changed, 138 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/APINamesMath.inc b/clang/lib/DPCT/RulesLang/APINamesMath.inc index 052947aab571..d411327b5d62 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..af6681824f33 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp @@ -170,6 +170,86 @@ 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, + CALL_FACTORY_ENTRY( + "__fmaf_ieee_rd", + CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd", + ARG(0), ARG(1), ARG(2)))), + 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, + CALL_FACTORY_ENTRY( + "__fmaf_ieee_rn", + CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn", + ARG(0), ARG(1), ARG(2)))), + 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, + CALL_FACTORY_ENTRY( + "__fmaf_ieee_ru", + CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru", + ARG(0), ARG(1), ARG(2)))), + 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, + CALL_FACTORY_ENTRY( + "__fmaf_ieee_rz", + CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz", + ARG(0), ARG(1), ARG(2)))), + 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 307cfdaa45e3..45a7193cc0d0 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/math/cuda-math-extension.cu b/clang/test/dpct/math/cuda-math-extension.cu index 24b2cec16b1c..84c114f5c257 100644 --- a/clang/test/dpct/math/cuda-math-extension.cu +++ b/clang/test/dpct/math/cuda-math-extension.cu @@ -181,6 +181,14 @@ __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: f2 = sycl::ext::intel::math::fmaf_rd(f0, f1, f2); + f2 = __fmaf_ieee_rd(f0, f1, f2); + // CHECK: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2); + f2 = __fmaf_ieee_rn(f0, f1, f2); + // CHECK: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2); + f2 = __fmaf_ieee_ru(f0, f1, f2); + // CHECK: 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 2de6f2d8fd7b..5a3bc17a7347 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: */ From b724c368f786f77edbf8b8425ba8b94678d0bea6 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 28 Apr 2025 11:37:21 +0800 Subject: [PATCH 2/6] Add warning Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/Diagnostics/Diagnostics.inc | 2 ++ .../RewriterSinglePrecisionIntrinsics.cpp | 36 ++++++++++++------- clang/test/dpct/math/cuda-math-extension.cu | 19 +++++++--- 3 files changed, 41 insertions(+), 16 deletions(-) diff --git a/clang/lib/DPCT/Diagnostics/Diagnostics.inc b/clang/lib/DPCT/Diagnostics/Diagnostics.inc index 9e132e474aa7..b820d50b4acb 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, 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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code.") +DEF_COMMENT(FTZ_BEHAVIOR, 1138, MEDIUM_LEVEL, "The flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code.") // clang-format on diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp index af6681824f33..18004ec6e4b1 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp @@ -175,10 +175,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { "__fmaf_ieee_rd", MATH_API_REWRITER_PAIR( math::Tag::math_libdevice, - CALL_FACTORY_ENTRY( + WARNING_FACTORY_ENTRY( "__fmaf_ieee_rd", - CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd", - ARG(0), ARG(1), ARG(2)))), + CALL_FACTORY_ENTRY("__fmaf_ieee_rd", + CALL(MapNames::getClNamespace() + + "ext::intel::math::fmaf_rd", + ARG(0), ARG(1), ARG(2))), + Diagnostics::FTZ_BEHAVIOR)), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -195,10 +198,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { "__fmaf_ieee_rn", MATH_API_REWRITER_PAIR( math::Tag::math_libdevice, - CALL_FACTORY_ENTRY( + WARNING_FACTORY_ENTRY( "__fmaf_ieee_rn", - CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn", - ARG(0), ARG(1), ARG(2)))), + CALL_FACTORY_ENTRY("__fmaf_ieee_rn", + CALL(MapNames::getClNamespace() + + "ext::intel::math::fmaf_rn", + ARG(0), ARG(1), ARG(2))), + Diagnostics::FTZ_BEHAVIOR)), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -215,10 +221,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { "__fmaf_ieee_ru", MATH_API_REWRITER_PAIR( math::Tag::math_libdevice, - CALL_FACTORY_ENTRY( + WARNING_FACTORY_ENTRY( "__fmaf_ieee_ru", - CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru", - ARG(0), ARG(1), ARG(2)))), + CALL_FACTORY_ENTRY("__fmaf_ieee_ru", + CALL(MapNames::getClNamespace() + + "ext::intel::math::fmaf_ru", + ARG(0), ARG(1), ARG(2))), + Diagnostics::FTZ_BEHAVIOR)), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -235,10 +244,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { "__fmaf_ieee_rz", MATH_API_REWRITER_PAIR( math::Tag::math_libdevice, - CALL_FACTORY_ENTRY( + WARNING_FACTORY_ENTRY( "__fmaf_ieee_rz", - CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz", - ARG(0), ARG(1), ARG(2)))), + CALL_FACTORY_ENTRY("__fmaf_ieee_rz", + CALL(MapNames::getClNamespace() + + "ext::intel::math::fmaf_rz", + ARG(0), ARG(1), ARG(2))), + Diagnostics::FTZ_BEHAVIOR)), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( diff --git a/clang/test/dpct/math/cuda-math-extension.cu b/clang/test/dpct/math/cuda-math-extension.cu index 84c114f5c257..1e0cf459f362 100644 --- a/clang/test/dpct/math/cuda-math-extension.cu +++ b/clang/test/dpct/math/cuda-math-extension.cu @@ -181,13 +181,24 @@ __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: f2 = sycl::ext::intel::math::fmaf_rd(f0, f1, f2); + // CHECK: DPCT1138:{{[0-9]+}}: The flush-to-zero behavior can only be controlled by the compiler option in SYCL. 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: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2); + // CHECK: /* + // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The flush-to-zero behavior can only be controlled by the compiler option in SYCL. 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: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2); + // CHECK: /* + // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The flush-to-zero behavior can only be controlled by the compiler option in SYCL. 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: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2); + // CHECK: /* + // CHECK-NEXT: DPCT1138:{{[0-9]+}}: The flush-to-zero behavior can only be controlled by the compiler option in SYCL. 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); From 993c9c796784797decf644cb9ce1102b8de1d9d8 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 28 Apr 2025 13:45:22 +0800 Subject: [PATCH 3/6] Fix Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/help_option_check/lin/help_option_check.cpp | 6 +++--- clang/test/dpct/help_option_check/win/help_option_check.cpp | 6 +++--- 2 files changed, 6 insertions(+), 6 deletions(-) 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_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 From 312d9ceda8e341f4b52e11c7970ed959ce3af9cc Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 28 Apr 2025 14:45:04 +0800 Subject: [PATCH 4/6] Fix lit Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/help_option_check/lin/help_all.txt | 2 +- clang/test/dpct/help_option_check/win/help_all.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) 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 bd46b8a23996..ae7d4818711c 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/win/help_all.txt b/clang/test/dpct/help_option_check/win/help_all.txt index bf32eaf4b3f2..a17102e21694 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). From 25e45fc83bc712c2de7d6b57065b5353ca44d2b6 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 30 Apr 2025 07:59:06 +0800 Subject: [PATCH 5/6] update warning Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/Diagnostics/Diagnostics.inc | 4 ++-- .../RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp | 8 ++++---- clang/test/dpct/math/cuda-math-extension.cu | 8 ++++---- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/lib/DPCT/Diagnostics/Diagnostics.inc b/clang/lib/DPCT/Diagnostics/Diagnostics.inc index b820d50b4acb..dd936ace57d0 100644 --- a/clang/lib/DPCT/Diagnostics/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics/Diagnostics.inc @@ -300,8 +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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code.") -DEF_COMMENT(FTZ_BEHAVIOR, 1138, MEDIUM_LEVEL, "The flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code.") +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/Math/RewriterSinglePrecisionIntrinsics.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp index 18004ec6e4b1..adf609188415 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp @@ -181,7 +181,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR)), + Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rd")), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -204,7 +204,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR)), + Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rn")), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -227,7 +227,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR)), + Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_ru")), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -250,7 +250,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR)), + Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rz")), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( diff --git a/clang/test/dpct/math/cuda-math-extension.cu b/clang/test/dpct/math/cuda-math-extension.cu index 1e0cf459f362..80f03703bcd3 100644 --- a/clang/test/dpct/math/cuda-math-extension.cu +++ b/clang/test/dpct/math/cuda-math-extension.cu @@ -181,22 +181,22 @@ __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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code. + // 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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code. + // 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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code. + // 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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code. + // 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); From 65e971666b016f3dfd346245ad5c3bfc3aceeb73 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 30 Apr 2025 09:50:45 +0800 Subject: [PATCH 6/6] Fix Signed-off-by: Jiang, Zhiwei --- .../RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp index adf609188415..cb3a928fa6d6 100644 --- a/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp +++ b/clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp @@ -181,7 +181,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rd")), + Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rd"))), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -204,7 +204,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rn")), + Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rn"))), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -227,7 +227,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_ru")), + Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_ru"))), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY( @@ -250,7 +250,7 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() { CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz", ARG(0), ARG(1), ARG(2))), - Diagnostics::FTZ_BEHAVIOR, "__fmaf_ieee_rz")), + Diagnostics::FTZ_BEHAVIOR, ARG("__fmaf_ieee_rz"))), MATH_API_REWRITER_PAIR( math::Tag::emulation, WARNING_FACTORY_ENTRY(