DPCT1123

Intel® DPC++ Compatibility Tool Developer Guide and Reference

Download PDF

ID 768918

Date 6/24/2024

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

Visible to Intel only — GUID: GUID-E8CD348E-D107-4B0B-893B-EDEB4DE42FD0

View Details

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:

__global__ void kernel(int *d) {
  int gtid = blockIdx.x * blockDim.x + threadIdx.x;
  d[gtid] = gtid;
}

void foo(int *d) {
  void *kernel_array[100];
  kernel_array[10] = (void *)&kernel;
  void *args[1] = {&d};
  cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0);
}

results in the following migrated SYCL code:

void kernel(int *d, const sycl::nd_item<3> &item_ct1) {
  int gtid = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
             item_ct1.get_local_id(2);
  d[gtid] = gtid;
}

void foo(int *d) {
  sycl::device dev_ct1;
  sycl::queue q_ct1(dev_ct1,
                    sycl::property_list{sycl::property::queue::in_order()});
  void *kernel_array[100];
  kernel_array[10] = (void *)&kernel;
  void *args[1] = {&d};
  /*
  DPCT1123:0: 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.
  */
  q_ct1.parallel_for(
      sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16),
                        sycl::range<3>(1, 1, 16)),
      [=](sycl::nd_item<3> item_ct1) {
        (kernel_array[10])();
      });
}

which needs to be rewritten to:

void kernel(int *d, const sycl::nd_item<3> &item_ct1) {
  int gtid = item_ct1.get_group(2) * item_ct1.get_local_range(2) +
             item_ct1.get_local_id(2);
  d[gtid] = gtid;
}

void foo(int *d) {
  sycl::device dev_ct1;
  sycl::queue q_ct1(dev_ct1,
                    sycl::property_list{sycl::property::queue::in_order()});
  q_ct1.parallel_for(
      sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16),
                        sycl::range<3>(1, 1, 16)),
      [=](sycl::nd_item<3> item_ct1) {
        kernel(d, item_ct1);
      });
}

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® DPC++ Compatibility Tool Developer Guide and Reference

DPCT1123

Message

Detailed Help

Suggestions to Fix