#dlang DCompute:ネイティブDによるOpenCLやCUDAのGPGPU【翻訳】

この記事は、 DCompute: GPGPU with Native D for OpenCL and CUDA – The D Blog を自分用に翻訳したものを許可を得て公開するものである。

ソース中にコメントの形で原文を残している。誤字や誤訳などを見つけたら今すぐ Pull requestだ!

Nicholas WilsonはMurdoch Universityの生徒です。 BEng (Hons)/BScのためにIndustrial Computer Systems (Hons)とInstrumentation & Control/ Molecular Biology & Genetics and Biomedical Scienceを学んでいます。彼は電界発光イメージングによるソーラーセルの欠陥のローコストな検出についての卒業論文を書き終わったので、 DComputeの作業をしてそれについてD Blogに書く時間ができました。彼はピアノとアイススケートをたしなみ、number bashing、オートマトン、その他様々なことをDでできるように7年をかけています。

DComputeはGPUやその他アクセラレータを使う計算集約型コードのためにOpenCLやCUDA用のネイティブカーネルをDで書くことをサポートするフレームワークでありコンパイラ拡張です。ハイパフォーマンスDライブラリやアプリケーションの高速な開発を可能にすることを目標にユーザーコードと退屈なものとエラーを起こしがちなコンピュートAPIとの相互作用を自動化するドライバが開発中です。

イントロダクション

昨年5月にOpenCL APIの使用をちょっとひどいものにするDのメタプログラミングを行いながらJohn ColvinのDConf 2016でのプレゼンテーションを見たあと、私は思いました。「Dでカーネルを書けるようになれば、これはOpenCL Cで文字列操作をするよりもっと簡単になるんじゃないか」当時は比較的忙しい学期の終わりが来ており、これはいい冬¹のプロジェクトになると思いました。なんといってもLDC、LLVM D コンパイラはLLVMのSPIR-VとPTXバックエンドへのアクセスができるので、私は「これは難しくない、ただのグルーコードだ」と考えました。私は少しばかりこれにかかる時間を甘く見ており、DCompute(ものに名前をつけるのは難しいことです)の最初のステージ、LDCへの私の変更のメインライニングができたのは2月、つまり8ヶ月後になりました。ちょうどDConfの登録締切の時期で、私はそこで作ったものの進捗について話しました。

LDCとDMDのフロントエンドコードベースと親しくなること以外にも、(関数がカーネルであると伝えるなどのための)特殊なメタデータやOpenCL Cの__global等やCUDAの__global__等を表現するためのアドレススペースを扱い、LDCにその概念を導入する必要があるため、ターゲットとするLLVM SPIR-VとPTXバックエンドも理解する必要がありました。

しかし一度コードに親しくなり上記の差異を整理すれば、OpenCLとCUDAの修飾子をコンパイラが認識できる属性に翻訳し命令を簡単で一貫性のあるインターフェースにラッピングする作業は順調に進みました。

作業が進みほとんどLDCのメインラインにマージする準備が整った頃、私はCIにおける問題にぶつかりました。 Khronosで開発されたSPIR-Vバックエンドはとても古いLLVM 3.6.1をベースにしており、私の思いつきにかかわらず、リリースがありませんでした。そこで私はバックエンドとコンバージョンユーティリティをLLVMのmasterブランチに向けてフォワードポートし自分でリリースしました。これは進行中であり、magic intrinsicsの適切なLLVM命令への変換や、LLVM Trunkへバックエンドをマージするための準備のためのTableGen-drivenアプローチへの移行が進んでいます。うまく行けばすぐに終わるはずです™。

DComputeの現状

現在のDComputeではカーネルをネイティブにDで書くことができ、テンプレートと静的イントロスペクション、 UFCS、スコープガード、レンジとそのアルゴリズム、 CTFEなどの言語機能にアクセスできます。気をつけるべきこととして、ハードウェアとパフォーマンスの観点から関数ポインタ、仮想関数、動的再帰、RTTI、例外、ガベージコレクタの使用はカーネル言語から除外されています。 OpenCL C++と違いカーネル関数をテンプレート化したりオーバーロードやデフォルト値をもたせたりできる点に注目するべきでしょう。 imageとpipeのサポートは進行中です。

コード例

Dでカーネルを書くには、-mdcompute-targets=<targets>をLDCに渡す必要があります。 <targets>はコンマで区切られたビルドしたいターゲットのリストで、たとえばOpenCL 1.2とCUDA compute capability 3.5の場合、それぞれ(そう、一度に全部書くことができます!)ocl-120,cuda-350と書きます。たとえば64ビットモードのときはkernels_ocl120_64.spvのように、そのデバイスのためのすべてのコードを含む1つのファイルが各ターゲットごとに得られます。

Dにおける「ベクトルの加算」カーネルは以下のようになります:

@compute(CompileFor.deviceOnly) module example;
import ldc.dcompute;
import dcompute.std.index;

alias gf = GlobalPointer!float;

@kernel void vadd(gf a, gf b, gf c) 
{
	auto x = GlobalIndex.x;
	a[x] = b[x]+c[x];
}

@compute属性がついたモジュールは各コマンドラインターゲットに向けてコンパイルされ、@kernelは関数をカーネルにし、GlobalPointerはOpenCLの__global修飾子と等価です。

カーネルは関数だけに制限されません。ラムダやテンプレートも動作します:

@kernel void map(alias F)(KernelArgs!F args)
{
    F(args);
}
//以下ホストコード内
AutoBuffer!float x,y,z; // y & z はデータで初期化されています
q.enqueue!(map!((a,b,c) => a=b+c))(x.length)(x, y, z);

KernelArgsはホストの型からデバイスの型 (たとえばポインタへのバッファ。またはこの例の場合AutoIndexed PointersへのAutoBuffers) に変換され、ホストやデバイスの種類の違いからカプセル化されています。

最後の行はカーネルを起動するのに必要な文q.enqueue!kernel(dimensions)(args)で、CUDAのkernel<<<dimensions,queue>>>(args)と近いものです。カーネル起動のためのライブラリは開発中です。

上記の式をホストのコードに変換するすべてのマジックがコンパイラにあるCUDAと違い、q.enqueue!func(sizes)(args)はDComputeのドライバライブラリの静的イントロスペクションによって処理されます。 Dにおいてそのようなことができるたった一つの理由は、コンパイラがシンボルに与える修飾名をシンボルの.mangleofプロパティで取得できるからです。これとDの簡単でパワフルなテンプレートの組み合わせによって、コンピュートAPIを使うことに関する心的オーバーヘッドを著しく減らすことができました。また、そのライブラリへの実装もシンプルになり、したがって同じふるまいをコンパイラにさせるのに比べて早く実装できました。 CUDAユーザーに向けた部分はまだ十分ではありませんが、OpenCLユーザーにとっては新風となることでしょう (OpenCLのベクトル加算ホストコード例のステップ7-11を見てみてください)。

まだDComputeでこのようなことはできませんが、開発は進み、うまく行けばまもなく現実のものになるでしょう。

最初のインスピレーションについてJohn Colvinに、編集についてMike Parkerに、そしてLDCのフォークにDavid Nadlinger、 Kai Nacke、 Martin Kinke、またスペシャルサンクスとしてLDCコードベースの理解の補助と作業のレビューについてJohan Engelenに謝辞を伝えたいです。

DComputeの開発を援助したい(または最新情報を受け取りたい)場合、どうぞ自由にlibmir Gitterに連絡してください。同様に、LLVMへ組み込むためのSPIR-V バックエンドの開発への参加も受け付けています。

南半球 ↩︎