DPCT1083#

Message#

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

Or

The size of <type name> in the migrated code may be different from the original code. You may need to adjust the code.

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 memory size is correct in the migrated code.

Suggestions to Fix#

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  }
10
11  void bar(char *data_handle, int width, int height) {
12    std::vector<uchar3> data;
13    data.assign(reinterpret_cast<uchar3 *>(data_handle),
14               reinterpret_cast<uchar3 *>(data_handle) + width * height);
15    ...
16  }

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__multi_ptr<sycl::access::decorated::no>()
21                       .get());
22          });
23    });
24  }
25
26  void bar(char *data_handle, int width, int height) {
27    std::vector<sycl::uchar3> data;
28    /*
29    DPCT1083:1: The size of uchar3 in the migrated code may be different from the original code. You may need to adjust the code.
30    */
31    data.assign (reinterpret_cast<sycl::uchar3 *>(data_handle),
32                 reinterpret_cast<sycl::uchar3 *>(data_handle) + width * height);
33    ...
34  }

which needs to be rewritten 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_multi_ptr<sycl::access::decorated::no>()
16                .get());
17          });
18    });
19  }
20
21  void bar(char *data_handle, int width, int height) {
22    std::vector<sycl::uchar3> data;
23    for (int Index = 0; Index < width * height; Index++) {
24      data.emplace_back(data_handle[Index * 3], data_handle[Index * 3 = 1],
25                        data_handle[Index * 3+2]);
26    }
27    ...
28  }