Intel® DPC++ Compatibility Tool Developer Guide and Reference
A newer version of this document is available. Customers should click here to go to the newest version.
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:
// header file
template <typename T> __global__ void k() {
__shared__ union {
T t;
...
} a;
...
}
// source file
void foo() { k<int><<<1, 1>>>(); }results in the following migrated SYCL* 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) {
/*
DPCT1054: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::local_accessor<uint8_t[sizeof(type_ct1)], 0> a_ct1_acc_ct1(cgh);
cgh.parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
k<int>(a_ct1_acc_ct1.get_pointer());
});
});
}which is rewritten to:
// header file
template <typename T>
union type_ct1 {
T t;
...
};
template <typename T> void k(uint8_t *a_ct1) {
type_ct1<T> &a = *(type_ct1<T> *)a_ct1;
...
}
// source file
void foo() {
dpct::get_default_queue().submit([&](sycl::handler &cgh) {
sycl::local_accessor<uint8_t[sizeof(type_ct1<T>)], 0> a_ct1_acc_ct1(cgh);
cgh.parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
k<int>(a_ct1_acc_ct1.get_pointer());
});
});
}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.