code emitter for GCN - w

なんか、「OpenCLのコンパイルって超速くならんの？(スーパー意訳というか全然違う話をしてたが僕がコンパイル速ければすべて解決すると勝手に結論出した)」みたいな感じになったので、

とりあえず現状どうなってるか測ったところ、なんもしてないプログラムで 30msec かかるとかだった。

じゃあclCreateProgramWithBinaryだとどうかと思って、適当に試したところ、

ぐらいだった。

llvmir / amdil / GCN ISA の出しかたは、コンパイル時に、

とかを指定すると好きなのを出せる。http://clcc.sourceforge.net/ とか使うと試しやすいと思う。

というわけで、結論としては、GCN ISAを直吐きすれば、実行時リアルタイムコード生成もできそうというところだったのでできそうか試してみた。

https://github.com/tanakamura/ai-isa-jit
ai-isa-jit (AMD IL & AMD ISA JIT : あいあいさーじっとと読む。なお、ILはまだ実装されてない)

test/run.c をビルドして、できた run に一緒に付いてる f.cl を渡して実行すると、 f.cl をビルド後、その中身を書きかえてそれを実行する。

が全然まだ揃ってないけど、真面目に使いそうなら揃えると思う。

やってみた感触としては、かなり難しい。ISAはドキュメントあるけど、ABIが全く説明無いので、なんかよくわからない挙動をする。

例えば、一緒に入れてる、f.cl だが、

__kernel void f(__global int *in, __global int *out)
{
    out[get_global_id(0)] = 8;
}

これを、

__kernel void f(__global int *in, __global int *out)
{
    out[0] = 8;
}

こうしたり、

__kernel void f(__global int *in, __global int *out)
{
    out[get_global_id(0)] = in[get_global_id(0)];
}

こうしたりすると、引数の渡されかたが変わる。

まあMantleの詳細がわかってがなんかこのあたりうまくやってくれるようならもうちょっと真面目に考える。

あとELFの構造はスーパー決め打ちなのでこれでいつでも動くかどうかは知らない。(.textにELF入ってる変態フォーマットなのが全部悪い)