Go to Qiita Advent Calendar 2024 Top
LoginSignup
13
12

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

More than 3 years have passed since last update.

@terryky

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

Last updated at Posted 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を使います。

gl2posenet_m.png

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負荷高い)
delegate_off.png

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

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 ソースコードを clone 
git clone https://github.com/tensorflow/tensorflow.git 
cd tensorflow
git checkout r2.1

# TensorFlow Lite ライブラリのビルドに必要な関連ライブラリのダウンロード
./tensorflow/lite/tools/make/download_dependencies.sh

# TensorFlow Lite ライブラリ本体のビルド
> make -j 4  -f ./tensorflow/lite/tools/make/Makefile BUILD_WITH_NNAPI=false EXTRA_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/Makefile BUILD_WITH_NNAPI=false TARGET=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 -> eglCreateSync
eglDestroySyncKHR -> eglDestroySync
eglWaitSyncKHR -> eglWaitSync
eglClientWaitSyncKHR -> 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::BuiltinOpResolver resolver;

model = FlatBufferModel::BuildFromFile(POSENET_MODEL_PATH);
InterpreterBuilder(*model, resolver)(&interpreter);

/* ----- GpuDelegate V2 用に追加するブロックここから ----- */
#if defined (USE_GPU_DELEGATEV2)
    const TfLiteGpuDelegateOptionsV2 options = {
        .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$ ./gl2posenet
INFO: Created TensorFlow Lite delegate for GPU.
ERROR: Failed to build program executable - Build program failure(11:0) : error : program scope variable not in constant address space
(12:0) : error : program scope variable not in constant address space
(13:0) : error : program scope variable not in constant 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 OpenGL
Segmentation fault

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

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

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

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

tensorflow/lite/delegates/gpu/cl/cl_program.cc
Status CreateCLProgram(const std::string& code,
                       const std::string& compiler_options,
                       const CLContext& context, const CLDevice& device,
                       CLProgram* result) {
  int error_code;
  const char* source = code.c_str();
  fprintf (stderr, "%s\n", source);      // ★ 追加 ★

  cl_program program = clCreateProgramWithSource(context.context(), 1, &source,
                                                 nullptr, &error_code);
  if (!program || error_code != CL_SUCCESS) {
    return UnknownError(absl::StrCat("Failed to create compute program - ",
                                     CLErrorCodeToString(error_code)));
  }

  *result = CLProgram(program, device.id());
  RETURN_IF_ERROR(BuildProgram(program, device, compiler_options));
  return OkStatus();
}

すると、下記のような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_imagef
const sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
const sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
const sampler_t smp_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;
__kernel void main_function(
__global float4* src_data,
    __read_only image2d_t filters0,
    __read_only image2d_t filters1,
    __read_only image2d_t filters2,
    __read_only image2d_t filters3,
    __read_only image2d_t biases,
    float relu_clip1,
__global float4* dst_data,
    int4 src_size,
    int4 dst_size,
    int2 kernel_size,
    int2 dilation,
    int2 stride,
    int2 padding
) {
  int X = get_global_id(0) * 2;
  int Y = get_global_id(1) * 2;
  int Z = get_global_id(2) * 2;
  if (X >= dst_size.x || Y >= dst_size.y || Z >= dst_size.z) return;
  int xc0 = (X +0) * stride.x + padding.x;
  int xc1 = (X +1) * stride.x + padding.x;
  int yc0 = (Y +0) * stride.y + padding.y;
  int yc1 = (Y +1) * stride.y + padding.y;
  ACCUM_FLT4 r0 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r1 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r2 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r3 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r4 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r5 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r6 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  ACCUM_FLT4 r7 = (ACCUM_FLT4)(0.0f, 0.0f, 0.0f, 0.0f);
  int cx0;
  int cx1;
  int cy0;
  int cy1;
  int filter_offset = 0;
  for (int y = 0; y < kernel_size.y; ++y) {
  cy0 = y * dilation.y + yc0;
  cy1 = y * dilation.y + yc1;
  bool in_y0 = cy0 >= 0 && cy0 < src_size.y;
    cy0 = clamp(cy0, 0, src_size.y - 1);
  bool in_y1 = cy1 >= 0 && cy1 < src_size.y;
    cy1 = clamp(cy1, 0, src_size.y - 1);
  for (int x = 0; x < kernel_size.x; ++x) {
  cx0 = x * dilation.x + xc0;
  cx1 = x * dilation.x + xc1;
  bool in_x0 = cx0 >= 0 && cx0 < src_size.x;
    cx0 = clamp(cx0, 0, src_size.x - 1);
  bool in_x1 = cx1 >= 0 && cx1 < src_size.x;
    cx1 = clamp(cx1, 0, src_size.x - 1);
  int addr_0 = cy0 * src_size.x + cx0;
  int addr_2 = cy1 * src_size.x + cx0;
  int addr_1 = cy0 * src_size.x + cx1;
  int addr_3 = cy1 * src_size.x + cx1;
  int dz = src_size.x * src_size.y;
  for (int s = 0; s < src_size.z; ++s) {
    FLT4 src0 = src_data[addr_0] * (FLT)(in_x0 && in_y0); addr_0 += dz;
    FLT4 src2 = src_data[addr_2] * (FLT)(in_x0 && in_y1); addr_2 += dz;
    FLT4 src1 = src_data[addr_1] * (FLT)(in_x1 && in_y0); addr_1 += dz;
    FLT4 src3 = src_data[addr_3] * (FLT)(in_x1 && in_y1); addr_3 += dz;
    FLT4 f0 = READ_IMAGE(filters0, smp_none, (int2)(Z + 0, filter_offset));
    FLT4 f1 = READ_IMAGE(filters1, smp_none, (int2)(Z + 0, filter_offset));
    FLT4 f2 = READ_IMAGE(filters2, smp_none, (int2)(Z + 0, filter_offset));
    FLT4 f3 = READ_IMAGE(filters3, smp_none, (int2)(Z + 0, filter_offset));
    FLT4 f4 = READ_IMAGE(filters0, smp_none, (int2)(Z + 1, filter_offset));
    FLT4 f5 = READ_IMAGE(filters1, smp_none, (int2)(Z + 1, filter_offset));
    FLT4 f6 = READ_IMAGE(filters2, smp_none, (int2)(Z + 1, filter_offset));
    FLT4 f7 = 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) {
    FLT4 bias_val = READ_IMAGE(biases, smp_none, (int2)(Z, 0));
    {
      int xc = X + 0;
      int yc = Y + 0;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 1;
      int yc = Y + 0;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 0;
      int yc = Y + 1;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 1;
      int yc = Y + 1;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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) {
    FLT4 bias_val = READ_IMAGE(biases, smp_none, (int2)(Z, 0));
    {
      int xc = X + 0;
      int yc = Y + 0;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 1;
      int yc = Y + 0;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 0;
      int yc = Y + 1;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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;

      }
    }
    {
      int xc = X + 1;
      int yc = Y + 1;
      if (xc < dst_size.x && yc < dst_size.y) {
        FLT4 res = 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::string GetCommonDefines(CalculationsPrecision precision) {
  std::string result;
（略）
  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::string GetCommonDefines(CalculationsPrecision precision) {
  std::string result;
（略）
  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$ ./gl2posenet
INFO: Created TensorFlow Lite delegate for GPU.
/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 exist
error : Failed to load the temporary preprocessed file /tmp/cl-6DE612._0_PPED to source buffer
error : Failed to preprocess the source string #0

ERROR: TfLiteGpuDelegate Prepare: delegate is not initialized
ERROR: 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.
class ToTensorConverter : public OpenClConverterImpl {
 public:
（略）
    std::string shader_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. ソースコード

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

13
12
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
13
12

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?