GPU Compute プログラミングガイド¶

KL Example
GPU Operator の呼び出し
型のサポート
Memory マネジメント
Parameter Passing

この文書では KLで GPU Compute を行うための ―現在NVIDIAデバイスのみで動作します― 実装の詳細についてと、それらをどのようにKLから使用するか、概要を示します。以下に書かれる全ては GPU Compute ガイドで示したように CUDA GPU サポート機能を正しく有効にしたものとします。

GPUコンピューティングは比較的に新しい機能であり、APIの追加や進化がこれからもやってくる可能性があります。私たちは、エンドユーザからのフィードバックや提案に常に耳を傾けます。

KL Example ¶

以下に、GPUコンピューティングを使用するKLコードを示し、さらなる議論の土台とします。ただし以下に示すコードはGPU上で実行可能ではありますが、とくにGPUに適したアルゴリズムの例というわけではありません。

/*
** Example: GPU Compute
*/

operator myComputeKernel<<<index>>>(Scalar input[], io Scalar output[])
{
  output[index] = input[index] + Scalar(log(index));
}

operator entry()
{
  Boolean runOnGPU = true;

  Scalar input[];
  input.resize(1024);
  for (Integer i=0; i<1024; i++)
    input[i] = i*i;
  input.convertToSVM();

  Scalar output[];
  output.convertToGPU();
  output.resize(1024);
  myComputeKernel<<<1024@runOnGPU>>>(input, output);
  output.convertToCPU();
  report(output);
}

GPU Operator の呼び出し ¶

GPUコンピューティング機能は、KLの標準的な並列実行記法並列実行文（PEX文）のオペレーターに @真偽値 true を後置することで有効となります。 @パラメータを指定しないか、@パラメータ直後に false と評価されるものを置くと、オペレータはGPUではなくCPUでの呼び出しとなります。

CUDAをサポートしない機器、もしくはCUDAを無効としたか、読み込みに失敗した環境で、並列実行記法に @true パラメータを与え実行すると、GPUでの実行とはならず、CPU上でGPU環境をシミュレートしてのオペレータ呼び出しとなる。これは特にテストに有用である。この擬似環境での実行では、operatorを @true で呼び出したにもかかわらず、コードがGPUで実行されていないことのユーザーへの告知としては以下のメッセージが表示されます：

[FABRIC:MT] Falling back to running GPU operator on CPU

型のサポート ¶

KLの型全てがGPU上で使用可能なわけではありません。とくに object、 interface そして MapReduce は現在サポートされません。GPU上ではサポートされない型を走らせようとすると、KLはエラーを出力します。ただし、Dict や String はGPUオペレータに対しパラメータとして受け渡す事はできないが、オペレータ内部で使用することはできます。これらの型についての制約は、使用実態のフィードバックや、妥当性によって今後のバージョンで撤廃されえます。

GPUで実行するオペレータにとって、最適な型は、可変長配列（Variable Array）です。GPUオペレータを並列実行する際、大規模な入力データ（または複数のデータ群）を、直接可変長配列としてマップし、使用する。これらの可変長配列には、より複雑な型 ― Vec3 や Xfo を含むことが多いでしょう。 input 、 output に配列を用いた例については KL Example を参照してください。

Memory マネジメント ¶

GPUコンピューティングでは、4つのメモリー利用モードがある。ユーザーは必要に応じ、最適なものを選択できる。このうち3つのモードは：

CPU: 標準的なCPUメモリ。malloc() より heap 領域に確保される。CPU上に作成される全ての型のメモリ確保先のディフォルトです。
GPU: 標準的なGPUに確保されるメモリ。GPUのみからアクセスが可能であり、CPUからアクセスを行うと、 guarded モードではエラーが返り unguarded モードではクラッシュする。
SVM: Shared Virtual Memory を意味する。CUDAでは、CUDA Managed Memory と呼ばれるものを指す。この種のメモリは、CPUからもGPUからもアクセスが可能である。NVIDIAのドライバによりデータは透過的に転送される。このため、メモリが実際にはどこに存在すべきか知るひつようのないような新参利用者にとって、一番簡単なメモリモードとなる。ただし、経験豊富なユーザーにとっては、よりきめ細やかな制御を必要とする可能性がある。
GLBuffer: OpenGLバッファーとして読み込まれ、GPUオペレータでの処理に費やされるメモリ。この種のメモリーについては、以下の gpupg-gl-binding セクションで詳述します。

Memory 確保先の変更 ¶

KLの全ての型はメモリの確保先を変更するため、いくつかの関数（method）をサポートしている。多くの型 ― shallow type を含む ― ではこれらの変更メソッドはその型自体が、メモリを所持している訳ではないので実際には操作が実行されない。しかし、可変長配列の場合では、これらの変更のためのメソッドは、その配列の要素が格納されている場所を、変更する。関連する３つのメッドは次の通り：

myVar.convertToCPU()
myVar.convertToGPU()
myVar.convertToSVM()

これらのメソッドは、メモリ確保先をあるモードから他の3つのモードに変換します。

GL Binding ¶

