如何从另一个 GPU 设备读取数据？答案

【问题标题】：How to read from another GPU device?如何从另一个 GPU 设备读取数据？
【发布时间】：2020-11-24 23:18:55
【问题描述】：

我想在两个 GPU 上同时运行以下简单代码。这里我有一个变量A[i]=[0 1 2 3 4 5 6 7 8 9]，想计算C[i]=A[i+1]+A[i]+A[i-1]。这就是答案：C[i]=[1 3 6 9 7 11 18 21 24 17]。粗体数字是错误的。对于两个设备，设备=1 的C[4] 需要访问设备=2 的A[5]。我怎样才能以最简单的方式做到这一点？

我的专长不是编程，我想使用多 GPU 来求解 PDE 方程。因此，我非常感谢为我当前的问题修改此代码的任何帮助。

谢谢。

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <stdlib.h>
#include<time.h>

__global__ void iKernel(float *A, float *C, const int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) C[i] = A[i-1] + A[i] + A[i+1];
}


int main(int argc, char **argv)
{
    int ngpus;

    printf("> starting %s", argv[0]);

    cudaGetDeviceCount(&ngpus);
    printf(" CUDA-capable devices: %i\n", ngpus);

    ngpus = 2;

    int size = 10;

    int    iSize = size / ngpus;
    size_t iBytes = iSize * sizeof(float);

    printf("> total array size %d M, using %d devices with each device "
        "handling %d M\n", size / 1024 / 1024, ngpus, iSize / 1024 / 1024);


    // allocate device memory
    float **d_A = (float **)malloc(sizeof(float *) * ngpus);
    float **d_C = (float **)malloc(sizeof(float *) * ngpus);

    float **h_A = (float **)malloc(sizeof(float *) * ngpus);
    float **gpuRef = (float **)malloc(sizeof(float *) * ngpus);
    cudaStream_t *stream = (cudaStream_t *)malloc(sizeof(cudaStream_t) * ngpus);

    for (int i = 0; i < ngpus; i++){
        // set current device
        cudaSetDevice(i);

        // allocate device memory
        cudaMalloc((void **)&d_A[i], iBytes);
        cudaMalloc((void **)&d_C[i], iBytes);

        // allocate page locked host memory for asynchronous data transfer
        cudaMallocHost((void **)&h_A[i], iBytes);
        cudaMallocHost((void **)&gpuRef[i], iBytes);

        // create streams for timing and synchronizing
        cudaStreamCreate(&stream[i]);
    }

    dim3 block(512);
    dim3 grid((iSize + block.x - 1) / block.x);

    //h_A[ngpus][index]
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);
        for (int j = 0; j < iSize; j++){
            h_A[i][j] = j + i*iSize;
            printf("%d %d %d %0.8f \n", i,j,iSize, h_A[i][j]);
        }
    }
    // record start time
    double iStart = clock();

    // distributing the workload across multiple devices
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);

        cudaMemcpyAsync(d_A[i], h_A[i], iBytes, cudaMemcpyHostToDevice, stream[i]);

        iKernel << <grid, block, 0, stream[i] >> >(d_A[i], d_C[i], iSize);

        cudaMemcpyAsync(gpuRef[i], d_C[i], iBytes, cudaMemcpyDeviceToHost,
            stream[i]);
    }

    // synchronize streams
    for (int i = 0; i < ngpus; i++){
        cudaSetDevice(i);
        cudaStreamSynchronize(stream[i]);
    }

    for (int i = 0; i < ngpus; i++){
        for (int j = 0; j < iSize; j++){
            printf("%d %d %0.8f \n", i,j,gpuRef[i][j]);
        }
    }
    return EXIT_SUCCESS;
}

【问题讨论】：

在您的内核中，您期望i=0 时的行为是什么？如果i=0 访问A[i-1]，您期望代码应该做什么？需要明确的是，我认为我要问的这个问题与 CUDA 没有太大关系。我并不是说这是您的代码的唯一问题。我建议的另一件事是您演示如何使用正确的 CUDA 错误检查并使用cuda-memcheck 运行您的代码。错误输出可能对您有指导意义，对那些试图帮助您的人很有用。最后，您不妨回顾一下 CUDA simpleMultiGPU 示例代码。
罗伯特，感谢您的快速回复。我知道对于 i=0 和 i=9，内核需要修改。但我的主要问题是位于每个设备边界的 C[4] 和 C[5]。
如答案中所述，您有几个选择。 1. 使用固定分配而不是cudaMalloc。 2. 使用托管内存 3. 在内核启动之间显式复制 GPU 之间的边界区域。 4. 如果系统拓扑支持，将两个 GPU 置于对等关系。然后，一个 GPU 上的内核可以通过中间总线（PCIE 或 NVLink）直接从另一个 GPU 的内存中读取。
This 可能感兴趣。

标签： c++ cuda gpu

【解决方案1】：

您必须将重叠区域上传到两个设备。您不能（轻松）从其他设备读取值，因此您必须根据需要复制并填充至少一些输入值。访问iSize + 2不同的输入值时，iSize显然输入大小不够。

如果这是一个多通道算法，您需要在通道之间显式执行相关区域的副本。

在尝试针对多 GPU 系统时，尝试在纸上正式建模数据依赖关系。

两个 GPU 都可以访问使用 cudaMallocHost 分配的内存，但通常不建议使用该内存类型，因为与设备本地内存相比，PCIe 总线上的性能相当糟糕。还有驱动程序管理的内存，但这也不适合两个 GPU 共享同一个活动工作集。

【讨论】：