CUDA：将数据从 GPU 发送到 GPU

Question

我有两张 GPU 卡 Tesla C2070（P2P 和 UAV 支持）我想发送和接收 使用 CUDA 的数据。

• 在 GPU A 中，我有一个矩阵：
  a11 a12 a13 a14
  a21 a22 a23 a24

• 在 GPU B 中，我有另一个矩阵：
  b11 b12 b13 b14
  b21 b22 b23 b24

我只能像下面的代码那样发送连续的元素：

int main(void) 
{

    float *d_a, *d_b;
    int N = 4;
    int M = 2;
    size_t pitch;   

    cudaSetDevice(0);   
    cudaMallocPitch(&d_a, &pitch, sizeof(float)*N, M);
    cudaDeviceEnablePeerAccess(1, 0);

    cudaSetDevice(1);       
    cudaMallocPitch(&d_b, &pitch, sizeof(float)*N, M);
    cudaDeviceEnablePeerAccess(0, 0);

    //Initialization for d_a
    //Initialization for d_b

    //Copy M*N/2 element from d_a to d_b, starting from d_a[1]
    cudaMemcpy(&d_b[1], &d_a[1], M*N/2*sizeof(float), cudaMemcpyDefault);

    //Print result d_b          
}

如何将矩阵的最后两列从GPU A直接发送到GPU B，所以在GPU B上我会得到：
b11 b12 a13 a14
b21 b22 a23 a24

同理，如何将矩阵的第一行从GPU A发送到GPU B，所以在GPU B上我会得到：
a11 a12 a13 a14
b21 b22 b23 b24

如果我有如下一维数组：a1 a2 a3 a4 a5 a6 a7 a8.....
如何从 GPU A 发送元素 1、4、7、...（每 3 个元素）以替换 GPU B 上的相同元素？

Answer 1

您需要查看的 API 调用是 cudaMemcpy2D。这允许相当直接地复制所有或部分音调数据，并且是cudaMallocPitch 的自然对应物。

如果我们暂时搁置您的问题的多 GPU 方面，而只专注于复制倾斜数据（在 UVA 平台中，如何处理 GPU 到 GPU 的传输基本上是您不需要知道的实现细节about)，只需要做三件事即可：

1. 使用指针算法计算源内存和目标内存的起始地址
2. 请记住，源内存和目标内存的间距始终是恒定的（由cudaMallocPitch 返回）。请注意，您应该为分配的每个指针保留一个音高。无法保证 API 将为相同大小的两个不同分配返回相同的间距，如果分配不在同一设备上，则尤其如此
3. 请记住，您需要以字节为单位计算任何传输的宽度，并且数字宽度始终是计数，而不是字节值。

这是一个基于您发布的代码的具体示例，该示例在假设列主要顺序的情况下执行两个倾斜分配之间的数据子集复制。请注意，为简洁起见，我将大部分寻址机制封装在一个简单的类中，该类可以在主机和设备上使用。分配了两个 5x10 间距数组，并将一个 3x3 子数组从一个复制到另一个。我已经使用内核printf 来显示复制动作：

#include <cstdio>

struct mat
{
    int m, n;
    size_t pitch;
    char *ptr;

    __device__ __host__
    mat(int _m, int _n, size_t _pitch, char *_ptr) : m(_m), n(_n), pitch(_pitch), ptr(_ptr) {};

    __device__ __host__ float * getptr(int i=0, int j=0) {
        float * col = (float*)(ptr + j*pitch);
        return col + i;
    };

    __device__ __host__ float& operator() (int i, int j) { 
        return *getptr(i,j);
    };

    __device__ __host__
    void print() {
        for(int i=0; i<m; i++) {
            for(int j=0; j<n; j++) {
                printf("%4.f ", (*this)(i,j));
            }
            printf("\n");
        }
    };
};

__global__ void printmat(struct mat x) { x.print(); }

int main(void) 
{

    const int M = 5, N = 10;
    const size_t hostpitch = M * sizeof(float);

    float *a = new float[M*N], *b = new float[M*N];
    mat A(M, N, hostpitch, (char *)(a));
    mat B(M, N, hostpitch, (char *)(b));
    for(int v=0, j=0; j<N; j++) {
        for(int i=0; i<M; i++) {
            A(i,j) = (float)v; B(i,j) = (float)(100+v++);
        }
    }

    char *d_a, *d_b;
    size_t pitch_a, pitch_b;   
    cudaMallocPitch((void **)&d_a, &pitch_a, sizeof(float)*M, N);
    cudaMallocPitch((void **)&d_b, &pitch_b, sizeof(float)*M, N);
    mat Ad(M, N, pitch_a, d_a); mat Bd(M, N, pitch_b, d_b);

    cudaMemcpy2D(Ad.getptr(), Ad.pitch, A.getptr(), A.pitch, 
            A.pitch, A.n, cudaMemcpyHostToDevice);
    printmat<<<1,1>>>(Ad);

    cudaMemcpy2D(Bd.getptr(), Bd.pitch, B.getptr(), B.pitch, 
            B.pitch, B.n, cudaMemcpyHostToDevice);
    printmat<<<1,1>>>(Bd);

    int ci = 3, cj = 3;
    cudaMemcpy2D(Ad.getptr(1,1), Ad.pitch, Bd.getptr(1,1), Bd.pitch, 
            ci*sizeof(float), cj, cudaMemcpyDeviceToDevice);
    printmat<<<1,1>>>(Ad); cudaDeviceSynchronize();

    return 0;
}

这是做什么的：

>nvcc -m32 -Xptxas="-v" -arch=sm_21 pitched.cu
pitched.cu
tmpxft_00001348_00000000-5_pitched.cudafe1.gpu
tmpxft_00001348_00000000-10_pitched.cudafe2.gpu
pitched.cu
ptxas : info : 0 bytes gmem, 8 bytes cmem[2]
ptxas : info : Compiling entry function '_Z8printmat3mat' for 'sm_21'
ptxas : info : Function properties for _Z8printmat3mat
    8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas : info : Used 23 registers, 48 bytes cmem[0]
tmpxft_00001348_00000000-5_pitched.cudafe1.cpp
tmpxft_00001348_00000000-15_pitched.ii

>cuda-memcheck a.exe
========= CUDA-MEMCHECK
   0    5   10   15   20   25   30   35   40   45
   1    6   11   16   21   26   31   36   41   46
   2    7   12   17   22   27   32   37   42   47
   3    8   13   18   23   28   33   38   43   48
   4    9   14   19   24   29   34   39   44   49
 100  105  110  115  120  125  130  135  140  145
 101  106  111  116  121  126  131  136  141  146
 102  107  112  117  122  127  132  137  142  147
 103  108  113  118  123  128  133  138  143  148
 104  109  114  119  124  129  134  139  144  149
   0    5   10   15   20   25   30   35   40   45
   1  106  111  116   21   26   31   36   41   46
   2  107  112  117   22   27   32   37   42   47
   3  108  113  118   23   28   33   38   43   48
   4    9   14   19   24   29   34   39   44   49
========= ERROR SUMMARY: 0 errors