DPCT1120#

Message#

<variable name> is in device memory and passed into device functions as an extra argument. As <variable name> is a template variable here, then the declaration of an extra parameter in device functions, the variable referenced in command group, and device functions may be incorrect as a result. You may need to adjust the code.

Detailed Help#

<variable name> is in device memory and passed into device functions as an extra argument. As <variable name> is a template variable here, then the declaration of an extra parameter in device functions, the variable referenced in command group, and device functions should be template variable also, the migration result may be incorrect. You may need to adjust the code.

Suggestions to Fix#

For example, this original CUDA* code:

 1struct Foo {
 2  int x;
 3  __host__ __device__ Foo() {}
 4  __host__ __device__ ~Foo() {}
 5};
 6
 7template <typename T> __constant__ T cmem;
 8
 9__global__ void kernel() {
10  ...
11  int x = cmem<Foo>.x;
12  ...
13}
14
15void foo() {
16  ...
17  kernel<<<1, 1>>>();
18  ...
19}

results in the following migrated SYCL* code:

 1struct Foo {
 2  int x;
 3  Foo() {}
 4  ~Foo() {}
 5};
 6
 7/*
 8DPCT1120:0: cmem is in device memory and passed into device functions as an
 9extra arguments. As cmem is a template variable here, then the declaration of
10an extra parameter in device functions, the variable referenced in command
11group, and device functions may be incorrect as a result. You may need to
12adjust the code.
13*/
14template <typename T> static dpct::constant_memory<T, 0> cmem;
15
16void kernel(T cmem) {
17  ...
18  int x = cmem<Foo>.x;
19  ...
20}
21
22void foo() {
23  ...
24  cmem.init();
25
26  dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
27    auto cmem_ptr_ct1 = cmem.get_ptr();
28
29    cgh.parallel_for(
30        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
31        [=](sycl::nd_item<3> item_ct1) {
32          kernel(*cmem_ptr_ct1);
33        });
34  });
35  ...
36}

which needs to be rewritten to:

 1struct Foo {
 2  int x;
 3  Foo() {}
 4  ~Foo() {}
 5};
 6
 7template <typename T> static dpct::constant_memory<T, 0> cmem;
 8
 9template <class T>
10void kernel(T cmem) {
11  ...
12  int x = cmem.x;
13  ...
14}
15
16void foo() {
17  ...
18  cmem<Foo>.init();
19
20  dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
21    auto cmem_ptr_ct1 = cmem<Foo>.get_ptr();
22
23    cgh.parallel_for(
24        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
25        [=](sycl::nd_item<3> item_ct1) {
26          kernel(*cmem_ptr_ct1);
27        });
28  });
29  ...
30}