マルチGPUの基本的な使用

Question

次のコード（ベクトルの合計）のパフォーマンスを向上させるために、2つのデバイスを使用するにはどうすればよいですか？ 複数のデバイスを同時に使用することはできますか？ はいの場合、異なるデバイスのグローバルメモリ上のベクトルの割り当てをどのように管理できますか？

#include <stdio.h> 
#include <stdlib.h> 
#include <math.h> 
#include <time.h> 
#include <cuda.h> 

#define NB 32 
#define NT 500 
#define N NB*NT 

__global__ void add(double *a, double *b, double *c); 

//=========================================== 
__global__ void add(double *a, double *b, double *c){ 

    int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    while(tid < N){ 
     c[tid] = a[tid] + b[tid]; 
     tid += blockDim.x * gridDim.x; 
    } 

} 

//============================================ 
//BEGIN 
//=========================================== 
int main(void) { 

    double *a, *b, *c; 
    double *dev_a, *dev_b, *dev_c; 

    // allocate the memory on the CPU 
    a=(double *)malloc(N*sizeof(double)); 
    b=(double *)malloc(N*sizeof(double)); 
    c=(double *)malloc(N*sizeof(double)); 

    // allocate the memory on the GPU 
    cudaMalloc((void**)&dev_a, N * sizeof(double)); 
    cudaMalloc((void**)&dev_b, N * sizeof(double)); 
    cudaMalloc((void**)&dev_c, N * sizeof(double)); 

    // fill the arrays 'a' and 'b' on the CPU 
    for (int i=0; i<N; i++) { 
     a[i] = (double)i; 
     b[i] = (double)i*2; 
    } 

    // copy the arrays 'a' and 'b' to the GPU 
    cudaMemcpy(dev_a, a, N * sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b, b, N * sizeof(double), cudaMemcpyHostToDevice); 

    for(int i=0;i<10000;++i) 
     add<<<NB,NT>>>(dev_a, dev_b, dev_c); 

    // copy the array 'c' back from the GPU to the CPU 
    cudaMemcpy(c, dev_c, N * sizeof(double), cudaMemcpyDeviceToHost); 

    // display the results 
    // for (int i=0; i<N; i++) { 
    //  printf("%g + %g = %g\n", a[i], b[i], c[i]); 
    // } 
    printf("\nGPU done\n"); 

    // free the memory allocated on the GPU 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
    // free the memory allocated on the CPU 
    free(a); 
    free(b); 
    free(c); 

    return 0; 
} 

ありがとうございます。 ミシェル

Answer 1

CUDA 4.0がリリースされて以来、お求めのタイプのマルチGPU計算は比較的簡単です。これに先立って、同じホストアプリケーション内で複数のGPUを使用するには、GPUごとに1つのホストスレッドとスレッド間通信システムを使用するマルチスレッドホストアプリケーションを使用する必要があります。

今、あなたのホストコードのメモリ割り当て部分のため、このような何かを行うことが可能である：

double *dev_a[2], *dev_b[2], *dev_c[2]; 
const int Ns[2] = {N/2, N-(N/2)}; 

// allocate the memory on the GPUs 
for(int dev=0; dev<2; dev++) { 
    cudaSetDevice(dev); 
    cudaMalloc((void**)&dev_a[dev], Ns[dev] * sizeof(double)); 
    cudaMalloc((void**)&dev_b[dev], Ns[dev] * sizeof(double)); 
    cudaMalloc((void**)&dev_c[dev], Ns[dev] * sizeof(double)); 
} 

（免責事項：ブラウザで書かれ、コンパイルされたことがない、テストされたことがない、自己責任で使用します）。

ここでの基本的な考え方は、デバイス上で操作を実行するときにcudaSetDeviceを使用してデバイスを選択することです。だから、上記のスニペットでは、私は2つのGPUを仮定し、それぞれの[（N/2）倍の最初のデバイスとN-（N/2）の2番目の]にメモリを割り当てました。

ホストからデバイスへのデータの転送は、のような単純なものでした：

// copy the arrays 'a' and 'b' to the GPUs 
for(int dev=0,pos=0; dev<2; pos+=Ns[dev], dev++) { 
    cudaSetDevice(dev); 
    cudaMemcpy(dev_a[dev], a+pos, Ns[dev] * sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b[dev], b+pos, Ns[dev] * sizeof(double), cudaMemcpyHostToDevice); 
} 

（免責事項：ブラウザで書かれ、コンパイルされたことがない、テストされたことがない、自己責任で使用します）。

あなたのコードのカーネル起動セクションには、その後のようなものになります：

for(int i=0;i<10000;++i) { 
    for(int dev=0; dev<2; dev++) { 
     cudaSetDevice(dev); 
     add<<<NB,NT>>>(dev_a[dev], dev_b[dev], dev_c[dev], Ns[dev]); 
    } 
} 

を（免責事項：ブラウザで書かれ、コンパイルされたことがない、テストされたことがない、自己責任で使用します）。

カーネルの各インスタンスが処理する配列要素の数が異なるため、カーネル呼び出しに追加の引数を追加しました。私はあなたに必要な修正を仕上げるためにあなたに任せます。 でも基本的な考え方は同じです：cudaSetDeviceを使用して特定のGPUを選択し、通常の方法でカーネルを実行します。各カーネルには独自の固有の引数があります。

単純なマルチGPUアプリケーションを作成するには、これらの部分を組み合わせる必要があります。最近のCUDAのバージョンやハードウェアでは、複数のGPUアプリケーション（ユニファイドアドレッシング、ピアツーピア機能など）をサポートするために使用できる他にもたくさんの機能がありますが、これで十分です。また、CUDA SDKには簡単なmuLti-GPUアプリケーションがあり、より多くのアイデアを探すことができます。