DPCT1125#

Message#

The type “<type name>” defined in function “<function name>” is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the definition location of the type.

Detailed Help#

SYCL* 2020 does not support declaring a local memory variable in place. A local accessor needs to be declared in the command group scope and passed to all functions using the local memory. So, the definition location of the local memory variable type needs to be changed. Note: The DPC++ compiler extension sycl_ext_oneapi_local_memory (intel/llvm) supports declaring a local memory variable in place, but only in the kernel function scope.

Suggestions to Fix#

For example, this original CUDA* code:

 1template <typename T> struct kernel_type_t {
 2  using Type = T;
 3};
 4
 5template <typename T> __global__ void device(int a, int b) {
 6  using Tk = typename kernel_type_t<T>::Type;
 7  __shared__ Tk mem[256];
 8  ...
 9}
10
11template <typename T> __global__ void kernel(int a, int b) { device<T>(a, b); }
12
13template <typename T> void foo() {
14  int i;
15  kernel<T><<<1, 1>>>(i, i);
16}

results in the following migrated SYCL code:

 1template <typename T> struct kernel_type_t {
 2  using Type = T;
 3};
 4
 5/*
 6DPCT1125:1: The type "Tk" defined in function "device" is used as the parameter
 7type in all functions in the call path from the corresponding
 8sycl::handler::parallel_for() to the current function. You may need to adjust
 9the definition location of the type.
10*/
11template <typename T> void device(int a, int b, Tk *mem) {
12  using Tk = typename kernel_type_t<T>::Type;
13
14  ...
15}
16
17/*
18DPCT1125:2: The type "Tk" defined in function "device" is used as the parameter
19type in all functions in the call path from the corresponding
20sycl::handler::parallel_for() to the current function. You may need to adjust
21the definition location of the type.
22*/
23template <typename T>
24void kernel(int a, int b, Tk *mem) {
25  device<T>(a, b, mem);
26}
27
28template <typename T> void foo() {
29  int i;
30  /*
31  DPCT1126:0: The type "Tk" defined in function "device" is used as the
32  parameter type in all functions in the call path from the
33  sycl::handler::parallel_for() to the function "device". You may need to adjust
34  the definition location of the type.
35  */
36  dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
37    sycl::local_accessor<Tk, 1> mem_acc_ct1(sycl::range<1>(256), cgh);
38
39    cgh.parallel_for(
40        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
41        [=](sycl::nd_item<3> item_ct1) {
42          kernel<T>(
43              i, i,
44              mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>()
45                  .get());
46        });
47  });
48}

which needs to be rewritten to:

 1template <typename T> struct kernel_type_t {
 2  using Type = T;
 3};
 4
 5template <typename T> void device(int a, int b, typename kernel_type_t<T>::Type *mem) {
 6  ...
 7}
 8
 9template <typename T>
10void kernel(int a, int b, typename kernel_type_t<T>::Type *mem) {
11  device<T>(a, b, mem);
12}
13
14template <typename T> void foo() {
15  int i;
16  using Tk = typename kernel_type_t<T>::Type;
17  dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
18    sycl::local_accessor<Tk, 1> mem_acc_ct1(sycl::range<1>(256), cgh);
19
20    cgh.parallel_for(
21        sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
22        [=](sycl::nd_item<3> item_ct1) {
23          kernel<T>(
24              i, i,
25              mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
26        });
27  });
28}