More than 5 years have passed since last update.

Coral Edge TPU Dev Board で TensorFlow Lite GPU Delegate V2 (OpenCL) を試す

Last updated at2019-12-25Posted at 2019-11-17

1. 概要

　前回記事では、Coral EdgeTPU Dev Board 上で、TensorFlow Lite の GPU Delegate (OpenGLES版)を試しました。
　一方で、TensorFlow r2.1 ブランチ以降、GPU Delegate の実装は V2 となり、デフォルト動作がOpenCL となっています。OpenCL が使えないプラットフォームに限り、OpenGLES がfallback 動作となるようです。

　また、別記事において、EdgeTPU DevBoard はOpenCL 1.2 Full Profile (DEVICE_TYPE=GPU) をサポートしていることが確認できていますので、今回は、EdgeTPU DevBoard でのGPU Delegate V2 (OpenCL) 動作に挑戦します。

　題材は Posenet を動かし、下記のような結果表示をOpenGLESで行います。
　つまり、認識は OpenCL、結果表示は OpenGLES により、(使用するAPIは違えど) ともにGPUを使います。

1.1 先に挑戦結果

OpenCL Delegate、無事に動きました。
　1) 無理くりだけど GPU Delegate V2 (OpenCL) を有効にして Posenet を動かすことができました。
　2) 処理速度はDelegate有無で有意差はありませんでしたが、CPU負荷は劇的に下がっていることが確認できました。

比較項目	GPU Delegateなし	GPU Delegateあり
処理時間	340[ms]	320 [ms]
CPU負荷	201 [%]	10 [%]

■（ご参考）GPU Delegate無し時の htop　(CPU負荷高い)

■（ご参考）GPU Delegate あり時の htop　(CPU負荷低い)

1.2 ソースコード

GitHubで公開しています。
https://github.com/terryky/tflite_gles_app/tree/master/gl2posenet

2 TensorFlow Lite GPU Delegate V2 ライブラリのビルド

　ここからは、TensorFlow Lite GPU Delegate V2　を動かすための具体手順を書いていきます。
　実際の作業手順に沿って書いているので、若干まわりくどい内容になっていますが、作業トレースされる方のお役に立てるかもとも思いからそのまま書いています。ご了承ください。

2.1 PC向けビルド

　GPU Delegate を使用するには、Tensorflow Lite 本体ライブラリのビルドに加えて、Delegateライブラリを追加でビルドする必要があります。公式ドキュメントには、Android/iOS 向けの手順しか記載されていませんが、そこからの類推で、下記手順で Linux PC 向けに Delegateライブラリをビルドすることができました。

# TensorFlow ソースコードを clonegit clone https://github.com/tensorflow/tensorflow.gitcdtensorflowgit checkout r2.1# TensorFlow Lite ライブラリのビルドに必要な関連ライブラリのダウンロード./tensorflow/lite/tools/make/download_dependencies.sh# TensorFlow Lite ライブラリ本体のビルド> make-j 4-f ./tensorflow/lite/tools/make/MakefileBUILD_WITH_NNAPI=falseEXTRA_CXXFLAGS="-march=native"# GPU Delegate ライブラリのビルド> bazel build-s-c opt--copt="-DMESA_EGL_NO_X11_HEADERS" tensorflow/lite/delegates/gpu:delegate

　公式ドキュメントには記載がありませんが、Delegate ライブラリビルド時にMESA_EGL_NO_X11_HEADERS オプション付与が必要でした。これがないと、X11ヘッダが global な namespace で定義する型が、TensorFlowの local namespace での型定義と猛烈に衝突してビルドがこけてしまいます。

上記手順により、X86_64 Linux 用の libtensorflow-lite.a および libdelegate.so が出来上がります。これをアプリとリンクします。

$ls-l tensorflow/lite/tools/make/gen/linux_x86_64/lib/-rw-rw-r-- 1 terryky terryky 4615918  11月 17 18:30 libtensorflow-lite.a$ls-l bazel-bin/tensorflow/lite/delegates/gpu/-r-xr-xr-x 1 terryky terryky 90896 11月 17 18:30 libdelegate.a*-r-xr-xr-x 1 terryky terryky 62328 11月 17 18:30 libdelegate.so*

(補足) download_dependencies.sh で下記のようなエラーになったら

$./tensorflow/lite/tools/make/download_dependencies.sh./tensorflow/lite/tools/make/download_dependencies.sh: line 59: 1: Usage: download_and_extract URL DIR

　2019/12/16 現在、download_dependencies.sh が指定するダウンロード先URLが正しくないようです。tensoflow/tensorflow/workspace.bzl に記載されているURLにあうように、download_dependencies.sh を直接編集します。

