カスタム 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() メソッドを呼び出します。

Python

 core = ov.Core() 
core.set_property("GPU", {"CONFIG_FILE": "<path_to_the_xml_file>"})

C++

 ov::Core core; 
// GPU 拡張をロード 
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

重要

Open Model Zoo が廃止されたため、OpenVINO IR 形式のモデルは Hugging Face で公開されるようになりました。

設定ファイルの形式#

構成ファイルは、提供するすべてのカスタム操作に対してタイプ CustomLayer のノードを含む .xml ファイル構造に従うことが期待されます。

以下のセクションで説明する定義では、次の表記を使用します:

表記規則	説明
(0/1)	このノードまたは属性のインスタンスを 0 つまたは 1 つ持つことができます
	このノードまたは属性のインスタンスは 1 つだけ必要です
(0+)	このノードまたは属性のインスタンスはいくつでも持つことができます
(1+)	このノードまたは属性のインスタンスを 1 つ以上持つことができます

CustomLayer ノードとサブノードの構造#

CustomLayer ノードには、単一のカスタム操作の構成全体が含まれています。

属性名	#	説明
`name`		使用する操作タイプの名前。この名前は、IR で使用されるタイプと同一である必要があります。
`type`		`SimpleGPU` である必要があります。
`version`		`1` である必要があります。

サブノード: Kernel (1), Buffers (1), CompilerOptions (0+), WorkSizes (0/1)

カーネルノードとサブノードの構造#

Kernel ノードには、すべてのカーネル・ソースコード構成が含まれています。

サブノード: Source (1+), Define (0+)

ソースノードとサブノードの構造#

Source ノードは、単一の OpenCL* ソースファイルを指します。

属性名	#	説明
`filename`		OpenCL* ソースコードを含むファイルの名前。パスは実行可能ファイルからの相対パスです。複数のソースノードには、ソースが順番に連結されます。

サブノード: なし

ノードとサブノード構造の定義#

Define ノードは、コンパイル (JIT) 中にソースに追加される単一の #define 命令を構成します。

属性名	#	説明
`name`		定義された JIT の名前。静的定数の場合、これには文字列として受け取られる値も含めることができます。
`param`	(0/1)	このパラメーター値は、この JIT 定義の値として使用されます。
`type`	(0/1)	パラメーターのタイプです。受け入れられる値: `int`、 `float`、と `int[]`、`float[]` の配列。
`default`	(0/1)	指定されたパラメーターが OpenVINO IR の操作に含まれていない場合に使用されるデフォルト値。

サブノード: なし

結果の JIT は次の形式になります: #define [name] [type] [value/default]。

バッファーノードとサブノードの構造#

Buffers ノードは、OpenCL* エントリー関数のすべての入出力バッファーを構成します。バッファーノード構造は存在しません。

サブノード: Data (0+), Tensor (1+)

データノードとサブノードの構造#

Data ノードは、重みやバイアスなど静的データを含む単一の入力を構成します。

属性名	#	説明
`name`		OpenVINO IR の操作にアタッチされた BLOB の名前。
`arg-index`		バインドされるエントリー関数引数内の 0 から始まるインデックス。

サブノード: なし

テンソルノードとサブノードの構造#

Tensor ノードは、単一の入力または出力テンソルを構成します。

属性名	#	説明
`arg-index`		バインドされるエントリー関数引数内の 0 から始まるインデックス。
`type`		`input` または `output`
`port-index`		OpenVINO IR の操作入出力ポートの 0 から始まるインデックス
`format`	(0/1)	テンソルのデータレイアウト宣言。受け入れられる値 : `BFYX`、`BYXF`、`YXFB`、`FYXB`、すべて小文字の同じ値。デフォルト値: `BFYX`。

CompilerOptions ノードとサブノードの構造#

CompilerOptions ノードは、OpenCL* ソースのコンパイルフラグを設定します。

属性名	#	説明
`options`		OpenCL* コンパイラーに渡されるオプション文字列

サブノード: なし

WorkSizes ノードとサブノードの構造#

WorkSizes ノードは、OpenCL* プログラムを実行するためキューに投入する際に、グローバル/ローカル・ワーク・サイズを構成します。

