11

最近、CUDA で文字列比較ジョブを行っていますが、探している正確な文字列が見つかったときに __global__ 関数がどのように値を返すことができるのだろうかと思います。

つまり、大きな大きな文字列プールの中から特定の文字列を同時に見つけるには、大量のスレッドを含む __global__ 関数が必要です。正確な文字列がキャッチされたら、__global__ 関数がすべてのスレッドを停止して戻ることができることを願っていますメイン関数に戻り、「彼がやった」と教えてくれます!

私はCUDA Cを使用しています.どうすればこれを達成できますか?

4

3 に答える 3

21

CUDA (または NVIDIA GPU) では、1 つのスレッドが実行中のすべてのスレッドの実行を中断する方法はありません。結果が見つかったらすぐにカーネルを終了することはできません。今日では不可能です。

ただし、1 つのスレッドが結果を見つけたら、できるだけ早くすべてのスレッドを終了させることができます。これを行う方法のモデルを次に示します。

__global___ void kernel(volatile bool *found, ...) 
{
    while (!(*found) && workLeftToDo()) {

       bool iFoundIt = do_some_work(...); // see notes below

       if (iFoundIt) *found = true;
    }
}

これに関するいくつかのメモ。

  1. の使用に注意してくださいvolatile。これは重要。
  2. カーネルを起動する前に、必ず初期化found(デバイス ポインタでなければなりません)を行ってください!false
  3. 別のスレッドが更新されても、スレッドはすぐには終了しませんfound。次に while ループの先頭に戻ったときにのみ終了します。
  4. どのように実装するかがdo_some_work重要です。作業が多すぎる (または変動しすぎる) 場合、結果が見つかった後に終了するまでの遅延が長くなります (または変動します)。作業が少なすぎると、スレッドはfound有用な作業を行うのではなく、ほとんどの時間をチェックに費やします。
  5. do_some_workまた、タスクの割り当て (つまり、インデックスの計算/インクリメント) も担当します。その方法は、問題によって異なります。
  6. 起動するブロックの数が、現在の GPU のカーネルの最大占有率よりも​​はるかに多く、実行中のスレッド ブロックの最初の「ウェーブ」で一致が見つからない場合、このカーネル (およびその下のカーネル) はデッドロックする可能性があります。 . 最初のウェーブで一致が見つかった場合、後のブロックは の後にのみ実行されますfound == true。つまり、起動してすぐに終了します。解決策は、同時に常駐できる数のブロックのみを起動し (別名「最大起動」)、それに応じてタスクの割り当てを更新することです。
  7. タスクの数が比較的少ない場合は、 を に置き換えて、whileタスクifの数をカバーするのに十分な数のスレッドを実行できます。その場合、デッドロックの可能性はありません (ただし、前のポイントの最初の部分が適用されます)。
  8. workLeftToDo()問題固有ですが、やるべき作業が残っていない場合は false を返すため、一致が見つからない場合にデッドロックすることはありません。

現在、特に L1 キャッシュのない古いアーキテクチャでは、上記の結果として過剰なパーティション キャンピング (すべてのスレッドが同じメモリをバンキングする) が発生する可能性があります。したがって、ブロックごとの共有ステータスを使用して、もう少し複雑なバージョンを作成することをお勧めします。

__global___ void kernel(volatile bool *found, ...) 
{
    volatile __shared__ bool someoneFoundIt;

    // initialize shared status
    if (threadIdx.x == 0) someoneFoundIt = *found;
    __syncthreads();

    while(!someoneFoundIt && workLeftToDo()) {

       bool iFoundIt = do_some_work(...); 

       // if I found it, tell everyone they can exit
       if (iFoundIt) { someoneFoundIt = true; *found = true; }

       // if someone in another block found it, tell 
       // everyone in my block they can exit
       if (threadIdx.x == 0 && *found) someoneFoundIt = true;

       __syncthreads();
    }
}

このように、ブロックごとに 1 つのスレッドがグローバル変数をポーリングし、一致したスレッドのみがグローバル変数に書き込むので、グローバル メモリ トラフィックが最小限に抑えられます。

余談ですが、__global__ 関数は無効です。これは、数千のスレッドから単一の CPU スレッドに値を返す方法を定義するのが難しいためです。ユーザーが自分の目的に合ったデバイスまたはゼロコピー メモリ内の戻り配列を考案するのは簡単ですが、一般的なメカニズムを作成するのは困難です。

免責事項: ブラウザーで記述されたコード、未テスト、未検証。

于 2012-09-20T04:22:56.240 に答える
5

冒険好きなら、カーネルの実行を停止するための別のアプローチは、単に実行することです

// (write result to memory here)
__threadfence();
asm("trap;");

答えが見つかれば。

これにはメモリのポーリングは必要ありませんが、カーネルがエラー状態で終了するという点で、Mark Harris が提案したソリューションよりも劣っています。これにより、実際のエラーが隠される可能性があり (そのため、実行が成功したこととエラーが明確にわかるように結果を書き出すようにしてください)、ドライバーがこれを例外として処理するため、他の問題が発生したり、全体的なパフォーマンスが低下したりする可能性があります。

安全でシンプルな解決策を探している場合は、代わりに Mark Harris の提案に従ってください。

于 2012-09-20T11:09:10.883 に答える
0

グローバル関数には、あなたが思っているほど多くのスレッドが含まれているわけではありません。これは、スレッド モデルを指定するパラメーターを渡すことによって呼び出される、デバイス上で実行される単なるカーネル (関数) です。CUDA が採用するモデルは、2D グリッド モデルと、グリッド上の各ブロック内の 3D スレッド モデルです。

あなたが抱えている問題のタイプでは、各ブロックで 1D のスレッドがオンになっている 1D グリッド以外のものを使用する必要はありません。これは、文字列プールが他の問題のように 2D に分割する意味がないためです (行列乗算など)。

文字列プール内の 100 個の文字列の簡単な例について説明します。これらの文字列はすべて、順次ではなく並列化された方法でチェックする必要があります。

//main
//Should cudamalloc and cudacopy to device up before this code
dim3 dimGrid(10, 1); // 1D grid with 10 blocks
dim3 dimBlocks(10, 1); //1D Blocks with 10 threads 
fun<<<dimGrid, dimBlocks>>>(, Height)
//cudaMemCpy answerIdx back to integer on host

//kernel (Not positive on these types as my CUDA is very rusty
__global__ void fun(char *strings[], char *stringToMatch, int *answerIdx)
{
    int idx = blockIdx.x * 10 + threadIdx.x;

    //Obviously use whatever function you've been using for string comparison
    //I'm just using == for example's sake
    if(strings[idx] == stringToMatch)
    { 
       *answerIdx = idx
    }
}

これは明らかに最も効率的ではなく、パラメータを渡してCUDAでメモリを操作する正確な方法ではない可能性が最も高いですが、ワークロードを分割することの要点を理解し、「グローバル」関数が多くの異なる上で実行されることを願っていますすべてのコアに停止を指示することはできません。私がよく知らない方法があるかもしれませんが、ワークロードをデバイスに分割するだけで得られる速度の向上 (もちろん賢明な方法で) は、すでに大幅なパフォーマンスの向上をもたらします。スレッド モデルを理解するには、Nvidia の CUDA サイトのドキュメントを読むことを強くお勧めします。それらは非常に役立ち、最適なパフォーマンスを得るためにグリッドとブロックを設定する最良の方法を教えてくれます.

于 2012-09-20T04:02:37.920 に答える