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:

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

results in the following migrated SYCL* code:

 1 void kernel(int c, int d) {}
 2
 3 void foo() {
 4 #define KERNEL_LAUNCH(NAME, a, b, c, d)                                        \
 5   dpct::get_default_queue().submit([&](sycl::handler &cgh) {                   \
 6     auto c_ct0 = c;                                                            \
 7     auto d_ct1 = d;                                                            \
 8                                                                                \
 9     cgh.parallel_for(                                                          \
10         sycl::nd_range<3>(sycl::range<3>(1, 1, a) * sycl::range<3>(1, 1, b),   \
11                           sycl::range<3>(1, 1, b)),                            \
12         [=](sycl::nd_item<3> item_ct1) { kernel(c_ct0, d_ct1); });             \
13   });
14   /*
15   DPCT1038:0: When the kernel function name is used as a macro argument, the
16   migration result may be incorrect. You need to verify the definition of the
17   macro.
18   */
19   KERNEL_LAUNCH(kernel, 3, 2, 1, 0)
20 }

which is rewritten to:

 1 void kernel(int c, int d) {}
 2
 3 void foo() {
 4 #define KERNEL_LAUNCH(NAME, a, b, c, d)                                        \
 5   dpct::get_default_queue().submit([&](sycl::handler &cgh) {                   \
 6     auto c_ct0 = c;                                                            \
 7     auto d_ct1 = d;                                                            \
 8                                                                                \
 9     cgh.parallel_for(                                                          \
10         sycl::nd_range<3>(sycl::range<3>(1, 1, a) * sycl::range<3>(1, 1, b),   \
11                           sycl::range<3>(1, 1, b)),                            \
12         [=](sycl::nd_item<3> item_ct1) { NAME(c_ct0, d_ct1); });               \
13   });
14   KERNEL_LAUNCH(kernel, 3, 2, 1, 0)
15 }