download_dependencies.sh

(修正前)EIGEN_URL="$(grep-o'http.*bitbucket.org/eigen/eigen/get/.*tar\.gz'"${BZL_FILE_PATH}" |grep-v mirror.tensorflow |head-n1)"(修正後)EIGEN_URL="$(grep-o'http.*gitlab.com/libeigen/eigen/-/.*tar\.gz'"${BZL_FILE_PATH}" |grep-v mirror.tensorflow |head-n1)"

2.2 aarch64 Linux 向けビルド

　本来なら、bazel でターゲットアーキを指定すれば aarch64 Linux 向けビルドができるはずだとは思うのですが、私のbazelスキルが低く、Webに記載のあった何通りかの方法で試してみても、どれもうまくいきませんでした（ライブラリのビルドがこけるとか、アプリとのリンク時に関数が足りないとか）。
　このあたり、前回記事から何も進化できていませんが、今回も、bazel がDelegateライブラリをビルドするためにコンパイルしているファイルを Tensorflow Lite 本体の Makefile に追記する、という力業でビルドしました。
　（どなたかスマートな方法を教えて頂ければ嬉しいです！）

　今回は、bazel のビルドログから、TensorFlow Lite の Makefile に追記すべき項目を抽出する安直ツールbazel2make を作ったので、ちょっとだけ楽に作業できました。

　bazel2make ツールは下記で公開しています。
　https://github.com/terryky/tflite_gles_app/tree/master/tools/bazel2make

具体作業としては、下記のとおりです。　

# bazel ビルドログをファイルに書き出す$bazel build-s-c opt--copt="-DMESA_EGL_NO_X11_HEADERS" tensorflow/lite/delegates/gpu:delegate &> bazel_log.txt# bazel ビルドログから makefile に追記すべき項目を抽出する# 例) CORE_CC_ALL_SRCS += tensorflow/lite/delegates/gpu/cl/cl_device.cc$bazel2log bazel_log.txt# makefile に追記する（ここは手作業）$emacs tensorflow/lite/tools/make/Makefile# aarch64 指定で libtensorflowlite.a をビルドする$make-j 4-f ./tensorflow/lite/tools/make/MakefileBUILD_WITH_NNAPI=falseTARGET=aarch64

　上記手順を行うことで、aarch64 Linux 向けの、TensorFlow Lite 本体と GPU Delegate とが一体化した大きなライブラリが１つ出来上がります。
　あとは、このライブラリを EdgeTPU Devboard へコピーし、EdgeTPU Devboard 上で、アプリをネイティブビルドします。

$ls-s tensorflow/lite/tools/make/gen/aarch64_armv8-a/lib/-rw-rw-r-- 1 terryky terryky 18649770 11月 17 18:30 libtensorflow-lite.a

(補足1) 「fp16.h がない」と言われてビルドがコケたら

　Bazelビルド時に自動でダウンロードされていたファイルがキャッシュディレクトリに残っているので、それを使います。

■TensorFlow を git clone したディレクトリに、Bazel キャッシュディレクトリへのリンクを張る。

$cd$TENSORFLOW_TOP_DIR$ln-s ~/.cache/bazel/_bazel_hoge/xxxxxxxxxx/external.

■Makefile にインクルードパスを追記

tensorflow/lite/tools/make/Makefile

CXXFLAGS +=-Iexternal/FP16/include/

(補足2) 「compiled_program_cache_generated.h がない」と言われてビルドがコケたら

　Bazelビルド時に生成されているファイルがキャッシュに残っているので、それをコピーします。

$cd$TENSORFLOW_TOP_DIR$cp ~/.cache/bazel/_bazel_hoge/xxxxxxxxxx/execroot/org_tensorflow/bazel-out/k8-opt/bin/tensorflow/lite/delegates/gpu/cl/compiled_program_cache_generated.h ./tensorflow/lite/delegates/gpu/cl/

(補足3) 「eglCreateSyncKHRなんて知らない」と言われてビルドがコケたら

　関数名を変更しましょう。

tensorflow/lite/delegates/gpu/cl/egl_sync.cc

eglCreateSyncKHR->eglCreateSynceglDestroySyncKHR->eglDestroySynceglWaitSyncKHR->eglWaitSynceglClientWaitSyncKHR->eglClientWaitSync

3. GPU Delegate V2 (OpenCL) を使うアプリのビルド

　ライブラリのビルドができたので、ここからはアプリをビルドしていきます。

3.1 アプリの修正

