この記事はC++ Advent Calendar 2015の6日目の記事です。
話題は、3日目のRiyaaaaaさんの記事を見たとある人に「お前もC++AMPerならC++AMPの記事書けよ！！」って煽られたので、急遽変更して、C++AMPのネタにします。
元々の予定だった『C++11になぜmake_uniqueがないのか』については別記事に、アドベントカレンダーと関係ない記事として投稿したので、そっちをご覧ください。
なお、この記事の内容は某所で既に発表済みの内容を（ほぼ）そのまま持ってきたので、読んだことある人いたらごめんなさい。
お題は、C++AMPとOpenCL CとSPIR(1.2)の主に相互変換についてです（＋OpenCL 2.1とSPIR-Vのちょっとした紹介）。
C++AMPについては、先述のRyaaaaさんの記事か、随分前に投稿した私の記事も合わせてお読みください。
C++AMP→OpenCL C
Kalmarを使うとC++ AMPからOpenCL Cに変換できます。
Kalmarの詳しい使い方については先述の記事が詳しいので割愛します。
ただ、Kalmarを使うと言っても普通に使うだけだとOpenCL Cのコードは出てこないので、Kalmarを修正して自力ビルドする必要があります。
具体的には、mcwamp_opencl.cppの337行目付近で、
std::cout << source << std::endl;
してやってから、ビルドすれば、実行時に（※OpenCL Cのコンパイラはオンラインコンパイラなので）OpenCL Cのソースを取得することができます。
なお、この時
export CLAMP_RUNTIME=CL
export CLAMP_NOSPIR=1
しておかないと、SPIR (LLVM IR bitcode)が吐かれるのでBCなんちゃらという意味不明なのが出てくるので注意してください。
SPIRについては後述します。
ともかくこれでOpenCL Cを取得できるので、試しに単純な足し算
concurrency::parallel_for_each(
aa.get_extent(),
[aa, bb, &cc] (const concurrency::index<1> idx) restrict(amp)
{
const int i = idx[0];
cc[i] = aa[i] + bb[i];
});
をビルドしてみます。そうすると、結果、色々出てくるんですが、カーネル関数が
__kernel void ZZ4mainEN3_EC__119__cxxamp_trampolineEPiiiiiiiiS0_iiiiiiiS0_iiiiiii( __global unsigned int *llvm_cbe_tmp__1, unsigned int llvm_cbe_tmp__2, unsigned int llvm_cbe_tmp__3, unsigned int llvm_cbe_tmp__4, unsigned int llvm_cbe_tmp__5, unsigned int llvm_cbe_tmp__6, unsigned int llvm_cbe_tmp__7, unsigned int llvm_cbe_tmp__8, __global unsigned int *llvm_cbe_tmp__9, unsigned int llvm_cbe_tmp__10, unsigned int llvm_cbe_tmp__11, unsigned int llvm_cbe_tmp__12, unsigned int llvm_cbe_tmp__13, unsigned int llvm_cbe_tmp__14, unsigned int llvm_cbe_tmp__15, unsigned int llvm_cbe_tmp__16, __global unsigned int *llvm_cbe_tmp__17, unsigned int llvm_cbe_tmp__18, unsigned int llvm_cbe_tmp__19, unsigned int llvm_cbe_tmp__20, unsigned int llvm_cbe_tmp__21, unsigned int llvm_cbe_tmp__22, unsigned int llvm_cbe_tmp__23, unsigned int llvm_cbe_tmp__24) {
unsigned long llvm_cbe_tmp__25;
unsigned long llvm_cbe_tmp__26;
unsigned int llvm_cbe_tmp__27;
unsigned int llvm_cbe_tmp__28;
llvm_cbe_tmp__25 = /*tail*/ amp_get_global_id(0u);
llvm_cbe_tmp__26 = ((signed long )(((signed long )(llvm_cbe_tmp__25 << 32ul)) >> ((signed long )32ul)));
llvm_cbe_tmp__27 = *((&llvm_cbe_tmp__1[((signed long )(((unsigned long )(((unsigned long )(((unsigned long )(((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__8))) + ((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__6))))))) + ((unsigned long )llvm_cbe_tmp__26)))))]));
llvm_cbe_tmp__28 = *((&llvm_cbe_tmp__9[((signed long )(((unsigned long )(((unsigned long )(((unsigned long )(((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__16))) + ((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__14))))))) + ((unsigned long )llvm_cbe_tmp__26)))))]));
*((&llvm_cbe_tmp__17[((signed long )(((unsigned long )(((unsigned long )(((unsigned long )(((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__24))) + ((unsigned long )(((signed long )(signed int )llvm_cbe_tmp__22))))))) + ((unsigned long )llvm_cbe_tmp__26)))))])) = (((unsigned int )(((unsigned int )llvm_cbe_tmp__28) + ((unsigned int )llvm_cbe_tmp__27))));
return;
}
というのが最後の方に出てくるはずです。
でも、LLVMっぽいのはともかく、なんだこれって感じです・・・。大量のキャスト・・・？
よくわからないのでキャストを消したり変数名を変えたりしてみると、こんなコードになります。
__kernel void func(
__global unsigned int *src1, unsigned int unused2, unsigned int unused3, unsigned int unused4, unsigned int unused5, unsigned int rightOffset1, unsigned int unused7, unsigned int rightOffset0,
__global unsigned int *src2, unsigned int unused10, unsigned int unused11, unsigned int unused12, unsigned int unused13, unsigned int leftOffset1, unsigned int unused15, unsigned int leftOffset0,
__global unsigned int *dst, unsigned int unused18, unsigned int unused19, unsigned int unused20, unsigned int unused21, unsigned int dstOffset1, unsigned int unused23, unsigned int dstOffset0)
{
unsigned long index = get_global_id(0);
unsigned int right = src1[rightOffset0 + rightOffset1 + index];
unsigned int left = src2[leftOffset0 + leftOffset1 + index];
dst[dstOffset0 + dstOffset1 + index] = left + right;
return;
}
なんか余計なオフセット計算入ってますね・・・。あと未使用引数多すぎです。
まぁともかくOpenCL Cに変換はできました。OpenCL C書かなくてももっと手軽にOpenCL Cが手に入るのは便利（？）。
ただし、
struct Vector4
{
public:
alignas(32) double data[4];
Vector4() restrict(amp)
{
data[0] = 0;
data[1] = 0;
data[2] = 0;
data[3] = 0;
};
};
concurrency::parallel_for_each(
x.get_extent(),
[f, &x, &v, dt, m](const concurrency::index<1> idx) restrict(amp)
{
const int i = idx[0];
// a = f/m
Vector4 a;
for (int j = 0; j < 4; j++)
{
a.data[j] = f[i].data[j] /m;
}
// x += v*dt + a*dt^2/2
for (int j = 0; j < 4; j++)
{
const double dxv = v[i].data[j] * dt;
const double dxa = a.data[j] * dt*dt/2;
const double dx = dxv + dxa;
x[i].data[j] += dx;
}
// v += a*dt
for (int j = 0; j < 4; j++)
{
const double dv = a.data[j] * dt;
v[i].data[j] += dv;
}
});
みたいなコードを書いたら、コンパイルエラーになってしまいました。
見てみると
double __attribute__((vector_size(16 ))) llvm_cbe_tmp__31;
llvm_cbe_tmp__31 = ((double __attribute__((vector_size(16 ))) )/*UNDEF*/{});
というOpenCL Cを吐いていて、double2は{}で初期化できないよって言われてた。
まぁそりゃそうという感じ・・・です。{0}ってすれば解決するのに。
あと、なぜdouble4じゃなくてdouble2なんだろう・・・（preferred vector size見てくれたんですかね？）。
と言った感じでちょっと微妙な結果ですが、先に書いた通り標準ではOpenCL Cをバックエンド（？）には使わないので、仕方ないです。
既定の動作であるSPIRでは動いたので、SPIRはどうなってるのか見てみましょう。
C++AMP→SPIR(LLVM IR)
SPIR（1.2まで。SPIR-Vもあるが別物。詳しくは後述）は、LLVMベースの中間表現で、その実態は「OpenCL CをどのようにLLVM IRと対応させるか」を決めた規格です。
なので、新しい言語というわけではないので、実際にはLLVM IRが分かる人ならわかります。
C++ Advent Calendar勢ならLLVM IRとか英語より読めると聞いたので、LLVM IRの詳しい説明は省略します。
さて、Kalmarでは、なにもしないと、SPIRをバックエンドとして使います（なぜならそっちの方が効率がいいからです）。
先にOpenCL Cで作ってしまった場合は
unset CLAMP_NOSPIR
とすればSPIRで作るようにできます。
知っての通りLLVM IRにはバイナリ表現(bitcode)と、可読な文字列表現(human-readable)の2種類があります。Kalmarはbitcodeを吐くらしく、先の通り普通に
std::cout << source << std::endl;
とすると、「BCなんちゃらかんちゃら」と出てきます。このBCなんちゃらはLLRM IR bitcode formatのマジックナンバーです。
ということで、以下のようにして.bcとしてファイルを吐き出してから、
{
std::ofstream ofs("spir.bc", std::ios::out|std::ios::binary|std::ios::trunc);
for(int i = 0; i < size; i++)
{
ofs << source[i];
}
}
llvm-disで.llに変換すると読めるようになります。
llvm-dis spir.bc
# これで同じフォルダにspir.llが生成される
先のベクトル加算のLLVM IRは
; Function Attrs: nounwind uwtable
define spir_kernel void @ZZ4mainEN3_EC__119__cxxamp_trampolineEPiiiiiiiiS0_iiiiiiiS0_iiiiiii(i32 addrspace(1)*, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32, i32, i32, i32, i32, i32, i32) #5 align 2 {
%25 = tail call i64 @amp_get_global_id(i32 0) #6
%sext = shl i64 %25, 32
%26 = ashr exact i64 %sext, 32
%27 = sext i32 %7 to i64
%28 = sext i32 %5 to i64
%.sum.i.i3.i = add nsw i64 %27, %28
%.sum1.i.i4.i = add nsw i64 %.sum.i.i3.i, %26
%29 = getelementptr inbounds i32 addrspace(1)* %0, i64 %.sum1.i.i4.i
%30 = load i32 addrspace(1)* %29, align 4, !tbaa !8
%31 = sext i32 %15 to i64
%32 = sext i32 %13 to i64
%.sum.i.i1.i = add nsw i64 %31, %32
%.sum1.i.i2.i = add nsw i64 %.sum.i.i1.i, %26
%33 = getelementptr inbounds i32 addrspace(1)* %8, i64 %.sum1.i.i2.i
%34 = load i32 addrspace(1)* %33, align 4, !tbaa !8
%35 = add nsw i32 %34, %30
%36 = sext i32 %23 to i64
%37 = sext i32 %21 to i64
%.sum.i.i.i = add nsw i64 %36, %37
%.sum1.i.i.i = add nsw i64 %.sum.i.i.i, %26
%38 = getelementptr inbounds i32 addrspace(1)* %16, i64 %.sum1.i.i.i
store i32 %35, i32 addrspace(1)* %38, align 4, !tbaa !8
ret void
}
といった感じ。やっぱりよく分からんオフセット計算してるように見えます・・・。
うーん？
OpenCL C→SPIR(LLVM IR)
普通のOpenCL Cも、もちろんSPIRに変換できるので、実際に普通のOpenCLだとどうなっているのか見てみましょう（もはやC++関係ないですが）。
OpenCL CをSPIRに変換するには特殊なclangを使います。
「OpenCL CはLLVMに対応している」みたいなことをちらっと聞いたことがある人もあるかもしれないが、つまりはこれです。
clang本流には取り込まれていないらしいです（SPIR-Vのことを考えるとたぶん取り込まれる予定もないんじゃないですかね）。
使い方はREADMEに書いてある通り。
# LLVM3.2のダウンロード
wget http://www.llvm.org/releases/3.2/llvm-3.2.src.tar.gz
tar xf llvm-3.2.src.tar.gz
# SPIR(1.2)向けclangのダウンロード
cd llvm-3.2.src/tools
wget https://github.com/KhronosGroup/SPIR/archive/spir_12.zip
unzip spir_12.zip
mv spir_12 clang
# ビルド
cd ../../
mkdir build
mkdir install
cd build -DCMAKE_INSTALL_PREFIX=/INTALL/PATH
cmake ..
make
make install
これで/INTALL/PATHの中に一式入ります。あとは
export PATH=/INTALL/PATH:${PATH}
して、試しにclang --versionなどしてみると、バージョンが3.2になっていたらおそらく成功です（どうせならKalmarみたいにバージョン名も変えてほしかったな）。
あとは、
kernel void VectorAdd(
global int dst[],
global const int left[],
global const int right[])
{
const int i = get_global_id(0);
dst[i] = left[i] + right[i];
}
みたいなファイルを適当にocl.clというファイル名で保存して、同じディレクトリで
wget https://raw.githubusercontent.com/KhronosGroup/SPIR-Tools/master/headers/opencl_spir.h
でopencl_spir.hを持ってくる必要があります。そうして
clang -cc1 -emit-llvm-bc -triple spir64-unknown-unknown -cl-spir-compile-options "" -include opencl_spir.h ocl.cl
とすると、LLVM IR bitcodeであるocl.bcが取得できます。bitcodeは読めないので前と同じように
llvm-dis ocl.bc
とすれば、人間が読めるLLM IR(ocl.ll)が取得できました！
覗いてみたらこんな感じになります。
define cc76 void @VectorAdd(i32 addrspace(1)* nocapture %dst, i32 addrspace(1)* nocapture %left, i32 addrspace(1)* nocapture %right) nounwind {
entry:
%call = tail call cc75 i64 @_Z13get_global_idj(i32 0) nounwind readnone
%sext = shl i64 %call, 32
%idxprom = ashr exact i64 %sext, 32
%arrayidx = getelementptr inbounds i32 addrspace(1)* %left, i64 %idxprom
%0 = load i32 addrspace(1)* %arrayidx, align 4, !tbaa !9
%arrayidx2 = getelementptr inbounds i32 addrspace(1)* %right, i64 %idxprom
%1 = load i32 addrspace(1)* %arrayidx2, align 4, !tbaa !9
%add = add nsw i32 %1, %0
%arrayidx4 = getelementptr inbounds i32 addrspace(1)* %dst, i64 %idxprom
store i32 %add, i32 addrspace(1)* %arrayidx4, align 4, !tbaa !9
ret void
}
C++ AMP版と比べて、変なオフセットとか別にないので素直なコードになってますね。
SPIR(LLVM IR)→OpenCLバイナリ
OpenCL 1.2の拡張機能で、このSPIR bitcodeを使ってOpenCLビルドもできます。
もう少し正確に言うと、SPIR(LLVM IR)をOpenCLのプログラムオブジェクトとして使うことができます。
使い方は、
- 通常はclCreateProgramWithSourceでソースコードの文字列を読み込むところを、*.bcを普通にファイル読み込む＆相当するbitcodeを作ってから、clCreateProgramWithBinaryに入力する。
- clBuildProgramでのビルド時のオプションに"–x spir -spir-std=1.2"を追加する
- 他の使い方は同じ
なので、OpenCL Cじゃないフロントエンドを自分で作れば、そのSPIR(LLVM IR)をそのままOpenCLのカーネルとして使うこともできるようになっているということです。
OpenCL 2.1とSPIR-V
以上、ここまでは現行のOpenCL 2.0世代のお話でした。
最後に、OpenCL 2.1の話を書いて終わりにします。
実は、現在は既に次規格であるOpenCL 2.1が、つい半年前に正式リリースされたばかりです。ただし今日現在では、まだOpenCL 2.1を実装した処理系は存在しません・・・。ただ、AMD・Intel・ARM等OpenCLの主力勢（？）は2.1のサポートを表明しているので2.1も引き続き使えるものになる予定です。
OpenCL 2.0から2.1の変更点
大きくは次の4点が変更されました
- 「サブグループ」ができた
- サブグループというのは、ワークアイテムとワークグループの中間にあるようなもの。1つのワークグループを更に分割したものとも言える。
- サブグループ内で同期を取ったりできるようになる。
- OpenCL Cから、OpenCL C++になった
- OpenCL C++は、C++14のサブセット。ホストと共通ヘッダを使えたり、随分と書きやすくなった。
- OpenCL Cも使えないわけではない。
- 後述のSPIR-Vのサポートが必須になったことで、フロントエンドは割と何でも良くなったからC++になったのだと思う。
- デバイス（とホスト）のタイマー時間が取得できるようになった
- 時間計測とかに使える・・・ものだと思う。
- SPIR-Vという中間表現のサポートが必須になった
- 先述の通り既にSPIR 1.2の時点でcl_khr_spirという拡張機能があったが、それが拡張機能ではなくなった。
- SPIR-VからコンパイルするAPIが追加された。
- SPIR 1.2の時はLLVM IRだったが、SPIR-Vは完全にLLVMとは無関係になった。詳しくは後述。
C++er的には、2.1になるとOpenCL CではなくOpenCL C++になって、カーネル側のコードにC++がそのまま使えるようになるという大きな変更があります！
これによって、ホスト側をC++で書いているけどデバイス側の構造体定義と別ファイルに・・・みたいなことをしなくても、ホスト・カーネル両方から同じC++のヘッダをインクルードしたすることができるようになります。
といっても、先述の通りまだ動かす環境が存在しないので、それがどれだけ嬉しいことかはちょっと断言できないところですが、そのうちAMDとかが実装を出した時にはboost勉強会あたりで話をしてみたいと思います。
SPIR-V
SPIR-Vは、SPIR1.2の次の規格です。先述の通り、SPIR-Vでは、LLVMは無関係で完全に独自の命令セットになりました。
現時点（？）では、OpenCL C2.1とOpenGL GLSL4.5、およびVulkanが、SPIR-Vを使う予定です。
つまり、Khronosとしては
- 高レイヤー
- OpenCL (OpenCL C)
- OpenGL (GLSL)
- OpenGL ES (ESSL)
- とか色々
- ↑全部が、同じ中間表現(SPIR-V)で記述される
- 中間レイヤー
- Vulkan
- 低レイヤー
- 各デバイスごとのドライバ（Vulkanを飛ばして、直接SPIR-Vを食べても良いっぽい）
というモデルにしたいっぽいです（※SYCLはどこ行った？）。
Whitepaperのpp. 5-6あたりに例が載っています。
LLVM IRよりもっとアセンブリ言語寄りに見えます。が、所詮IRなのでたぶん対応付けは簡単そうな印象です。
実際に、LLVM IRからSPIR-Vへの翻訳機も開発中だと噂で聞きました（残念ながら明確なソースは見つけられませんでした・・・知ってたらコメントで教えてください）。
まとめ
個人的には、C++AMPでもCUDAでもOpenCLでもいいから、そろそろどれかに統一して欲しいので、少なくとも中間表現がSPIR-Vで統一される未来がくるととても嬉しいです。
明日のC++ Advent Calendarはegtraさんによる可変長テンプレートについての話だそうです。楽しみですね！