DPCT1125#

メッセージ#

関数 “<関数名>” で定義されたタイプ “<タイプ名>” は、対応する sycl::handler::parallel_for() から現在の関数までの呼び出しパス内のすべての関数のパラメーター・タイプとして使用されます。タイプの定義場所を調整する必要があるかもしれません。

詳細な説明#

SYCL* 2020 は、ローカルメモリー変数のインプレース宣言をサポートしていません。ローカルアクセサーはコマンド・グループ・スコープで宣言し、ローカルメモリーを使用するすべての関数に渡す必要があります。そのため、ローカルメモリー変数他ぷの定義場所を変更する必要があります。注: DPC++ コンパイラー拡張 sycl_ext_oneapi_local_memory (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_local_memory.asciidoc) は、カーネル関数スコープ内でのみ、ローカルメモリー変数のインプレース宣言をサポートします。

修正方法の提案

例えば、以下のオリジナル CUDA* コードについて考えてみます。

1  template <typename T> struct kernel_type_t { 
2   using Type = T; 
3  }; 
4 
5  template <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 
11 template <typename T> __global__ void kernel(int a, int b) { device<T>(a, b); } 
12 
13 template <typename T> void foo() { 
14  int i; 
15  kernel<T><<<1, 1>>>(i, i); 
16 }

このコードは、以下の SYCL* コードに移行されます。

1  template <typename T> struct kernel_type_t { 
2   using Type = T; 
3  }; 
4 
5  /* 
6  DPCT1125:1: The type "Tk" defined in function "device" is used as the parameter 
7  type in all functions in the call path from the corresponding 
8  sycl::handler::parallel_for() to the current function.You may need to adjust 
9  the definition location of the type.
10 */ 
11 template <typename T> void device(int a, int b, Tk *mem) { 
12  using Tk = typename kernel_type_t<T>::Type; 
13 
14  ... 
15 } 
16 
17 /* 
18 DPCT1125:2: The type "Tk" defined in function "device" is used as the parameter 
19 type in all functions in the call path from the corresponding 
20 sycl::handler::parallel_for() to the current function.You may need to adjust 
21 the definition location of the type.
22 */ 
23 template <typename T> 
24 void kernel(int a, int b, Tk *mem) { 
25  device<T>(a, b, mem); 
26 } 
27 
28 template <typename T> void foo() { 
29  int i; 
30  /* 
31  DPCT1126:0: The type "Tk" defined in function "device" is used as the parameter 
32 type in all functions in the call path from the corresponding 
33 sycl::handler::parallel_for() to the current function.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 }

上記は次のように書き換える必要があります。

1  template <typename T> struct kernel_type_t { 
2   using Type = T; 
3  }; 
4 
5  template <typename T> void device(int a, int b, typename kernel_type_t<T>::Type *mem) { 
6   ... 
7  } 
8 
9  template <typename T> 
10 void kernel(int a, int b, typename kernel_type_t<T>::Type *mem) { 
11  device<T>(a, b, mem); 
12 } 
13 
14 template <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 }