DPCT1123#

Message#

The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required

Detailed Help#

Since SYCL* 2020 does not support calling a function pointer in a device function (SYCL 2020 Spec, 5.4. Language restrictions for device functions), the user needs to adjust the code to directly call the function pointed by that function pointer. Furthermore, the tool cannot do the analysis well for the dimension of sycl::nd_item (if option –assume-nd-range-dim=1 is used during migration) and the arguments, so the user may also need to adjust the related code.

Suggestions to Fix#

For example, this original CUDA* code:

 1__global__ void kernel(int *d) {
 2  int gtid = blockIdx.x * blockDim.x + threadIdx.x;
 3  d[gtid] = gtid;
 4}
 5
 6void 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}

results in the following migrated SYCL code:

 1void 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
 7void 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}

which needs to be rewritten to:

 1void 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
 7void 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}