カスタム GPU 操作の実装方法¶
OpenVINO™ でサポートされていない操作をそのまま有効にするには、OpenVINO 操作セットの拡張機能と、対象となるデバイスのカスタムカーネルが必要になる場合があります。ここでは、GPU デバイスのカスタムカーネルのサポートについて説明します。
GPU コードパスは、OpenCL* に関する多くの詳細を抽象化します。OpenCL* C のカーネルコードと、そのパラメーターを操作のパラメーターに接続する XML 構成ファイルを用意する必要があります。
カスタム操作構成ファイルを使用するには、次の 2 つのオプションがあります。
カーネルを含むセクションを、自動的にロードされる
<lib_path>/cldnn_global_custom_kernels/cldnn_global_custom_kernels.xmlファイルに含めます。カスタム操作を使用するネットワークをプラグインにロードする前に、
"CONFIG_FILE"キーと構成ファイル名を値として使用し、アプリケーションからov::Core::set_property()メソッドを呼び出します。
core = ov.Core()
core.set_property("GPU", {"CONFIG_FILE": "<path_to_the_xml_file>"})
ov::Core core;
// Load GPU Extensions
core.set_property("GPU", {{ "CONFIG_FILE", "<path_to_the_xml_file>" }});
簡単な hello_classification を除くすべての OpenVINO サンプルと、ほとんどの Open Model Zoo のデモには、カスタムカーネルをロードする専用のコマンドライン・オプション -c があります。例えば、分類サンプルのカスタム操作を読み込むには、以下のコマンドを実行します。
$ ./classification_sample -m <path_to_model>/bvlc_alexnet_fp16.xml -i ./validation_set/daily/227x227/apron.bmp -d GPU
-c <absolute_path_to_config>/custom_layer_example.xml
設定ファイルの形式¶
構成ファイルは、提供するすべてのカスタム操作に対してタイプ CustomLayer のノードを含む .xml ファイル構造に従うことが期待されます。
以下のセクションで説明する定義では、次の表記を使用します。
表記規則 |
説明 |
|---|---|
(0/1) |
このノードまたは属性のインスタンスを 0 つまたは 1 つ持つことができます |
|
|
このノードまたは属性のインスタンスは 1 つだけ必要です |
(0+) |
このノードまたは属性のインスタンスはいくつでも持つことができます |
(1+) |
このノードまたは属性のインスタンスを 1 つ以上持つことができます |
CustomLayer ノードとサブノードの構造¶
CustomLayer ノードには、単一のカスタム操作の構成全体が含まれています。
属性名 |
# |
説明 |
|---|---|---|
|
|
使用する操作タイプの名前。この名前は、IR で使用されるタイプと同一である必要があります。 |
|
|
|
|
|
|
サブノード: Kernel (1)、Buffers (1), CompilerOptions (0+)、WorkSizes (0/1)
ソースノードとサブノードの構造¶
Source ノードは、単一の OpenCL* ソースファイルを指します。
属性名 |
# |
説明 |
|---|---|---|
|
|
OpenCL* ソースコードを含むファイルの名前。パスは実行可能ファイルからの相対パスです。複数のソースノードには、ソースが順番に連結されます。 |
サブノード: なし
ノードとサブノード構造の定義¶
Define ノードは、コンパイル (JIT) 中にソースに追加される単一の #define 命令を構成します。
属性名 |
# |
説明 |
|---|---|---|
|
|
定義された JIT の名前。静的定数の場合、これには文字列として受け取られる値も含めることができます。 |
|
(0/1) |
このパラメーター値は、この JIT 定義の値として使用されます。 |
|
(0/1) |
パラメーターのタイプです。受け入れる値: |
|
(0/1) |
指定されたパラメーターが OpenVINO IR の操作に含まれていない場合に使用されるデフォルト値。 |
サブノード: なし
結果の JIT は次の形式になります: #define [name] [type] [value/default]。
バッファーノードとサブノードの構造¶
Buffers ノードは、OpenCL* エントリー関数のすべての入出力バッファを構成します。バッファーノード構造は存在しません。
サブノード: Data (0+)、Tensor (1+)
データノードとサブノード構造¶
Data ノードは、重みやバイアスなど静的データを含む単一の入力を構成します。
属性名 |
# |
説明 |
|---|---|---|
|
|
OpenVINO IR の操作にアタッチされた BLOB の名前。 |
|
|
バインドされるエントリー関数引数内の 0 から始まるインデックス。 |
サブノード: なし
テンソルノードとサブノード構造¶
Tensor ノードは、単一の入力または出力テンソルを構成します。
属性名 |
# |
説明 |
|---|---|---|
|
|
バインドされるエントリー関数引数内の 0 から始まるインデックス。 |
|
|
|
|
|
OpenVINO IR の操作入出力ポートの 0 から始まるインデックス。 |
|
(0/1) |
テンソルのデータレイアウト宣言。受け入れる値: |
CompilerOptions ノードとサブノードの構造¶
CompilerOptions ノードは、OpenCL* ソースのコンパイルフラグを設定します。
属性名 |
# |
説明 |
|---|---|---|
|
|
OpenCL* コンパイラーに渡されるオプション文字列。 |
サブノード: なし
WorkSizes ノードとサブノードの構造¶
WorkSizes ノードは、OpenCL* プログラムを実行するためキューに投入する際に、グローバル/ローカル・ワーク・サイズを構成します。
属性名 |
# |
説明 |
|---|---|---|
|
(0/1) (0/1) |
実行中に使用される OpenCL* のワークサイズを定義する、最大 3 つの整数または式の配列。数式では B、F、Y、X 次元の値を使用でき、演算子 +、-、/、*、% を含めることができます。すべての演算子は整数算術で評価されます。デフォルト値: |
|
(0/1) |
ワークサイズを取得するテンソル。受け入れる値: |
サブノード: なし
設定ファイルの例¶
次のサンプルコードは、XML 形式の構成ファイルの例を示しています。設定ファイルの構造については、設定ファイルの形式を参照してください。
<CustomLayer name="ReLU" type="SimpleGPU" version="1">
<Kernel entry="example_relu_kernel">
<Source filename="custom_layer_kernel.cl"/>
<Define name="neg_slope" type="float" param="negative_slope" default="0.0"/>
</Kernel>
<Buffers>
<Tensor arg-index="0" type="input" port-index="0" format="BFYX"/>
<Tensor arg-index="1" type="output" port-index="0" format="BFYX"/>
</Buffers>
<CompilerOptions options="-cl-mad-enable"/>
<WorkSizes global="X,Y,B*F"/>
</CustomLayer>
カスタムレイヤーの組み込み定義¶
次の表には、ユーザーソースの前に付加される定義が含まれています。
例については、カーネルの例をご覧ください。
名前 |
値 |
|---|---|
|
このカーネルにバインドされた入力テンソルの数 |
|
このカーネルの実行に使用されるグローバル・ワーク・サイズの配列 |
|
|
|
このカーネルの実行に使用されるローカル・ワーク・サイズの配列 |
|
|
|
テンソル次元サイズの配列。順番は常に |
|
|
|
テンソルのデータタイプ: |
|
テンソルの形式、BFYX、BYXF、YXFB、FYXB、または ANY。形式は定義された名前に連結されます。テンソル形式を使用して、 |
|
開始前にテンソル次元に使用されるパディング要素の配列。常に BFYX の順番になります。 |
|
|
|
テンソル次元の終了後に使用されるパディング要素の配列。常に BFYX の順番になります。 |
|
|
|
各次元に隣接する要素間のオフセット (要素単位)。常に BFYX の順番になります。 |
|
|
|
低位のパディングをバイパスした、テンソルの開始から最初の有効な要素までの要素の数 |
次の例に示すように、すべての <TENSOR> 値は、INPUT0、INPUT1、および OUTPUT0 など、この操作にバインドされたすべてのテンソルに対して自動的に定義されます。
#define INPUT0_DIMS_SIZE 4
#define INPUT0_DIMS (int []){ 1,96,55,55, }
カーネルの例¶
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
__kernel void example_relu_kernel(
const __global INPUT0_TYPE* input0,
__global OUTPUT0_TYPE* output)
{
const uint idx = get_global_id(0);
const uint idy = get_global_id(1);
const uint idbf = get_global_id(2); // batches*features, as OpenCL supports 3D nd-ranges only
const uint feature = idbf % OUTPUT0_DIMS[1];
const uint batch = idbf / OUTPUT0_DIMS[1];
//notice that pitches are in elements, not in bytes!
const uint in_id = batch*INPUT0_PITCHES[0] + feature*INPUT0_PITCHES[1] + idy*INPUT0_PITCHES[2] + idx*INPUT0_PITCHES[3] + INPUT0_OFFSET;
const uint out_id = batch*OUTPUT0_PITCHES[0] + feature*OUTPUT0_PITCHES[1] + idy*OUTPUT0_PITCHES[2] + idx*OUTPUT0_PITCHES[3] + OUTPUT0_OFFSET;
INPUT0_TYPE value = input0[in_id];
// neg_slope (which is non-zero for leaky ReLU) is put automatically as #define, refer to the config xml
output[out_id] = value < 0 ? value * neg_slope : value;
}
注
前のセクションで説明したように、INPUT0_TYPE などすべての項目は、効率上の理由から、OpenVINO によって実際には OpenCL* (プリ) コンパイラー入力として定義されます。結果のデバッグについては、以下のデバッグのヒントを参照してください。