https://github.com/ShigekiKarita/d-nvrtc

d-nvrtc (dnv) という CUDA のD言語用ラッパーを作りました。この規模のライブラリを D 言語で書いたのは初めてで、面白いノウハウが溜まったので久々に記事を書いています。 d-nvrtc の売り (目標) は D 言語の持つ強い静的型付けと GC の手軽さで簡単に高速な CUDA プログラミングができることです。現在はこんな感じで動作します。

import dnv;

import std.stdio;
import std.random;
import std.range;

int n = 10;
auto gen = () => new Array!float(generate!(() => uniform(-1f, 1f)).take(n).array()); 
auto a = gen();
auto b = gen();
auto c = new Array!float(n); // CUDA メモリも GC で管理

enum code = Code(
  "saxpy", q{float *A, float *B, float *C, int numElements},
  q{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < numElements) C[i] = A[i] + B[i];
  });
auto saxpy = new TypedKernel!(code);
saxpy(a, b, c, n); // 引数が型が違うとコンパイルエラー

foreach (ai, bi, ci; zip(a.to_cpu(), b.to_cpu(), c.to_cpu())) {
  assert(ai + bi == ci);
}

属性を持つクラスの冗長性

CUDAカーネルクラス Kernel には 2 つの属性があります。1 つは NVRTC をラップしたクラス Compiler の型検査による属性 (静的・動的・無し) です。もう 1 つはCUDAカーネルを呼び出すクラス Launcher による属性 (e.g., 要素数によるもの・動的共有メモリを使うもの・ユーザ定義) といった具合です。

ここで問題となるのは、例に上げただけでも $3 \times 3 = 9$ 通りのパターンを実装する必要があります。継承なりで重複をさけることはできるものの、 冗長であり命名の面倒さが発生します。

class StaticCompiler : Compiler { ... }
class DynamicSharedLauncher : Launcher { ... }

// ひどい名前のクラス
class StaticDynamicSharedKernel : Kernel {
  Compiler compiler = new StaticCompiler;
  Launcher launcher = new SharedLauncher;
}

// 何通りもの似たようなクラス
...

Policy とは

Policy による冗長性の解決

Policy 間のブラックリスト

D言語における依存性注入

Scala などでは依存性注入 (DI) というパターンがでてくる。 これは D 言語の interface とよくにた trait (D言語のそれとは違う) の多重継承によって 柔軟な設計と疎結合性を実現している。 先に述べたブラックリスト的な機能を実現する方法は