　GPU Delegate を使うためには、アプリ側にも若干修正が必要です。
　Delegate を使わないアプリがすでに動いているなら、Interpreter 構築時に、下記おまじないを追加すればOKです。

/* 追加するヘッダファイル */# include "tensorflow/lite/delegates/gpu/delegate.h"/* 通常通りInterpreter構築 */unique_ptr<FlatBufferModel>model;unique_ptr<Interpreter>interpreter;ops::builtin::BuiltinOpResolverresolver;model=FlatBufferModel::BuildFromFile(POSENET_MODEL_PATH);InterpreterBuilder(*model,resolver)(&interpreter);/* ----- GpuDelegate V2 用に追加するブロックここから ----- */# if defined (USE_GPU_DELEGATEV2)constTfLiteGpuDelegateOptionsV2options={.is_precision_loss_allowed=1,// FP16.inference_preference=TFLITE_GPU_INFERENCE_PREFERENCE_FAST_SINGLE_ANSWER};auto*delegate=TfLiteGpuDelegateV2Create(&options);interpreter->ModifyGraphWithDelegate(delegate);# endif/* ----- GpuDelegate 用に追加するブロックここまで ----- */interpreter->AllocateTensors();

3.2 アプリのビルド

　上記修正したアプリソースをビルドし、aarch64 向けにクロスビルドした TensorFlow Lite ライブラリとリンクします。
　なお、EdgeTPU Devboard 上で画面表示するには、EdgeTPU Devboard の WindowSystem (wayland) 環境にあうように OpenGLES を初期化する必要がありますが、そのあたり、コマンド一発でビルドできるようにスクリプトを用意しました。

mendel@green-rabbit:~/tflite_gles_app/gl2posenet$./build_edgetpu_devboard.sh

4. 実行結果＆エラー回避

　アプリのビルドに成功したので、さっそく、EdgeTPU Devboard 上で動かします。
　が、下記のようなエラーメッセージが表示され終了してしまいました。

　悲しみをこらえつつ、解析してみます。

mendel@green-rabbit:~/tflite_gles_app/gl2posenet$./gl2posenetINFO: Created TensorFlow Lite delegateforGPU.ERROR: Failed to build program executable - Build program failure(11:0) : error : program scope variable notinconstant address space(12:0) : error : program scope variable notinconstant address space(13:0) : error : program scope variable notinconstant address space(85:0) : error : undefined identifier:'smp_none'(86:0) : error : undefined identifier:'smp_none'(87:0) : error : undefined identifier:'smp_none'(88:0) : error : undefined identifier:'smp_none'(89:0) : error : undefined identifier:'smp_none'(90:0) : error : undefined identifier:'smp_none'(91:0) : error : undefined identifier:'smp_none'(92:0) : error : undefined identifier:'smp_none'(106:0) : error : undefined identifier:'smp_none'(150:0) : error : undefined identifier:'smp_none'ERROR: Falling back to OpenGLSegmentation fault

4.1 OpenCL カーネルのランタイムビルドエラーその１ (__constant 修飾子)

　エラーログから推測するに、OpenCL で用いるカーネルコードをランタイムコンパイラがビルドする時にエラーとなっているようです。「カーネルの11行目で、プログラムスコープ変数が constant アドレス空間にない」と怒られているようです。

　といわれても、ここまで、OpenCL Delegate がどんなOpenCLカーネルを使うのか全く意識しないまま作業を進めてきたので、「カーネルの11行目」といわれても何のこっちゃわかりません。

　OpenCLカーネルのコンパイルは下記で行っているようなので、どんなカーネルソースコードなのか、printf() してみます。

tensorflow/lite/delegates/gpu/cl/cl_program.cc

