5

私はマルチ GPU プログラミングの初心者で、マルチ GPU コンピューティングについていくつか質問があります。たとえば、内積の例を見てみましょう。2 つの大きな配列 A[N] と B[N] を作成する CPU スレッドを実行しています。これらの配列のサイズのため、内積の計算を 2 つの GPU に分割する必要があります。どちらも Tesla M2050 (計算能力 2.0) です。問題は、CPU スレッドによって制御される do ループ内でこれらの内積を数回計算する必要があることです。各内積には、前の内積の結果が必要です。2 つの異なる GPU を別々に制御する 2 つの異なるスレッドの作成について読んだことがありますが (cuda の例で説明されているように)、それらの間でデータを同期して交換する方法についての手がかりがありませんでした。別の代替手段はありますか？あらゆる種類のヘルプ/例をいただければ幸いです。よろしくお願いします!

4

2 に答える 2

7

CUDA 4.0 より前は、マルチ GPU プログラミングにはマルチスレッド CPU プログラミングが必要でした。これは、特にスレッドまたは GPU 間で同期および/または通信する必要がある場合に困難になる可能性があります。また、すべての並列処理が GPU コードにある場合、複数の CPU スレッドを使用しても、GPU の機能を超えてパフォーマンスを向上させることなく、ソフトウェアの複雑さが増す可能性があります。

そのため、CUDA 4.0 以降では、シングルスレッドのホスト プログラムから複数の GPU を簡単にプログラムできます。 これについて、私が昨年発表したスライドをいくつか示します。

複数の GPU のプログラミングは、次のように簡単に行うことができます。

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    kernel<<<blocks, threads>>>(args);
}

内積の特定の例についてはthrust::inner_product、出発点として使用できます。私はプロトタイピングのためにそれをします。ただし、帯域幅のボトルネックに関する私のコメントを最後に参照してください。

ドット積を複数回実行する外側のループについて十分な詳細を提供していないため、それについては何もしようとしませんでした。

// assume the deviceIDs of the two 2050s are dev0 and dev1.
// assume that the whole vector for the dot product is on the host in h_data
// assume that n is the number of elements in h_vecA and h_vecB.

int numDevs = 0;
cudaGetNumDevices(&numDevs);
...
float result = 0.f;
for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1);
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1);
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f);
}

(n が numDevs の偶数倍でない場合、上記のインデックス付けが正しくないことは認めますが、読者の演習としてそれを修正しておきます。:)

これは簡単で、素晴らしいスタートです。最初に機能させてから、最適化します。

それが機能したら、デバイスで行っているのがドット積だけである場合、帯域幅が制限されていることがわかります-ほとんどがPCI-eによって、また、デバイス間の同時実行性が得られないことがわかりますinner_product。結果を返すために読み返すため..したがって、cudaMemcpyAsync（device_vectorコンストラクターは cudaMemcpy を使用します)。しかし、より簡単でおそらくより効率的なアプローチは、「ゼロ コピー」を使用することです。つまり、ホスト メモリに直接アクセスします (上記のマルチ GPU プログラミング プレゼンテーションでも説明されています)。各値を 1 回読み取って合計に加算するだけなので (共有メモリのコピーで並列再利用が行われます)、ホストからデバイスにコピーしてから読み取るよりも、ホストから直接読み取った方がよいでしょう。カーネルのデバイスメモリから。また、最大の同時実行性を確保するために、各 GPU でカーネルを非同期で起動する必要があります。

次のようなことができます。

int bytes = sizeof(float) * n;
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable);
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable);
// ... then fill your input arrays h_vecA and h_vecB


for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    cudaEventCreate(event[d]));
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0);
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0);
    cudaHostGetDevicePointer(&dresults[d], results, 0);
}

...

for (int d = 0; d < numDevs; d++) {
    cudaSetDevice(d);
    int first = d * (n/d);
    int last   = (d+1)*(n/d)-1;
    my_inner_product<<<grid, block>>>(&dresults[d], 
                                      vecA+first, 
                                      vecA+last, 
                                      vecB+first, 0.f);
    cudaEventRecord(event[d], 0);
}

// wait for all devices
float total = 0.0f;
for (int d = 0; d < devs; d++) {
    cudaEventSynchronize(event[d]);
    total += results[numDevs];
}
于 2012-03-05T05:25:04.220 に答える
1

複数のスレッドを作成するには、OpenMPまたはpthreadのいずれかを使用できます。あなたが話していることを行うには、2つのスレッド（omp parallel section、またはpthread_create）を作成して起動し、それぞれに計算の一部を実行させ、その中間結果を別々のプロセス全体の変数に格納する必要があるようです。 （グローバル変数はプロセスのスレッド間で自動的に共有されるため、元のスレッドは2つの生成されたスレッドによって行われた変更を確認できます）。元のスレッドが他のスレッドの完了を待機するようにするには、（グローバルバリアまたはスレッド結合操作を使用して）同期し、2つの生成されたスレッドが完了した後に結果を元のスレッドに結合します（

MPIまたはフォークを使用することもできます。その場合、通信はネットワークプログラミングと同様の方法で行うことができます...パイプ/ソケットまたは送信と受信（ブロック）を介した通信と同期。

于 2012-03-04T14:16:46.267 に答える