コグノスケ


link 未来から過去へ表示  link 過去から未来へ表示(*)

link もっと前
2021年7月2日 >>> 2021年7月2日
link もっと後

2021年7月2日

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() とLLVM IR Bitcode生成関数の呼び出し関係
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からターゲットデバイスのバイナリに変換しています(あと何度も変換しなくて良いように生成したバイナリをキャッシュする)。

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

ホストCPU向け実装の呼び出し関係
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() など)。

編集者:すずき(2023/09/24 11:58)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



link もっと前
2021年7月2日 >>> 2021年7月2日
link もっと後

管理用メニュー

link 記事を新規作成

<2021>
<<<07>>>
----123
45678910
11121314151617
18192021222324
25262728293031

最近のコメント5件

  • link 24年10月1日
    すずきさん (10/06 03:41)
    「xrdpで十分動作しているので、Wayl...」
  • link 24年10月1日
    hdkさん (10/03 19:05)
    「GNOMEをお使いでしたら今はWayla...」
  • link 24年10月1日
    すずきさん (10/03 10:12)
    「私は逆にVNCサーバーに繋ぐ使い方をした...」
  • link 24年10月1日
    hdkさん (10/03 08:30)
    「おー、面白いですね。xrdpはすでに立ち...」
  • link 14年6月13日
    2048player...さん (09/26 01:04)
    「最後に、この式を出すのに紙4枚(A4)も...」

最近の記事3件

  • link 24年10月28日
    すずき (10/30 23:49)
    「[Linuxからリモートデスクトップ] 目次: Linux開発用のLinuxマシンの画面を見るにはいろいろな手段がありますが、...」
  • link 23年4月10日
    すずき (10/30 23:46)
    「[Linux - まとめリンク] 目次: Linux関係の深いまとめリンク。目次: RISC-V目次: ROCK64/ROCK...」
  • link 24年10月24日
    すずき (10/25 02:35)
    「[ONKYOからM-AUDIOのUSB DACへ] 目次: PCかれこれ10年以上(2013年3月16日の日記参照)活躍してく...」
link もっとみる

こんてんつ

open/close wiki
open/close Linux JM
open/close Java API

過去の日記

open/close 2002年
open/close 2003年
open/close 2004年
open/close 2005年
open/close 2006年
open/close 2007年
open/close 2008年
open/close 2009年
open/close 2010年
open/close 2011年
open/close 2012年
open/close 2013年
open/close 2014年
open/close 2015年
open/close 2016年
open/close 2017年
open/close 2018年
open/close 2019年
open/close 2020年
open/close 2021年
open/close 2022年
open/close 2023年
open/close 2024年
open/close 過去日記について

その他の情報

open/close アクセス統計
open/close サーバ一覧
open/close サイトの情報

合計:  counter total
本日:  counter today

link About www2.katsuster.net
RDFファイル RSS 1.0

最終更新: 10/30 23:49