DPCT1111#

Message#

Please verify the input arguments of <migrated API> base on the target function <kernel function name>.

Detailed Help#

SYCLomatic cannot deduce some arguments from a function pointer or CUfunction variable. You need verify the argument value manually.

Suggestions to Fix#

For example, this original CUDA* code:

 1extern __shared__ float a[];
 2__global__ void kernel() {
 3  __shared__ int b[10];
 4  __syncthreads();
 5}
 6
 7void foo() {
 8  int numBlocks;
 9  cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, kernel, 128,
10                                                sizeof(float) * 20);
11  kernel<<<1, 128>>>();
12}

results in the following migrated SYCL* code:

 1void kernel(const sycl::nd_item<3> &item_ct1, int *b) {
 2
 3  item_ct1.barrier(sycl::access::fence_space::local_space);
 4}
 5
 6void foo() {
 7  int numBlocks;
 8  /*
 9  DPCT1111:0: Please verify the input arguments of
10  dpct::experimental::calculate_max_active_wg_per_xecore base on the target
11  function "kernel".
12  */
13  dpct::experimental::calculate_max_active_wg_per_xecore(
14      &numBlocks, 128,
15      sizeof(float) * 20 +
16          dpct_placeholder /* total shared local memory size */);
17  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
18    sycl::local_accessor<int, 1> b_acc_ct1(sycl::range<1>(10), cgh);
19
20    cgh.parallel_for(
21        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
22        [=](sycl::nd_item<3> item_ct1) {
23          kernel(item_ct1, b_acc_ct1.get_pointer());
24        });
25  });
26}

which is rewritten to:

 1void kernel(const sycl::nd_item<3> &item_ct1, int *b) {
 2  item_ct1.barrier(sycl::access::fence_space::local_space);
 3}
 4
 5void foo() {
 6  int numBlocks;
 7  dpct::experimental::calculate_max_active_wg_per_xecore(
 8      &numBlocks, 128, sizeof(float) * 20 + sizeof(int) * 10, 32, true, false);
 9  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
10    sycl::local_accessor<int, 1> b_acc_ct1(sycl::range<1>(10), cgh);
11
12    cgh.parallel_for(
13        sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
14        [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
15          kernel(item_ct1, b_acc_ct1.get_pointer());
16        });
17  });
18}