DPCT1069#

Message#

The argument <argument name> of the kernel function contains virtual pointer(s), which cannot be dereferenced. Try to migrate the code with usm-level=restricted.

Detailed Help#

SYCLomatic cannot process inner virtual pointer(s) when using option --usm-level=none.

Suggestions to Fix#

Re-migrate the code without specifing --usm-level=none.

For example, this original CUDA* code:

 1struct AAA {
 2  int *a;
 3};
 4
 5__global__ void k(AAA obj) {
 6  obj.a[2] = 123;
 7}
 8
 9void foo() {
10  AAA obj;
11  int *a;
12  cudaMalloc(&a, sizeof(int) * 10);
13  obj.a = a;
14  k<<<1, 1>>>(obj);
15}

results in the following migrated SYCL* code with option --usm-level=none:

 1#define DPCT_USM_LEVEL_NONE
 2#include <sycl/sycl.hpp>
 3#include <dpct/dpct.hpp>
 4struct AAA {
 5  int *a;
 6};
 7
 8void k(AAA obj) {
 9  obj.a[2] = 123;
10}
11
12void foo() {
13  AAA obj;
14  int *a;
15  a = (int *)dpct::dpct_malloc(sizeof(int) * 10);
16  obj.a = a;
17  /*
18  DPCT1069:0: The argument 'obj' of the kernel function contains virtual
19  pointer(s), which cannot be dereferenced. Try to migrate the code with
20  "usm-level=restricted".
21  */
22  dpct::get_out_of_order_queue().parallel_for(
23      sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
24      [=](sycl::nd_item<3> item_ct1) {
25        k(obj);
26      });
27}

Re-migrated code without option --usm-level=none:

 1#include <sycl/sycl.hpp>
 2#include <dpct/dpct.hpp>
 3struct AAA {
 4  int *a;
 5};
 6
 7void k(AAA obj) {
 8  obj.a[2] = 123;
 9}
10
11void foo() {
12  sycl::device dev_ct1;
13  sycl::queue q_ct1(dev_ct1,
14                    sycl::property_list{sycl::property::queue::in_order()});
15  AAA obj;
16  int *a;
17  a = sycl::malloc_device<int>(10, q_ct1);
18  obj.a = a;
19  q_ct1.parallel_for(
20      sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
21      [=](sycl::nd_item<3> item_ct1) {
22        k(obj);
23      });
24}