私は現在CUDAを学んでおり、アルゴリズムは入力データに基づいて重い計算を行う必要があります。これらの計算は、最大 1024 ラウンドのループで行われます。カーネルごとに少数のスレッド (< 100´000) がある限り、すべて正常に動作しますが、より多くのスレッドを使用したい場合は、完了するまでに時間がかかりすぎるため、カーネルが Windows によって中断されます。
私の解決策は、重い計算をいくつかのカーネル呼び出しに分割することです。
- 入力データを準備し、最初のx ラウンド (ループ展開) を計算するメインカーネル。これは、入力ごとに 1 回だけ呼び出されます。
- 次のx ラウンド (ループ展開) を実行する作業カーネル。これは、必要なすべてのラウンドを計算するために必要なだけ呼び出されます。
各カーネル呼び出し (1 つのmain、多くのwork ) の間に、次の呼び出しで使用される 16 + length バイトのデータを保存する必要があります (length は入力の長さであり、メイン呼び出しごとに固定されます)。メインカーネルは最初にこれらのバイトを書き込み、作業カーネルはそれらを読み取り、次の計算を実行し、新しい結果で元のデータを書き込みます。デバイス上のこれらのデータのみが必要です。ホスト アクセスは必要ありません。これにはどの種類のメモリを使用する必要がありますか? カーネル呼び出し中に保持される唯一の書き込み可能なメモリであるため、少なくともグローバル メモリである必要がありますね。しかし、その後、何ですか？正しい記憶 (および最高のパフォーマンス) で作業を進めるにはどうすればよいか、アドバイスをいただけますか?
「疑似コード」では、次のようになります。
prepare memory to hold threads * (16 + length) bytes
for length = 1 to x step 1
call mainKernel
rounds = 1024 - rounds_done_in_main
for rounds to 0 step rounds_done_in_work
call workKernel
end for
end for
cleanup memory
--------
template <unsigned char length> __global__ mainKernel() {
unsigned char input[length];
unsigned char output[16];
const int tid = ...;
devPrepareInput<length>(input);
calc round 1: doSomething<length>(output, input)
calc round 2: doSomething<length>(output, output + input) // '+' == append
write data to memory based on tid // data == output + input
}
template <unsigned char length, remaining rounds> __global__ workKernel() {
unsigned char *input;
unsigned char *output;
const int tid = ...;
read data from memory based on tid
ouput = data
input = data+16
if rounds >= 1
calc round x : doSomething<length>(output, output + input)
if rounds >= 2
calc round x+1: doSomething<length>(output, output + input) // '+' == append
if rounds == x // x is the number of rounds in the last work call
do final steps on output
else
write ouput + input to memory based on tid (for next call)
}