DPCT1083#

Message#

The size of <placeholder> in the migrated code may be different from the original code. Check that the allocated memory size in the migrated code is correct.

Detailed Help#

Some types have a different size in the migrated code than in the original code, for example sycl::float3 compared to float3. Check if the allocated size of memory is correct in the migrated code.

In the example below, 3*sizeof(float) is used to represent the size of float3 in the original code. In the migrated code the size of sycl::float3 is different, so the allocated size needs adjustment.

For example, this original CUDA* code:

1  _global_ void kernel() {   extern __shared__ float3 shared_memory[]; }
2
3  int main()
4  {   size_t shared_size = 3 * sizeof(float);   kernel<<<1, 1, shared_size>>>();  ... }

results in the following migrated SYCL* code:

 1  void kernel(uint8_t *dpct_local)
 2  {   auto shared_memory = (float3 *)dpct_local; }
 3  int main() {
 4    /*
 5    DPCT1083
 6    */
 7    size_t shared_size = 3 * sizeof(float);
 8    get_default_queue().submit([&](handler &cgh) {
 9      accessor<...> dpct_local_acc_ct1(range<1>(shared_size), cgh);
10      cgh.parallel_for(...,
11          [=](nd_item<3> item_ct1) {
12          kernel(dpct_local_acc_ct1.get_pointer());         });
13    });
14    ...
15  }

which is manually adjusted to:

 1  void kernel(uint8_t *dpct_local) {   auto shared_memory = (float3 *)dpct_local; }
 2
 3  int main() {
 4    size_t shared_size = 1 * sizeof(float3);
 5    get_default_queue().submit([&](handler &cgh)
 6      Unknown macro: {
 7        accessor<...> dpct_local_acc_ct1(range<1>(shared_size), cgh);    cgh.parallel_for(...,        [=](nd_item<3> item_ct1) {
 8          kernel(dpct_local_acc_ct1.get_pointer());
 9        });
10      }
11    );
12    ...
13  }

Suggestions to Fix#

Check the allocated size of memory and replace it with the correct size if necessary.