StatusCreateCLProgram(conststd::string&code,conststd::string&compiler_options,constCLContext&context,constCLDevice&device,CLProgram*result){interror_code;constchar*source=code.c_str();fprintf(stderr,"%s\n",source);// ★ 追加 ★cl_programprogram=clCreateProgramWithSource(context.context(),1,&source,nullptr,&error_code);if(!program||error_code!=CL_SUCCESS){returnUnknownError(absl::StrCat("Failed to create compute program - ",CLErrorCodeToString(error_code)));}*result=CLProgram(program,device.id());RETURN_IF_ERROR(BuildProgram(program,device,compiler_options));returnOkStatus();}

すると、下記のようなOpenCLカーネルコードが出力されます。

OpenCL-kernel

# define ACCUM_FLT4 float4# define FLT float# define FLT2 float2# define FLT3 float3# define FLT4 float4# define TO_FLT4 convert_float4# define TO_ACCUM_TYPE convert_float4# define TO_ACCUM_FLT convert_float# define READ_IMAGE read_imagef# define WRITE_IMAGE write_imagefconstsampler_tsmp_edge=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP_TO_EDGE|CLK_FILTER_NEAREST;constsampler_tsmp_none=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_NONE|CLK_FILTER_NEAREST;constsampler_tsmp_zero=CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP|CLK_FILTER_NEAREST;# define CONV0(R, S)    \R += S.x * f0; \R += S.y * f1; \R += S.z * f2; \R += S.w * f3;# define CONV1(R, S)    \R += S.x * f4; \R += S.y * f5; \R += S.z * f6; \R += S.w * f7;__kernelvoidmain_function(__globalfloat4*src_data,__read_onlyimage2d_tfilters0,__read_onlyimage2d_tfilters1,__read_onlyimage2d_tfilters2,__read_onlyimage2d_tfilters3,__read_onlyimage2d_tbiases,floatrelu_clip1,__globalfloat4*dst_data,int4src_size,int4dst_size,int2kernel_size,int2dilation,int2stride,int2padding){intX=get_global_id(0)*2;intY=get_global_id(1)*2;intZ=get_global_id(2)*2;if(X>=dst_size.x||Y>=dst_size.y||Z>=dst_size.z)return;intxc0=(X+0)*stride.x+padding.x;intxc1=(X+1)*stride.x+padding.x;intyc0=(Y+0)*stride.y+padding.y;intyc1=(Y+1)*stride.y+padding.y;ACCUM_FLT4r0=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r1=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r2=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r3=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r4=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r5=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r6=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);ACCUM_FLT4r7=(ACCUM_FLT4)(0.0f,0.0f,0.0f,0.0f);intcx0;intcx1;intcy0;intcy1;intfilter_offset=0;for(inty=0;y<kernel_size.y;++y){cy0=y*dilation.y+yc0;cy1=y*dilation.y+yc1;boolin_y0=cy0>=0&&cy0<src_size.y;cy0=clamp(cy0,0,src_size.y-1);boolin_y1=cy1>=0&&cy1<src_size.y;cy1=clamp(cy1,0,src_size.y-1);for(intx=0;x<kernel_size.x;++x){cx0=x*dilation.x+xc0;cx1=x*dilation.x+xc1;boolin_x0=cx0>=0&&cx0<src_size.x;cx0=clamp(cx0,0,src_size.x-1);boolin_x1=cx1>=0&&cx1<src_size.x;cx1=clamp(cx1,0,src_size.x-1);intaddr_0=cy0*src_size.x+cx0;intaddr_2=cy1*src_size.x+cx0;intaddr_1=cy0*src_size.x+cx1;intaddr_3=cy1*src_size.x+cx1;intdz=src_size.x*src_size.y;for(ints=0;s<src_size.z;++s){FLT4src0=src_data[addr_0]*(FLT)(in_x0&&in_y0);addr_0+=dz;FLT4src2=src_data[addr_2]*(FLT)(in_x0&&in_y1);addr_2+=dz;FLT4src1=src_data[addr_1]*(FLT)(in_x1&&in_y0);addr_1+=dz;FLT4src3=src_data[addr_3]*(FLT)(in_x1&&in_y1);addr_3+=dz;FLT4f0=READ_IMAGE(filters0,smp_none,(int2)(Z+0,filter_offset));FLT4f1=READ_IMAGE(filters1,smp_none,(int2)(Z+0,filter_offset));FLT4f2=READ_IMAGE(filters2,smp_none,(int2)(Z+0,filter_offset));FLT4f3=READ_IMAGE(filters3,smp_none,(int2)(Z+0,filter_offset));FLT4f4=READ_IMAGE(filters0,smp_none,(int2)(Z+1,filter_offset));FLT4f5=READ_IMAGE(filters1,smp_none,(int2)(Z+1,filter_offset));FLT4f6=READ_IMAGE(filters2,smp_none,(int2)(Z+1,filter_offset));FLT4f7=READ_IMAGE(filters3,smp_none,(int2)(Z+1,filter_offset));CONV0(r0,src0);CONV0(r1,src1);CONV0(r2,src2);CONV0(r3,src3);CONV1(r4,src0);CONV1(r5,src1);CONV1(r6,src2);CONV1(r7,src3);filter_offset++;}}}if(Z<dst_size.z){FLT4bias_val=READ_IMAGE(biases,smp_none,(int2)(Z,0));{intxc=X+0;intyc=Y+0;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r0)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+1;intyc=Y+0;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r1)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+0;intyc=Y+1;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r2)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+1;intyc=Y+1;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r3)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}}Z++;if(Z<dst_size.z){FLT4bias_val=READ_IMAGE(biases,smp_none,(int2)(Z,0));{intxc=X+0;intyc=Y+0;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r4)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+1;intyc=Y+0;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r5)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+0;intyc=Y+1;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r6)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}{intxc=X+1;intyc=Y+1;if(xc<dst_size.x&&yc<dst_size.y){FLT4res=TO_FLT4(r7)+bias_val;res=clamp(res,(FLT)(0.0f),(FLT)(relu_clip1));dst_data[(((Z)*dst_size.y+(yc))*dst_size.x+(xc))]=res;}}}Z++;}

　上記カーネルコードの11行目でコンパイルエラーだと怒られているので、const アドレス空間修飾子が、OpenCL C 規約に沿ってないことが問題のようです。
　TensorFlow Lite の OpenCL Delegate でカーネルコードを生成している処理を下記のように修正することで、この問題は回避することができました。

OpenCL カーネルコンパイルエラー回避方法（その１）

tensorflow/lite/delegates/gpu/cl/kernels/util.cc（修正前）

std::stringGetCommonDefines(CalculationsPrecisionprecision){std::stringresult;（略）result+="const sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | ""CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n";（略）

tensorflow/lite/delegates/gpu/cl/kernels/util.cc（修正後）

std::stringGetCommonDefines(CalculationsPrecisionprecision){std::stringresult;（略）result+="__constant sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | ""CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n";（略）

4.2 OpenCL カーネルのランタイムビルドエラーその２ (extension FP16)

　問題を一つ解決して少し前進したものの、次は下記のようなエラーメッセージを吐いて終了してしまいました。

mendel@green-rabbit:~/tflite_gles_app/gl2posenet$./gl2posenetINFO: Created TensorFlow Lite delegateforGPU./tmp/cl-6DE611.:2:38: error: can't enable a non-supported extension.# pragma OPENCL EXTENSION cl_khr_fp16 : enable                                     ^1 error generated.ERROR: TfLiteGpuDelegate Init: Failed to build program executable - Build program failureerror : File /tmp/cl-6DE612._0_PPED does not existerror : Failed to load the temporary preprocessed file /tmp/cl-6DE612._0_PPED to source buffererror : Failed to preprocess the source string #0ERROR: TfLiteGpuDelegate Prepare: delegate is not initializedERROR: Node number 31 (TfLiteGpuDelegateV2) failed to prepare.ERROR: Restored previous execution plan after delegate application failure.Segmentation fault

　今回のエラーはわかりやすいです。EdgeTPU Devboard がサポートしてない OpenCL extention を有効化しようとしてエラーとなっていました。
　実際、EdgeTPU Devboard の OpenCL は、fp16 をサポートしてないことは [前回記事] (https://qiita.com/terryky/items/714b29af1d405175211d) で確認済みです。

　GPU実行において、fp16 が使えないのは何とも歯がゆいですが、EdgeTPU Devboard に搭載されている GPU (VIVANTE GC7000L) の OpenCL がサポートしてないのでどうしようもないです。fp16 実行を諦めざるを得ません。

　TensorFlow Lite の OpenCL Delegate でカーネルコードを生成している処理を下記のように fp16 extensionを使わないように修正することで、この問題も回避することができました。

OpenCL カーネルコンパイルエラー回避方法（その２）

tensorflow/lite/delegates/gpu/cl/kernels/converter.cc（修正後）

// Implements conversion from BHWC to OpenCL-specific tensor layout.classToTensorConverter:publicOpenClConverterImpl{public:（略）std::stringshader_src=R"(// #pragma OPENCL EXTENSION cl_khr_fp16 : enable    ★ この行をコメントアウト ★（略）

　上記２つの問題を修正することで、無事に、EdgeTPU Devboard 上で TensorFlow Lite GPU Delegate V2 (OpenCL) を動作させることができました。

　Posenet の姿勢推定結果も正しく得られています。
　処理速度は冒頭に記載の通りです。

5. 最後に

　いろいろ苦労はしましたが、何とか OpenCL Delegate を動かすことができました。

　OpenCL Delegate を使っても、CPU実行と比べてそんなに速くならなかったのですが、EdgeTPU Devboard に搭載されている GPU (VIVANTE GC7000L) だと、そんなものかなぁ、とも思います。(OpenCL カーネルを fp16 で動かせない時点で期待薄)。

　ただし、処理速度は速くならないけれど、純粋にCPU負荷を下げるオフロードエンジンとして使えるという意味では有効だと思います。

6. ソースコード