多GPU的基本用法多GPU的基本用法(multi-GPU basic usage)

我怎么能以改善例如下面的代码（矢量的总和）的性能使用两个设备？是否有可能使用更多的设备“在同一时间”？如果是的话，我怎么能管理载体的分配在不同的设备的全局内存？

#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;
}

先感谢您。米歇尔

由于CUDA 4.0发布了，你是想询问类型的多GPU计算是比较容易的。在此之前，你必须需要以使用相同的主机应用程序中的GPU复式使用，每个GPU一个主机线程和某种线程间通信系统的多线程主机应用程序。

现在可以为您的主机代码的内存分配部分做这样的事情：

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设备之间进行选择，当你的设备上预成型操作。因此，在上面的代码片段，我假定两个GPU并在每个分配的内存[（N / 2）的两倍的第一设备上和在所述第二N-（N / 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简单的多GPU的应用程序，你可以看看更多的想法。