DPCT1038

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 10/31/2024

Version

Public

Visible to Intel only — GUID: GUID-74ADF0C6-C67D-49D5-9AB3-B8A05575A2A4

View Details

DPCT1038

Message

When the kernel function name is used as a macro argument, the migration result may be incorrect. You need to verify the definition of the macro.

Detailed Help

When a kernel function call is used in a macro definition and the function name is passed as a macro argument, the tool cannot determine the type of function parameters. This may lead to incorrectly generated DPC++ code.

Suggestions to Fix

Review the kernel call inside the macro definition and adjust it manually, if needed.

For example, this original CUDA* code:

 __global__ void kernel(int c, int d) {}

 void foo() {
 #define KERNEL_LAUNCH(NAME, a, b, c, d) NAME<<<a, b, 0>>>(c, d);
   KERNEL_LAUNCH(kernel, 3, 2, 1, 0)
 }

results in the following migrated SYCL* code:

 void kernel(int c, int d) {}

 void foo() {
 #define KERNEL_LAUNCH(NAME, a, b, c, d)                                        \
   dpct::get_default_queue().submit([&](sycl::handler &cgh) {                   \
     auto c_ct0 = c;                                                            \
     auto d_ct1 = d;                                                            \
                                                                                \
     cgh.parallel_for(                                                          \
         sycl::nd_range<3>(sycl::range<3>(1, 1, a) * sycl::range<3>(1, 1, b),   \
                           sycl::range<3>(1, 1, b)),                            \
         [=](sycl::nd_item<3> item_ct1) { kernel(c_ct0, d_ct1); });             \
   });
   /*
   DPCT1038:0: When the kernel function name is used as a macro argument, the
   migration result may be incorrect. You need to verify the definition of the
   macro.
   */
   KERNEL_LAUNCH(kernel, 3, 2, 1, 0)
 }

which is rewritten to:

 void kernel(int c, int d) {}

 void foo() {
 #define KERNEL_LAUNCH(NAME, a, b, c, d)                                        \
   dpct::get_default_queue().submit([&](sycl::handler &cgh) {                   \
     auto c_ct0 = c;                                                            \
     auto d_ct1 = d;                                                            \
                                                                                \
     cgh.parallel_for(                                                          \
         sycl::nd_range<3>(sycl::range<3>(1, 1, a) * sycl::range<3>(1, 1, b),   \
                           sycl::range<3>(1, 1, b)),                            \
         [=](sycl::nd_item<3> item_ct1) { NAME(c_ct0, d_ct1); });               \
   });
   KERNEL_LAUNCH(kernel, 3, 2, 1, 0)
 }

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1038

Message

Detailed Help

Suggestions to Fix