可変長配列にはさらに convertToGLBuffer() メソッドによる変換メソッドがあります。GPUコンピューティングだけではなく、さらには描画にも使用可能な GLバッファとして格納されます。このメソッドを呼ぶ際には、有効なGLコンテキストにバインドする必要があります。そうでない場合、以下の例外が発生します：

Exception: glewInit() call failed; is there a valid GL context bound?

GL binding の短いですが利用例：

/*
** Example: GL binding in KL
*/

require FabricOGL;

operator entry()
{
  Scalar input[];
  input.resize(1024);

  // the convertToGLBuffer() call returns the GL buffer ID
  Integer bufferId = input.convertToGLBuffer(GL_ARRAY_BUFFER, GL_DYNAMIC_DRAW);

  // the GL buffer ID can also be retrieved later via the getBufferId() method
  bufferId = input.getBufferId();

  report('GL buffer ID is: '+bufferId);
  myGPUOperator<<<1024@true>>>(input);
}

convertToGLBuffer() に渡すパラメータは、通常 glBufferData() に渡す target, usage パラメータと同じで、同じ値です。

Array Memory の確保先の決定 ¶

KLでは、可変長配列の要素が現在どこに存在するのかについて、 getElementsMemType() メソッドを使い、明らかにすることができます。ここでは簡単な使用例で説明します：

/*
** Example: Variable Array Memory Types in KL
*/

function printMemoryType(Vec3 myArray[])
{
  if (myArray.getElementsMemType() == Fabric_MemType_CPU)
    report('Array elements are in CPU memory.');
  else if (myArray.getElementsMemType() == Fabric_MemType_GPU)
    report('Array elements are in GPU memory.');
  else if (myArray.getElementsMemType() == Fabric_MemType_SVM)
    report('Array elements are in Shared Virtual memory.');
  else if (myArray.getElementsMemType() == Fabric_MemType_GLBuffer)
    report('Array elements are stored in a GL buffer.');
}

Memory and Resizing Arrays ¶

GPU処理をおこなう際、配列の resize() メソッドを使用するとメモリが存在するGPUデバイス上の確保されたメモリをリサイズします。GL buffer にバインドされた配列では、 resize() メソッドを使用することはできません。エラーが投げられます。

CPUからアクセス可能なメモリ―つまり CPU, SVM メモリでは、新規にサイズ変更された配列は、その配列の要素をその要素のディフォルトコンストラクタに使用し初期化します。（例： Vec3() コンストラクタによる初期化）GPUメモリでは、新規にサイズ変更された配列では、ゼロに初期化します。これは、CPUとGPUでの配列の動作の違いを示すものとして、極めて重要な差異となります。

/*
** Example: The resize() method
*/

operator entry()
{
  Vec3 a[];

  // array is resized to 1024 elements in CPU memory, all elements are
  // initialized with the Vec3() constructor
  a.resize(1024);

  // all 1024 elements are removed from CPU memory and transferred to GPU
  a.convertToGPU();

  // the array in GPU memory is resized to 2048 elements, the first 1024
  // hold their previous values while the new elements are initialized with 0s
  a.resize(2048);
}

The copyTo() Array Method ¶

copyTo() メソッドは、メモリ確保先に関係なく配列間でデータを転送に使用することができます。このメソッドは、転送先配列を転送元配列と同じサイズになるようリサイズします。転送先配列のメモリ確保先は不変なママです。

/*
** Example: The copyTo() method
*/

operator entry()
{
  Vec3 a[];
  a.resize(1024);
  // ... fill 'a' with data ...

  Vec3 b[];
  b.resize(128);
  // ... fill 'b' with data ...
  b.convertToGPU();

  // resizes 'b' to 1024 elements (on GPU) and copies the values from 'a'
  a.copyTo(b);
}

Parameter Passing ¶

Parameters to a KL GPU operator can be of any supported type and as with any normal KL operator they can be passed as in (the default) or io parameters. Shallow type parameters (such as an Integer or a struct) will have their values copied to and from the GPU before and after each parallel operator invocation; passing these shallow types as in parameters has very little overhead, whereas passing them as io or out requires an allocation on the GPU and is much slower. Variable Arrays on the other hand will only have a pointer to their values passed into the operator. The backing memory for the Variable Array elements must live in memory accessible to the target device.

guarded モードでは KL は可変長配列パラメータの確保先が、メモリ空間からアクセス可能であるか確認します。もし不可能であれば、そのことをユーザに例外の送信を通じてしらせます。例えば、以下のKLコードのでは：

/*
** Example: GPU Compute - Memory
*/

operator entry()
{
  Integer input[];
  input.resize(1024);
  Integer output[];
  output.resize(1024);
  myComputeKernel<<<1024@true>>>(input, output);
}

これはエラーメッセージなります：

Error: input: data not available in GPU memory

付け加えると、 input や output を convertToGPU() や convertToSVM() メソッドを呼びこのエラーを解決することができます KL Example を参照してください。

unguarded モードでは、デバイスや、CPUメモリを正しく使わないとクラッシュを引き起こします。