CUDA でコーディングを始めたばかりで、GPU を最大限に活用するために、スレッドの実行方法とメモリへのアクセス方法の概念を理解しようとしています。CUDA のベスト プラクティス ガイド、本 CUDA by Example、およびここにあるいくつかの投稿を読みました。また、Mark Harris によるリダクションの例は非常に興味深く有用であることがわかりましたが、すべての情報にもかかわらず、詳細についてかなり混乱しました。
列単位の操作を行う大きな 2D 配列 (N*M) があるとします。配列をブロックに分割して、各ブロックに 32 の倍数のスレッド数が含まれるようにします (すべてのスレッドが複数のワープに収まります)。各ブロックの最初のスレッドは、追加のメモリ (初期配列のコピーですが、独自の次元のサイズのみ) を割り当て、_shared_ を使用してポインターを共有します。同じブロックのすべてのスレッドが同じメモリにアクセスできるように変数。スレッド数は 32 の倍数であるため、1 回の読み取りでアクセスするにはメモリも 32 の倍数である必要があります。ただし、配列の幅が (32*x)+2 列になるように、メモリ ブロックの周囲に余分なパディング、境界線が必要です。境界線は大きな配列を分解することで得られるため、隣接する配列のコピーが一時的に利用できる重複領域があります。
Coleased メモリ アクセス:
ブロックのスレッドがローカル メモリ ブロックにアクセスしていると想像してください。
1 int x = threadIdx.x;
2
3 for (int y = 0; y < height; y++)
4 {
5 double value_centre = array[y*width + x+1]; // remeber we have the border so we need an offset of + 1
6 double value_left = array[y*width + x ]; // hence the left element is at x
7 double value_right = array[y*width + x+2]; // and the right element at x+2
8
9 // .. do something
10 }
さて、私の理解では、避けられないオフセット (+1,+2) があるため、ワープごとおよび割り当てごとに少なくとも 2 つの読み取り (左側の要素を除く) があるか、それとも問題ではないということです。最初のスレッドの後のメモリが完全に整列している限り、どこから読み始めますか? また、そうでない場合は、配列の幅が (32*x)+2 であり、32 バイトで整列されていないため、最初の行の後に各行の配列へのアクセスが整列されていないことに注意してください。ただし、さらにパディングすると、新しい行ごとに問題が解決します。
質問:上記の例では、オフセットなしでアクセスされるのは最初の行のみであり、配列内の左の要素のみが許可されるという私の理解は正しいですか?
ワープで実行されるスレッド:
ワープ内のスレッドは、すべての命令が同じである場合にのみ並列で実行されます (リンクによる)。条件付きステートメント/分岐実行がある場合、その特定のスレッドは、他のスレッドとのワープ内ではなく、単独で実行されます。
たとえば、配列を初期化すると、次のようなことができます
1 int x = threadIdx.x;
2
3 array[x+1] = globalArray[blockIdx.x * blockDim.x + x]; // remember the border and therefore use +1
4
5 if (x == 0 || x == blockDim.x-1) // border
6 {
7 array[x] = DBL_MAX;
8 }
ワープのサイズは 32 で、3 行目まで並列に実行され、その後他のすべてのスレッドで停止し、最初と最後のスレッドのみがさらに実行されて境界線を初期化しますか、それとも最初から他のすべてのスレッドから分離されますか?他のすべてのスレッドが満たさない if ステートメントがありますか?
質問:糸はどのようにして単一のたて糸に集められますか? ワープ内の各スレッドは、同じ命令を共有する必要があります。関数全体でこれを有効にする必要がありますか? これは、スレッド 1 (x=0) には当てはまりません。これは、ボーダーも初期化するため、他とは異なるためです。私の理解では、スレッド 1 は 1 つのワープで実行され、スレッド (2-33 など) は別のワープで実行され、ミスアラインメントのために単一の読み取りでメモリにアクセスしません。もう一方の境界線により、単一の縦糸に糸が通されます。あれは正しいですか？
行ごとにどちらかのメモリを完全に整列させるベストプラクティスは何だろうか (その場合、各ブロックを (32*x-2) スレッドで実行して、境界線のある配列が (32*x-2)+ になるようにする) 2 新しい行ごとに 32 の倍数) または、上記で示したように、ブロックごとに 32 の倍数のスレッドを使用して、アライメントされていないメモリをそのまま使用します。この種の質問は必ずしも単純ではなく、特定のケースに依存することが多いことは承知していますが、特定のことは悪い習慣であり、習慣にすべきではない場合があります.
少し試してみたところ、実行時間の違いはあまりわかりませんでしたが、私の例が単純すぎたのかもしれません。ビジュアル プロファイラーから情報を取得しようとしましたが、提供されるすべての情報を十分に理解できていません。しかし、占有レベルが 17% であるという警告が表示されました。これは非常に低いに違いないと思います。したがって、何か問題があると思います。スレッドがどのように並行して実行されるか、およびメモリ アクセスがどれほど効率的であるかについての情報を見つけることができませんでした。
-編集-
2 つの質問を追加して強調表示しました。1 つはメモリ アクセスに関するもので、もう 1 つはスレッドが単一のワープにどのように収集されるかに関するものです。