属性名	#	説明
`global` `local`	(0/1) (0/1)	実行中に使用される OpenCL* のワークサイズを定義する、最大 3 つの整数または式の配列。数式では B、F、Y、X 次元の値を使用でき、次の演算子を含めることができます: +、-、/、、%。すべての演算子は整数算術で評価されます。デフォルト値: `global=”B\F\Y\X” local=””`
`dim`	(0/1)	ワークサイズを取得するテンソル。受け入れられる値: `input N`、`output`、ここで `N` は 0 で始まる入力テンソルのインデックスです。デフォルト値: `output`

サブノード: なし

設定ファイルの例#

次のサンプルコードは、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>

カスタムレイヤーのビルトイン定義#

次の表には、ユーザーソースの前に付加される定義が含まれています。

例については、カーネルの例をご覧ください。

名前	値
`NUM_INPUTS`	このカーネルにバインドされた入力テンソルの数
`GLOBAL_WORKSIZE`	このカーネルの実行に使用されるグローバル・ワーク・サイズの配列
`GLOBAL_WORKSIZE_SIZE`	`GLOBAL_WORKSIZE` 配列のサイズ
`LOCAL_WORKSIZE`	このカーネルの実行に使用されるローカル・ワーク・サイズの配列
`LOCAL_WORKSIZE_SIZE`	`LOCAL_WORKSIZE` 配列のサイズ
`<TENSOR>_DIMS`	テンソル次元サイズの配列。順番は常に `BFYX` です
`<TENSOR>_DIMS_SIZE`	`<TENSOR>_DIMS` 配列のサイズ。
`<TENSOR>_TYPE`	テンソルのデータタイプ: `float`、`half`、または `char`
`<TENSOR>_FORMAT_<TENSOR_FORMAT>`	テンソルの形式、 BFYX、BYXF、YXFB、FYXB、または ANY。形式は定義された名前に連結されます。テンソル形式を使用して、`#ifdef/#endif` でコード内のコードパスを定義できます。
`<TENSOR>_LOWER_PADDING`	開始前にテンソル次元に使用されるパディング要素の配列。常に BFYX の順番になります。
`<TENSOR>_LOWER_PADDING_SIZE`	`<TENSOR>_LOWER_PADDING` 配列のサイズ
`<TENSOR>_UPPER_PADDING`	テンソル次元の終了後に使用されるパディング要素の配列。常に BFYX の順番になります。
`<TENSOR>_UPPER_PADDING_SIZE`	`<TENSOR>_UPPER_PADDING` 配列のサイズ
`<TENSOR>_PITCHES`	各次元に隣接する要素間のオフセット (要素単位)。常に BFYX の順番になります。
`<TENSOR>_PITCHES_SIZE`	`<TENSOR>_PITCHES` 配列のサイズ
`<TENSOR>_OFFSET`	低位のパディングをバイパスした、テンソルの開始から最初の有効な要素までの要素の数。

次の例に示すように、すべての <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、OpenCL は 3D nd レンジのみをサポートしているため 
    const uint feature = idbf % OUTPUT0_DIMS[1]; 
    const uint batch = idbf / OUTPUT0_DIMS[1]; 
    // ピッチはバイトではなく要素単位であることに注意 
    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 (リーキー ReLU の場合はゼロ以外) は #define として自動的に配置されます。config xml を参照 
    output[out_id] = value < 0 ? value * neg_slope : value; 
}

注

前のセクションで説明したように、INPUT0_TYPE などすべての項目は、効率上の理由から、OpenVINO によって実際には OpenCL* (プリ) コンパイラー入力として定義されます。結果のデバッグについては、以下のデバッグのヒントを参照してください。

デバッグのヒント#

OpenCL™ カーネルで printf を使用します。

特定の値をデバッグするには、カーネルで printf を使用します。ただし、過度に出力すると、大量のデータが生成される可能性があるので注意してください。printf の出力は典型的なものであり、バッファーに収まるように出力を切り詰めることもできます。また、バッファリングにより、実行終了時に実際にはバッファー全体の出力が取得されます。

詳細は、printf 関数を参照してください。