DPCT1123#
メッセージ#
カーネル関数ポインターはデバイスコードでは使用できません。正しい引数を指定してカーネル関数を直接呼び出す必要があります。カーネル関数の定義によって、sycl::nd_item の次元を調整することも必要になるかもしれません。
詳細な説明#
SYCL* 2020 ではデバイス関数内の関数ポインターの呼び出しをサポートしていないため (SYCL* 2020 仕様の 5.4. デバイス関数の言語制限)、ユーザーはその関数ポインターが指す関数を直接呼び出すようにコードを調整する必要があります。さらに、ツールは sycl::nd_item の次元 (移行中にオプション –assume-nd-range-dim=1 が使用された場合) と引数を適切に解析できないため、ユーザーは関連するコードを調整する必要がある場合もあります。
修正方法の提案#
例えば、以下のオリジナル CUDA* コードについて考えてみます。
1 __global__ void kernel(int *d) {
2 int gtid = blockIdx.x * blockDim.x + threadIdx.x;
3 d[gtid] = gtid;
4 }
5
6 void foo(int *d) {
7 void *kernel_array[100];
8 kernel_array[10] = (void *)&kernel;
9 void *args[1] = {&d};
10 cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0);
11 }このコードは、以下の SYCL* コードに移行されます。
1 void kernel(int *d, const sycl::nd_item<3> &item_ct1) {
2 int gtid = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
3 item_ct1.get_local_id(2);
4 d[gtid] = gtid;
5 }
6
7 void foo(int *d) {
8 sycl::device dev_ct1;
9 sycl::queue q_ct1(dev_ct1,
10 sycl::property_list{sycl::property::queue::in_order()});
11 void *kernel_array[100];
12 kernel_array[10] = (void *)&kernel;
13 void *args[1] = {&d};
14 /*
15 DPCT1123:0: The kernel function pointer cannot be used in the device code. You
16 need to call the kernel function with the correct argument(s) directly.
17 According to the kernel function definition, adjusting the dimension of the
18 sycl::nd_item may also be required.
19 */
20 q_ct1.parallel_for(
21 sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16),
22 sycl::range<3>(1, 1, 16)),
23 [=](sycl::nd_item<3> item_ct1) {
24 (kernel_array[10])();
25 });
26 }上記は次のように書き換える必要があります。
1 void kernel(int *d, const sycl::nd_item<3> &item_ct1) {
2 int gtid = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
3 item_ct1.get_local_id(2);
4 d[gtid] = gtid;
5 }
6
7 void foo(int *d) {
8 sycl::device dev_ct1;
9 sycl::queue q_ct1(dev_ct1,
10 sycl::property_list{sycl::property::queue::in_order()});
11 q_ct1.parallel_for(
12 sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16),
13 sycl::range<3>(1, 1, 16)),
14 [=](sycl::nd_item<3> item_ct1) {
15 kernel(d, item_ct1);
16 });
17 }