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

results in the following migrated SYCL* code:

 1  // header file
 2  template <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
12  void foo() {
13    dpct::get_default_queue().submit([&](sycl::handler &cgh) {
14      /*
15      DPCT1053: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::accessor<uint8_t[sizeof(type_ct1)], 0, sycl::access::mode::read_write,
20                     sycl::access::target::local> a_ct1_acc_ct1(cgh);
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.