首页 文章

多GPU基本用法

提问于
浏览
22

如何使用两个设备来改善例如以下代码的性能(向量之和)?是否可以“同时”使用更多设备?如果是,我如何管理不同设备的全局内存上的向量分配?

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

先感谢您 . 米歇尔

1 回答

  • 36

    自从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中还有一个简单的muLti-GPU应用程序,您可以查看更多想法 .

相关问题