コグノスケ

OpenCLのOSS実装poclを調べる その7 - poclとLLVMとOpenCLカーネルビルド

目次: OpenCL

今回は少し話題を変えてpoclとClang/LLVMの関係について説明します。poclはOpenCL C言語をビルドするため内部的にClang/LLVMを呼び出します。ビルドはおおまかに2段階に分かれています。設定次第で多少変わりますが、ざっくり各段階で何をしているのか説明します。まだ話題にしていない関数も出てきますが、その話は追々やっていこうと思います。

第1段階OpenCL C -> LLVM IR

最初はclBuildProgram() です。やっていることはOpenCL C -> LLVM IR Bitcode（コンパイル＋OpenCLランタイムとのリンク）です。処理に関連する関数と呼び出し関係は下記のとおりです。

clBuildProgram()
  compile_and_link_program()
    pocl_llvm_build_program(): プリプロセス、コンパイル（出力LLVM IR Bitcode）
      link(): OpenCLランタイムとリンク（未定義シンボルの検出）
      pocl_write_module(): 出力LLVM IR Bitcode

この関数の実行中に3つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、

• tempfile-xx-xx-xx-xx-xx.cl: オリジナルのOpenCL Cコード（xx-xx-... の部分はハッシュ値）
• tempfile-xx-xx-xx-xx-xx.preproc.cl: プリプロセス後のOpenCL Cコード（xx-xx-... の部分はハッシュ値）
• ~/.cache/pocl/kcache/XX/XXXXXXXX/program.bc: リンク後のLLVM IR Bitcode（XXはハッシュ値なので名前は様々）

が作成されます。link() は全てのシンボルを問答無用で追加するのではなくて、OpenCLカーネルコードから参照されているシンボル（＝未定義のシンボル）のみを追加します。

第2段階LLVM IR -> バイナリ

次はclEnqueueNDRangeKernel() です。やることはLLVM IR Bitcode -> ターゲットバイナリです。

API名はコンパイルやビルドと関係なさそうに見えますが、poclのホストCPU向け実装を見る限り、NDRangeKernelから起因してLLVM IR Bitcodeからターゲットデバイスのバイナリに変換しています（あと何度も変換しなくて良いように生成したバイナリをキャッシュする）。

処理に関連する関数と呼び出し関係は下記のとおりです。

clEnqueueNDRangeKernel()
  pocl_command_enqueue()
    pocl_pthread_submit()
      pthread_scheduler_push_command(): キューにコマンドを追加する、ワーカースレッドが起床してコマンドを得る

start_thread(): ワーカースレッド
  pocl_pthread_driver_thread()
    pthread_scheduler_get_work()
    check_cmd_queue_for_device(): キューからコマンドを得る、コマンドがNDRangeKernelの実行要求だったら、下記を呼ぶ
      pocl_pthread_prepare_kernel()
        pocl_check_kernel_dlhandle_cache(): dlopen()
          pocl_check_kernel_disk_cache(): キャッシュ
            llvm_codegen(): ターゲットデバイスバイナリ生成
              pocl_llvm_codegen(): .so.oバイナリ作成
              pocl_invoke_clang(): .soバイナリ作成

独自アクセラレータ向け実装もホストCPU向け実装に習いclEnqueueNDRangeKernel() を起点にバイナリに変換を行います。現状はキャッシュ機構はスキップし、毎回ターゲットバイナリを生成する実装です。このままだと非常に遅いので今後、改善する必要があります。

clEnqueueNDRangeKernel()
  pocl_command_enqueue()
    pocl_accel_submit(): ここはターゲットデバイスごとに実装が異なる
      scheduleCommands()
        pocl_exec_command()
          pocl_accel_run(): ここから先はターゲットデバイスごとに実装が異なる
            llvm_codegen(): ターゲットデバイスバイナリ生成
              pocl_llvm_codegen(): .so.oバイナリ作成
              pocl_invoke_clang(): .soバイナリ作成

この関数の実行中に2つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、

• tempfile-xx-xx-xx-xx-xx.so.o: 中間バイナリ
• XX/XXXXXXXX/kernelName/xx-x-x/kernelName.so: 最終的なOpenCLカーネルバイナリ

が作成されます。中間バイナリと最終バイナリの差は、C言語のコンパイルのときに生成する *.oと実行ファイルとの違いと同様に、前者はアドレスが解決されておらず、後者はアドレス解決済みという差がありました。他にも何か違いがあるかもしれません。

第1段階で紹介した通りclBuildProgram() のlink() にてOpenCLカーネルに必要なシンボルが集められます。そのため中間バイナリと最終バイナリを比較しても、含まれる関数やシンボルはほぼ変わりません（例外はclang_rt.builtinsに依存した関数、例えば __adddf3() など）。