Skip to content

Commit b724c36

Browse files
zhiweij1zhiweij1
committed
Add warning
Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>
1 parent ad1f6c2 commit b724c36

3 files changed

Lines changed: 41 additions & 16 deletions

File tree

‎clang/lib/DPCT/Diagnostics/Diagnostics.inc‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -300,6 +300,8 @@ DEF_WARNING(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Imag
300300
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()).")
301301
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.")
302302
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.")
303+
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.")
304+
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.")
303305

304306
// clang-format on
305307

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

Lines changed: 24 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -175,10 +175,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
175175
"__fmaf_ieee_rd",
176176
MATH_API_REWRITER_PAIR(
177177
math::Tag::math_libdevice,
178-
CALL_FACTORY_ENTRY(
178+
WARNING_FACTORY_ENTRY(
179179
"__fmaf_ieee_rd",
180-
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rd",
181-
ARG(0), ARG(1), ARG(2)))),
180+
CALL_FACTORY_ENTRY("__fmaf_ieee_rd",
181+
CALL(MapNames::getClNamespace() +
182+
"ext::intel::math::fmaf_rd",
183+
ARG(0), ARG(1), ARG(2))),
184+
Diagnostics::FTZ_BEHAVIOR)),
182185
MATH_API_REWRITER_PAIR(
183186
math::Tag::emulation,
184187
WARNING_FACTORY_ENTRY(
@@ -195,10 +198,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
195198
"__fmaf_ieee_rn",
196199
MATH_API_REWRITER_PAIR(
197200
math::Tag::math_libdevice,
198-
CALL_FACTORY_ENTRY(
201+
WARNING_FACTORY_ENTRY(
199202
"__fmaf_ieee_rn",
200-
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rn",
201-
ARG(0), ARG(1), ARG(2)))),
203+
CALL_FACTORY_ENTRY("__fmaf_ieee_rn",
204+
CALL(MapNames::getClNamespace() +
205+
"ext::intel::math::fmaf_rn",
206+
ARG(0), ARG(1), ARG(2))),
207+
Diagnostics::FTZ_BEHAVIOR)),
202208
MATH_API_REWRITER_PAIR(
203209
math::Tag::emulation,
204210
WARNING_FACTORY_ENTRY(
@@ -215,10 +221,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
215221
"__fmaf_ieee_ru",
216222
MATH_API_REWRITER_PAIR(
217223
math::Tag::math_libdevice,
218-
CALL_FACTORY_ENTRY(
224+
WARNING_FACTORY_ENTRY(
219225
"__fmaf_ieee_ru",
220-
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_ru",
221-
ARG(0), ARG(1), ARG(2)))),
226+
CALL_FACTORY_ENTRY("__fmaf_ieee_ru",
227+
CALL(MapNames::getClNamespace() +
228+
"ext::intel::math::fmaf_ru",
229+
ARG(0), ARG(1), ARG(2))),
230+
Diagnostics::FTZ_BEHAVIOR)),
222231
MATH_API_REWRITER_PAIR(
223232
math::Tag::emulation,
224233
WARNING_FACTORY_ENTRY(
@@ -235,10 +244,13 @@ RewriterMap dpct::createSinglePrecisionIntrinsicsRewriterMap() {
235244
"__fmaf_ieee_rz",
236245
MATH_API_REWRITER_PAIR(
237246
math::Tag::math_libdevice,
238-
CALL_FACTORY_ENTRY(
247+
WARNING_FACTORY_ENTRY(
239248
"__fmaf_ieee_rz",
240-
CALL(MapNames::getClNamespace() + "ext::intel::math::fmaf_rz",
241-
ARG(0), ARG(1), ARG(2)))),
249+
CALL_FACTORY_ENTRY("__fmaf_ieee_rz",
250+
CALL(MapNames::getClNamespace() +
251+
"ext::intel::math::fmaf_rz",
252+
ARG(0), ARG(1), ARG(2))),
253+
Diagnostics::FTZ_BEHAVIOR)),
242254
MATH_API_REWRITER_PAIR(
243255
math::Tag::emulation,
244256
WARNING_FACTORY_ENTRY(

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

Lines changed: 15 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -181,13 +181,24 @@ __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);
184+
// 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.
185+
// CHECK-NEXT: */
186+
// CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rd(f0, f1, f2);
185187
f2 = __fmaf_ieee_rd(f0, f1, f2);
186-
// CHECK: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2);
188+
// CHECK: /*
189+
// 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.
190+
// CHECK-NEXT: */
191+
// CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rn(f0, f1, f2);
187192
f2 = __fmaf_ieee_rn(f0, f1, f2);
188-
// CHECK: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2);
193+
// CHECK: /*
194+
// 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.
195+
// CHECK-NEXT: */
196+
// CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_ru(f0, f1, f2);
189197
f2 = __fmaf_ieee_ru(f0, f1, f2);
190-
// CHECK: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
198+
// CHECK: /*
199+
// 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.
200+
// CHECK-NEXT: */
201+
// CHECK-NEXT: f2 = sycl::ext::intel::math::fmaf_rz(f0, f1, f2);
191202
f2 = __fmaf_ieee_rz(f0, f1, f2);
192203
// CHECK: f2 = sycl::ext::intel::math::fmul_rd(f0, f1);
193204
f2 = __fmul_rd(f0, f1);

0 commit comments

Comments
 (0)