DPCT1038
目次
DPCT1038#
メッセージ#
マクロ引数にカーネル関数名を使用すると、移行の結果が正しくない場合があります。マクロの定義を確認する必要があります。
詳細な説明#
カーネル関数呼び出しがマクロ定義で使用され、関数名がマクロ引数として渡される場合、ツールは関数パラメーターのタイプを判断できません。これにより、DPC++ コードが正しく生成されない可能性があります。
修正方法の提案#
マクロ定義内のカーネル呼び出しを確認して、必要に応じて手動で調整してください。
例えば、以下のオリジナル CUDA* コードについて考えてみます。
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 }
このコードは、以下の SYCL* コードに移行されます。
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 }
このコードは次のように書き換えられます。
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 }