KOTET'S PERSONAL BLOG

この記事は、 DCompute: Running D on the GPU – 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ドライバはユーザーコードと退屈なものとエラーを起こしがちなAPIの間の相互作用を自動化します。

イントロダクション

これはDComputeの2つ目の記事です。 前回の記事¹ では、DComputeの発展と小さな例を扱いました。 カーネルをビルドすることには成功しましたが、それを既存のフレームワークや自力で簡単に実行する方法はありませんでした。 しかし今はそんなことはありません。 v0.1.0では、DComputeはOpenCLとCUDAのネイティブラッパーであると同時に、CUDAよりも簡単なカーネルディスパッチを実現するものでもあります。

カーネルを実行するためにはCUDAかOpenCL、どちらか適切なコンピュートAPIにそれを渡さなければなりません。 どちらのAPIも同じことができますが、パフォーマンスのためにそれぞれのAPIを別々に扱うのに十分なほど違ってもいます。 しかしそれでも2つの間には適度な一貫性をもったインターフェースを作ることができるほどには重なる部分があります。 しかしこれらのAPIのCバインディングは非常にローレベルで、それを使うことは非常に退屈かつエラーを引き起こしやすいです(yay void*)。 退屈さとエラー性に加えて、冗長な情報の指定法が組み合わさってきます。 幸運にもこれはDであり、イントロスペクションとコード生成によって冗長性を大きく減らすことができます。

このドライバはC APIをラップし、クリーンで一貫性があり、簡単に使えるインターフェースを提供します。 ドキュメントはいまのところちょっと不足気味ですが、ソースコードの殆どは素直に書かれています(C APIに慣れているなら、関数が使われているところから読んでみるといいでしょう)。 ときおり、まともなAPIを実現するためにちょっとしたマジックがあったりします。

野獣を飼いならす

OpenCLのclGet*Info関数はvoid*に隠れたクラスのプロパティにアクセスするためのものです。 典型的にはこのようになります:

enum CL_FOO_REFERENCE_COUNT = 0x1234;
cl_foo* foo = ...; 
cl_int refCount;
clGetFooInfo(foo, CL_FOO_REFERENCE_COUNT, refCount.sizeof, &refCount,null);

And that’s not even one for which you have to call, to figure out how much memory you need to allocate, then call again with the allocated buffer (and $DEITY help you if you want to get a cl_program’s binaries).

Dを使うとこのようになります:

struct Foo
{
    void* raw;
    static struct Info
    {
        @(0x1234) int referenceCount;
        ...
    }
    mixin(generateGetInfo!(Info, clGetFooInfo));
}

Foo foo  = ...;
int refCount = foo.referenceCount;

すべてのマジックはFoo.Infoの各メンバに対しプロパティを生成し、スケーラビリティとドキュメントを与えるgenerateGetInfoにあります。

CUDAにも同じようなやり方で公開されるプロパティがありますが、(OpenCLとは違い)これは必須ではないため、開発は先送りされています。

OpenCLと(酷さはそんなに変わらない)CUDAの両方のC APIを扱うとき、型安全の欠如とvoid*への&オペレータの使用のために、カーネルの起動は苦労するものです。 DComputeではその呪文がOpenCLの場合は

Event e = q.enqueue!(saxpy)([N])(b_res, alpha, b_x, b_y, N);

このように(1D with N work items)、そしてCUDAの場合

q.enqueue!(saxpy)([N, 1, 1], [1 ,1 ,1])(b_res, alpha, b_x, b_y, N);

このように(saxpy<<<N,1,0,q>>>(b_res,alpha,b_x,b_y, N);と等価)シンプルになります。

この例でqはキュー、Nはバッファーの長さ、(b_res, b_x & b_y)とsaxpy(single-precision a x plus y)はカーネルです。 OpenCLと CUDA のエンキュー関数のマジックに関する完全な例はここで見つかるでしょう。

DComputeの未来

DComputeは機能的ですが、まだ十分ではありません。 ドライバーは洗練とユーザーテストが必要で、継続的インテグレーションもセットアップしなければなりません。 異なるコンピュートAPIをまとめるドライバも業界のクラスプラットフォーム標準よりもっとクロスプラットフォームにできるので、作業中です。

SPIR-VをSPIRへ変換できれば1.xと2.0のCL実装でcl_khr_spirをターゲットにする能力が得られ、D カーネルコードを実行できるデバイスの数は急激に増えます(他のカーネルのOpenCLドライバーを使わない理由はありません)。

コンパイラ側のものとしては、OpenCLのimageやCUDAのtextureやsurface操作をLDCがサポートすれば書けるカーネルの幅が広がります。 私は現在LLVM IRからSPIR-Vを生成するためのKhronos’s SPIR-V LLVMのフォワードポートフォークを管理しています。 私はLLVM trunkへそれをマージするために取り組みを調整するためにIWOCLを使おうと考えており、そうするといくらかSPIR-Vの古いところに対処するためのハックが必要なくなります。

プロジェクトでDComputeを使う

DComputeを使うつもりなら、NVPTX (CUDA用)とSPIRV (OpenCL 2.1+用)ターゲットのどちらかまたは両方が有効化されたLLVMでLDCをビルドし、"dcompute": "~>0.1.0"をdub.jsonに追記する必要があります。 1.4以降のLDCのリリースではNVPTXは有効化されています。 OpenCLをターゲットとする場合、私のLLVMフォークでLDCをビルドしてください。