DPCT1054

メッセージ

<変数名> の型は、デバイスの関数内で <型> という名前で宣言されています。アクセサー宣言時に <型> 宣言が見えるようにコードを調整します。

説明

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

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

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

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

  • 呼び出し側では、ローカルアクセサーを定義する際に、最初のテンプレート引数が uint8_t[sizeof(<オリジナルの型 — 例: 「type_ct1」)] に設定されます。

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

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
 // original code:

 // header file
 template <typename T> __global__ void k() {
   __shared__ union {
     T t;
     ...
   } a;
   ...
 }

 // source file
 void foo() { k<int><<<1, 1>>>(); }

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

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
 // migrated DPC++ code:

 // header file
 template <typename T> void k(uint8_t *a_ct1) {
   union type_ct1 {
     T t;
     ...
   };
   type_ct1 *a = (type_ct1 *)a_ct1;
   ...
 }

 // source file
 void foo() {
   dpct::get_default_queue().submit([&](sycl::handler &cgh) {
     /*
     DPCT1053:0: The type of variable a is declared in device function with the
     name type_ct1. Adjust the code to make the type_ct1 declaration visible at
     the accessor declaration point.
     */
     sycl::accessor<uint8_t[sizeof(type_ct1)], 0, sycl::access::mode::read_write,
                    sycl::access::target::local> a_ct1_acc_ct1(cgh);
     ...
   });
 }

修正方法の提案

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