[SYCLomatic] Migrate __fmaf_ieee_r* to sycl::ext::intel::math::fmaf_r* #2808
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
| @@ -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 flush-to-zero behavior can only be controlled by the compiler option in SYCL. Please verify the correctness of the migrated code. | |||
There was a problem hiding this comment.
The API __fmaf_ieee_rd() ignores the flush to zero compiler option when handling denormalized data, while the migrated SYCL API is impacted by it. Please...
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
There was a problem hiding this comment.
Pull Request Overview
This PR adds support for migrating the four IEEE rounding variants of __fmaf_ieee_r* intrinsics to the corresponding sycl::ext::intel::math::fmaf_r* calls and cleans up test invocation scripts.
- Introduces mappings for
__fmaf_ieee_rd,__fmaf_ieee_rn,__fmaf_ieee_ru, and__fmaf_ieee_rzin the single-precision intrinsic rewriter. - Removes redundant output redirection (
>> %T/diff.txt) from both Windows and Linux help-option tests.
Reviewed Changes
Copilot reviewed 3 out of 10 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| clang/test/dpct/help_option_check/win/help_option_check.cpp | Simplified RUN lines by dropping >> %T/diff.txt suffix |
| clang/test/dpct/help_option_check/lin/help_option_check.cpp | Simplified RUN lines by dropping >> %T/diff.txt suffix |
| clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp | Added mappings for __fmaf_ieee_r* intrinsics to new SYCL API |
Files not reviewed (7)
- clang/lib/DPCT/Diagnostics/Diagnostics.inc: Language not supported
- clang/lib/DPCT/RulesLang/APINamesMath.inc: Language not supported
- clang/lib/DPCT/SrcAPI/APINames.inc: Language not supported
- clang/test/dpct/help_option_check/lin/help_all.txt: Language not supported
- clang/test/dpct/help_option_check/win/help_all.txt: Language not supported
- clang/test/dpct/math/cuda-math-extension.cu: Language not supported
- clang/test/dpct/math/cuda-math-intrinsics.cu: Language not supported
Comments suppressed due to low confidence (1)
clang/lib/DPCT/RulesLang/Math/RewriterSinglePrecisionIntrinsics.cpp:173
- Consider adding unit or integration tests for the new
__fmaf_ieee_r*mappings to verify that each rounding mode (rd,rn,ru,rz) is correctly rewritten and that associated warnings (FTZ or unsupported rounding) behave as expected.
__fmaf_ieee_rd
Comment on lines
+174
to
+195
| 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))) |
There was a problem hiding this comment.
[nitpick] The four nearly identical blocks for __fmaf_ieee_r* could be refactored into a shared helper or macro to reduce duplication and simplify future additions or adjustments.
Suggested change
| 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))) | |
| DEFINE_FMAF_IEEE_REWRITER("__fmaf_ieee_rd", "fmaf_rd", Diagnostics::FTZ_BEHAVIOR) |
Copilot uses AI. Check for mistakes.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
No description provided.