DPCT1054#

メッセージ#

<variable name> のタイプは、デバイスの関数内で <type> という名前で宣言されています。アクセサー宣言時に <type> 宣言が参照できるようにコードを調整します。

詳細な説明#

この警告は、__shared__ 変数の型がデバイス関数内で宣言されている場合に出力されます。

移行したコードで次のことが行われます。

  • ツールは、デバイス関数に以下を追加します。

    # uint8_t* 入力パラメーター # タイプへの名前 (名前がない場合) — 例: type_ct1 # uint8_t* からオリジナルのタイプへのタイプキャスト (オリジナルの型宣言の後)

  • 呼び出し側では、ローカルアクセサーを定義する際に、最初のテンプレート引数が uint8_t[sizeof(<original type, for example "type_ct1">)] に設定されます。

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

1  // header file 
2  template <typename T> __global__ void k() { 
3   __shared__ union { 
4   T t; 
5      ...
6   } a; 
7      ... 
8  } 
9 
10 // source file 
11 void foo() { k<int><<<1, 1>>>(); }

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

1  // header file 
2  template <typename T> void k(uint8_t *a_ct1) { 
3   union type_ct1 { 
4   T t; 
5      ... 
6   }; 
7   type_ct1 &a = *(type_ct1 *)a_ct1; 
8    ... 
9   } 
10 
11 // source file 
12 void foo() { 
13  dpct::get_default_queue().submit([&](sycl::handler &cgh) { 
14  /* 
15  DPCT1054:0: The type of variable a is declared in device function with the 
16  name type_ct1.Adjust the code to make the type_ct1 declaration visible at 
17  the accessor declaration point.
18  */ 
19  sycl::local_accessor<uint8_t[sizeof(type_ct1)], 0> a_ct1_acc_ct1(cgh); 
20 
21  cgh.parallel_for( 
22  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
23     [=](sycl::nd_item<3> item_ct1) { 
24     k<int>(a_ct1_acc_ct1.get_pointer()); 
25     }); 
26  }); 
27 }

このコードは次のように書き換えられます。

1  // header file 
2  template <typename T> 
3  union type_ct1 { 
4   T t; 
5      ... 
6  }; 
7  template <typename T> void k(uint8_t *a_ct1) { 
8   type_ct1<T> &a = *(type_ct1<T> *)a_ct1; 
9      ... 
10 } 
11 
12 // source file 
13 void foo() { 
14  dpct::get_default_queue().submit([&](sycl::handler &cgh) { 
15  sycl::local_accessor<uint8_t[sizeof(type_ct1<T>)], 0> a_ct1_acc_ct1(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     k<int>(a_ct1_acc_ct1.get_pointer()); 
21     }); 
22  }); 
23 }

修正方法の提案#

型宣言をアクセサー宣言位置で見えるように移動するか sizeof(<original type> をオリジナルのタイプに必要なサイズ (バイト) に置き換えます。