DPCT1054#

Message#

The type of variable <variable name> is declared in device function with the name <type>. Adjust the code to make the <type> declaration visible at the accessor declaration point.

Detailed Help#

This warning will be emitted when the type of a __shared__ variable is declared in a device function.

In the migrated code:

  • in the device function, the tool adds:

    # uint8_t* input parameter # name to the type, if unnamed, for example type_ct1 # a type cast from uint8_t* to the original type, after the original type declaration.

  • in the call side, when defining local accessor, the first template argument will be set to uint8_t[sizeof(<original type, for example "type_ct1">)].

For example, this original CUDA* code:

 1// header file
 2template <typename T> __global__ void k() {
 3  __shared__ union {
 4    T t;
 5    ...
 6  } a;
 7  ...
 8}
 9
10// source file
11void foo() { k<int><<<1, 1>>>(); }

results in the following migrated SYCL* code:

 1// header file
 2template <typename T> void k(uint8_t *a_ct1) {
 3  union type_ct1 {
 4    T t;
 5    ...
 6  };
 7  type_ct1 &a = *(type_ct1 *)a_ct1;
 8  ...
 9}
10
11// source file
12void foo() {
13  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
14    /*
15    DPCT1054:0: The type of variable a is declared in device function with the
16    name type_ct1. Adjust the code to make the type_ct1 declaration visible at
17    the accessor declaration point.
18    */
19    sycl::local_accessor<uint8_t[sizeof(type_ct1)], 0> a_ct1_acc_ct1(cgh);
20
21    cgh.parallel_for(
22        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
23        [=](sycl::nd_item<3> item_ct1) {
24          k<int>(a_ct1_acc_ct1.get_pointer());
25        });
26  });
27}

which is rewritten to:

 1// header file
 2template <typename T>
 3union type_ct1 {
 4  T t;
 5  ...
 6};
 7template <typename T> void k(uint8_t *a_ct1) {
 8  type_ct1<T> &a = *(type_ct1<T> *)a_ct1;
 9  ...
10}
11
12// source file
13void foo() {
14  dpct::get_default_queue().submit([&](sycl::handler &cgh) {
15    sycl::local_accessor<uint8_t[sizeof(type_ct1<T>)], 0> a_ct1_acc_ct1(cgh);
16
17    cgh.parallel_for(
18        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
19        [=](sycl::nd_item<3> item_ct1) {
20          k<int>(a_ct1_acc_ct1.get_pointer());
21        });
22  });
23}

Suggestions to Fix#

Move the type declaration so that it will be visible at the accessor declaration point or replace the sizeof(<original type> with size in bytes needed for the original type.