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:

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

 // source file
 void foo() { k<int><<<1, 1>>>(); }
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
 // Migrated code:
 // header file
 template <typename T> void k(uint8_t *a_ct1) {
   union type_ct1 {
     T t;
     ...
   };
   type_ct1 *a = (type_ct1 *)a_ct1;
   ...
 }

 // source file
 void foo() {
   dpct::get_default_queue().submit([&](sycl::handler &cgh) {
     /*
     DPCT1053:0: The type of variable a is declared in device function with the
     name type_ct1. Adjust the code to make the type_ct1 declaration visible at
     the accessor declaration point.
     */
     sycl::accessor<uint8_t[sizeof(type_ct1)], 0, sycl::access::mode::read_write,
                    sycl::access::target::local> a_ct1_acc_ct1(cgh);
     ...
   });
 }

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.