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

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


1// header file 
2template <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 
12void 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 
2template <typename T> 
3union type_ct1 { 
4    T t; 
5       ... 
6}; 
7template <typename T> void k(uint8_t *a_ct1) { 
8    type_ct1<T> &a = *(type_ct1<T> *)a_ct1; 
9       ... 
10} 
11 
12// source file 
13void 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> をオリジナルのタイプに必要なサイズ (バイト) に置き換えます。