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: カーネル関数ポインターはデバイスコードでは使用できません。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 }