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() {
2    extern __shared__ float3 shared_memory[];
3    ...
4  }
5
6  void foo() {
7    size_t shared_size = 3 * sizeof(float);
8    kernel<<<1, 1, shared_size>>>();
9  }

results in the following migrated SYCL* code:

 1  void kernel(uint8_t *dpct_local) {
 2    auto shared_memory = (sycl::float3 *)dpct_local;
 3    ...
 4  }
 5
 6  void foo() {
 7    /*
 8    DPCT1083:0: The size of local memory in the migrated code may be different
 9    from the original code. Check that the allocated memory size in the migrated
10    code is correct.
11    */
12    size_t shared_size = 3 * sizeof(float);
13    dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
14      sycl::local_accessor<uint8_t, 1> dpct_local_acc_ct1(
15          sycl::range<1>(shared_size), 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            kernel(dpct_local_acc_ct1.get_pointer());
21          });
22    });
23  }

which is manually adjusted to:

 1  void kernel(uint8_t *dpct_local) {
 2    auto shared_memory = (sycl::float3 *)dpct_local;
 3    ...
 4  }
 5
 6  void foo() {
 7    size_t shared_size = 1 * sizeof(sycl::float3);
 8    dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
 9      sycl::local_accessor<uint8_t, 1> dpct_local_acc_ct1(
10          sycl::range<1>(shared_size), cgh);
11
12      cgh.parallel_for(
13          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
14          [=](sycl::nd_item<3> item_ct1) {
15            kernel(dpct_local_acc_ct1.get_pointer());
16          });
17    });
18  }

Suggestions to Fix#

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