OpenCL.jlでGPGPU

この記事は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上で開発されています

https://github.com/JuliaGPU

GPUを使う場合にはOpenCLでもCUDAでも、ホスト側の処理とデバイス側の処理を記述する必要があります。最近ではデバイス側の処理をプログラマではなくライブラリが生成してくれるものも多いですが、両方が必要であることは変わりません。

今回紹介する[OpenCL.jl]はホスト側の処理を行うためのJuliaバインディングで、デバイス側には[OpenCL C]と呼ばれる言語を用います(あるいはOpenCL 2.2で追加された[OpenCL C++]が動くかもしれません）。

一方（今回紹介しない）CUDA用には、JuliaのコードからGPUのデバイスコードに変換するライブラリ[CUDAnative.jl]が開発されています。これは以前[RustでCUDAカーネルを書く](https://qiita.com/termoshtt/items/b98d5c46ab9c1ab1f7b6)で解説した方法とだいたい同じで、JuliaのコードをLLVM IRにコンパイルした後、通常のようにホストのアーキテクチャ（x86_64)にコンパイルせずにPTXとよばれるGPU上のアセンブラのような形式にコンパイルします。そしてコンパイルされたデバイス側の処理を実行するためのホスト側のAPIがCUDA Deriver APIと呼ばれているものですが、これのJuliaバインディングが[CUDAdrv.jl]です。

さらにハイレベルのAPIとして[CLArray.jl]や[GPUArray.jl]があるようですが、これについては次回以降に調べたいと思います（次回があるとは言ってない）。

[GPUArray.jl]: https://github.com/JuliaGPU/GPUArrays.jl
[CLArray.jl]: https://github.com/JuliaGPU/CLArrays.jl
[OpenCL.jl]: https://github.com/JuliaGPU/OpenCL.jl
[CUDAnative.jl]: https://github.com/JuliaGPU/CUDAnative.jl
[CUDAdrv.jl]: https://github.com/JuliaGPU/CUDAdrv.jl
[Transpiler.jl]: https://github.com/SimonDanisch/Transpiler.jl
[OpenCL C]: https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_C.html
[OpenCL C++]: https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_Cxx.html

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用のパッチがマージされている](https://github.com/JuliaGPU/OpenCL.jl/pull/170)のでmasterの[OpenCL.jl]は使用できます。

```
git clone https://github.com/JuliaGPU/OpenCL.jl
cd OpenCL.jl
julia
```

`]`でpkgモードに入って

```julia
(v1.0) pkg> add .
```

のようにインストールできます（[Gitリポジトリからインストールできる](https://docs.julialang.org/en/v1/stdlib/Pkg/index.html#Adding-unregistered-packages-1)と思うんですが、上手くインストールできなかった）。

これでOpenCLがインストールされました

```julia
using OpenCL
```

でインポートできます（システムにOpenCLが見つからないとここでプリコンパイルが失敗します）。

では公式のサンプルを見ていきましょう

```julia
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のコードを生成して呼び出しているようです。

```julia
a_buff = cl.Buffer(Float32, ctx, (:r, :copy), hostbuf=a)
```

のように`:`から始まっているものは`Symbol`と呼ばれ、以下で詳しく解説されています。

- [Juliaのシンボルとは? - りんごがでている](http://bicycle1885.hatenablog.com/entry/2015/08/15/143620)

`: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上の処理を記述する部分を調べる。