0_Simple__MultiGPU

使用多台 GPU 进行计算

▶ 源代码。使用不同的流来控制不同 GPU 上的运算任务。

#include <stdio.h>
#include <timer.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_functions.h>
#include <helper_cuda.h>
#include "simpleMultiGPU.h"

const int MAX_GPU_COUNT = 32;
const int DATA_N        = 1048576 * 32;

__global__ static void reduceKernel(float *d_Result, float *d_Input, int N)
{
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int threadN = gridDim.x * blockDim.x;
    float sum = 0;
    
    for (int pos = tid; pos < N; pos += threadN)
        sum += d_Input[pos];

    d_Result[tid] = sum;
}

int main(int argc, char **argv)
{
    printf("
	Start.
");

    const int BLOCK_N = 32, THREAD_N = 256;
    const int ACCUM_N = BLOCK_N * THREAD_N;
    int i, j, GPU_N;
    float sumGPU;
    TGPUplan plan[MAX_GPU_COUNT];
    
    cudaGetDeviceCount(&GPU_N);
    GPU_N = MIN(GPU_N, MAX_GPU_COUNT);
    printf("
	Device count: %i
", GPU_N);

    // 准备计算数据
    for (i = 0; i < GPU_N; i++)
        plan[i].dataN = DATA_N / GPU_N;

    // 计算数据量与设备数量没有对齐的部分
    for (i = 0; i < DATA_N % GPU_N; i++)
        plan[i].dataN++;

    // 申请内存，初始化 h_data
    for (i = 0; i < GPU_N; i++)
    {
        cudaSetDevice(i);
        cudaStreamCreate(&plan[i].stream);
        cudaMalloc((void **)&plan[i].d_data, plan[i].dataN * sizeof(float));
        cudaMalloc((void **)&plan[i].d_sum, ACCUM_N * sizeof(float));
        cudaMallocHost((void **)&plan[i].h_sum_from_device, ACCUM_N * sizeof(float));
        cudaMallocHost((void **)&plan[i].h_data, plan[i].dataN * sizeof(float));

        for (j = 0; j < plan[i].dataN; j++)
            plan[i].h_data[j] = (float)rand() / (float)RAND_MAX;
    }

    StartTimer();// 计时

    // 调用各 GPU 进行计算，plan[i].d_data -> plan[i].d_sum -> plan[i].h_sum_from_device
    for (i = 0; i < GPU_N; i++)
    {
        cudaSetDevice(i);
        cudaMemcpyAsync(plan[i].d_data, plan[i].h_data, plan[i].dataN * sizeof(float), cudaMemcpyHostToDevice, plan[i].stream);
        reduceKernel<<<BLOCK_N, THREAD_N, 0, plan[i].stream>>>(plan[i].d_sum, plan[i].d_data, plan[i].dataN);
        cudaMemcpyAsync(plan[i].h_sum_from_device, plan[i].d_sum, ACCUM_N *sizeof(float), cudaMemcpyDeviceToHost, plan[i].stream);
    }

    // 处理 GPU 计算结果，plan[i].h_sum_from_device -> plan[i].h_sum -> sumGPU
    for (i = 0; i < GPU_N; i++)
    {
        cudaSetDevice(i);
        cudaStreamSynchronize(plan[i].stream);

        for (j = 0, plan[i].h_sum = 0.0f; j < ACCUM_N; j++)
            plan[i].h_sum += plan[i].h_sum_from_device[j];
    }
    for (i = 0, sumGPU = 0.0f; i < GPU_N; i++)// CPU 最后规约
        sumGPU += plan[i].h_sum;
    printf("
	GPU Processing time: %f (ms)
", GetTimer());

    // 使用 CPU 计算，plan[i].h_data -> sumCPU
    double sumCPU = 0;
    for (i = 0; i < GPU_N; i++)
    {
        for (j = 0; j < plan[i].dataN; j++)
            sumCPU += plan[i].h_data[j];
    }

    // 检查结果
    double diff = fabs(sumCPU - sumGPU) / fabs(sumCPU);
    printf("
	GPU sum: %f
	CPU sum: %f
", sumGPU, sumCPU);
    printf("
	Relative difference: %E, %s
", diff, (diff < 1e-5) ? "Passed" : "Failed");
    
    //回收工作
    for (i = 0; i < GPU_N; i++)
    {
        cudaSetDevice(i);
        cudaFreeHost(plan[i].h_data);
        cudaFreeHost(plan[i].h_sum_from_device);
        cudaFree(plan[i].d_sum);
        cudaFree(plan[i].d_data);
        cudaStreamDestroy(plan[i].stream);
    }

    getchar();
    return 0;
}

▶ 输出结果

    Start.

    Device count: 1

    GPU Processing time: 13.726471 (ms)

    GPU sum: 16779778.000000
    CPU sum: 16779776.312309
    Relative difference: 1.005789E-07, Passed

▶ 涨姿势

● 在使用不同的设备执行相关函数（包括 cudaFree 等主机函数）时要注意，使用函数 cudaSetDevice() 来切换设备。