この記事はJulia Advent Calender 2018 2日目の記事です
OpenCLはKhronos Groupにより策定されている、マルチコアCPUやGPU等の計算資源を利用した並列コンピューティングのためのフレームワークです。この記事ではこのOpenCLのJuliaバインディングであるOpenCL.jlを使ってJuliaからGPUで簡単な計算をしてみます。
JuliaGPU
JuliaからGPUを使う方法として、上で述べたOpenCL.jlに加えてnvidiaのGPU向けフレームワークであるCUDAのためのライブラリがGitHub上で開発されています
GPUを使う場合にはOpenCLでもCUDAでも、ホスト側の処理とデバイス側の処理を記述する必要があります。最近ではデバイス側の処理をプログラマではなくライブラリが生成してくれるものも多いですが、両方が必要であることは変わりません。
今回紹介するOpenCL.jlはホスト側の処理を行うためのJuliaバインディングで、デバイス側にはOpenCL Cと呼ばれる言語を用います(あるいはOpenCL 2.2で追加されたOpenCL C++が動くかもしれません）。
一方（今回紹介しない）CUDA用には、JuliaのコードからGPUのデバイスコードに変換するライブラリCUDAnative.jlが開発されています。これは以前RustでCUDAカーネルを書くで解説した方法とだいたい同じで、JuliaのコードをLLVM IRにコンパイルした後、通常のようにホストのアーキテクチャ（x86_64)にコンパイルせずにPTXとよばれるGPU上のアセンブラのような形式にコンパイルします。そしてコンパイルされたデバイス側の処理を実行するためのホスト側のAPIがCUDA Deriver APIと呼ばれているものですが、これのJuliaバインディングがCUDAdrv.jlです。
さらにハイレベルのAPIとしてCLArray.jlやGPUArray.jlがあるようですが、これについては次回以降に調べたいと思います（次回があるとは言ってない）。
OpenCL.jl
まずはOpenCL.jlをインストールしましょう。
- Julia v"0.4.x" is supported on the release-0.4 branch and the OpenCL.jl versions v"0.4.x". Only bug-fixes will be applied.
- Julia v"0.5.x" is supported on the master branch and the OpenCL.jl versions v"0.5.x".
- Julia v"0.6.x" is experimentally supported on the master branch and the OpenCL.jl versions v"0.5.x"
READMEに書いてあるサポートバージョンは以上の通りです。私の環境はArchLinuxなのですが、この記事を書いている段階でpacmanで入るJuliaは1.0.1なのでこれでは使えません。
幸い最近Julia 1.0用のパッチがマージされているのでmasterのOpenCL.jlは使用できます。
git clone https://github.com/JuliaGPU/OpenCL.jl
cd OpenCL.jl
julia
]でpkgモードに入って
(v1.0) pkg> add .
のようにインストールできます（Gitリポジトリからインストールできると思うんですが、上手くインストールできなかった）。
これでOpenCLがインストールされました
using OpenCL
でインポートできます（システムにOpenCLが見つからないとここでプリコンパイルが失敗します）。
では公式のサンプルを見ていきましょう
const sum_kernel = "
__kernel void sum(__global const float *a,
__global const float *b,
__global float *c)
{
int gid = get_global_id(0);
c[gid] = a[gid] + b[gid];
}
"
a = rand(Float32, 50_000)
b = rand(Float32, 50_000)
device, ctx, queue = cl.create_compute_context()
a_buff = cl.Buffer(Float32, ctx, (:r, :copy), hostbuf=a)
b_buff = cl.Buffer(Float32, ctx, (:r, :copy), hostbuf=b)
c_buff = cl.Buffer(Float32, ctx, :w, length(a))
p = cl.Program(ctx, source=sum_kernel) |> cl.build!
k = cl.Kernel(p, "sum")
queue(k, size(a), nothing, a_buff, b_buff, c_buff)
r = cl.read(queue, c_buff)
if isapprox(norm(r - (a+b)), zero(Float32))
info("Success!")
else
error("Norm should be 0.0f")
end
カーネル側の言語は完全にOpenCL/Cで、OpenCL.jlの仕事はホスト側のOpenCL C APIのハンドリングと、Juliaの管理しているベクトル(a, b)の橋渡しですね。カーネル側をJuliaで書く試みはCLArray.jlがTranspiler.jl経由で行っていて、JuliaのASTからOpenCL/Cのコードを生成して呼び出しているようです。
a_buff = cl.Buffer(Float32, ctx, (:r, :copy), hostbuf=a)
のように:から始まっているものはSymbolと呼ばれ、以下で詳しく解説されています。
:rの部分には:r, :w, :rwが入り、それぞれ読み込み専用/書き込み専用/読み書き可能の意味です。
:copyの部分には:allocと:useが入り、それぞれclCreateBufferのフラグであるCL_MEM_COPY_HOST_PTR/CL_MEM_COPY_HOST_PTR/CL_MEM_COPY_HOST_PTRに対応します（フラグの意味はOpenCLのドキュメントを参照してください）。
まとめ
JuliaからOpenCLで記述された処理をGPUに実行させることが出来た。次はJuliaでGPU上の処理を記述する部分を調べる。