CUDA：从GPU发送数据到GPU

我有两个GPU卡特斯拉C2070（P2P & UAV的支持），我想发送和使用CUDA接收数据。
CUDA：从GPU发送数据到GPU

在GPU A，I具有矩阵：
A11 A12 A13 A14
A21 A22 A23 A24
在GPU B，I有另一个矩阵：
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 BI将获得：
B11 B12 A13 A14
B21 B22 A23 A24

同样，如何从GPU A发送矩阵的第一行至GPU B，等等GPU BI将获得：
A11 A12 A13 A14
B21 B22 B23 B24

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

来源

2014-01-26 Phong

欢迎堆栈溢出！不幸的是，我们不能真正帮助你，直到你展示你到目前为止尝试过的东西。 –

你想看看CUBLAS，你想告诉我们一些代码:) –

我已经添加了一些代码。非常感谢你。 – Phong

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

如果我们搁置您的问题的multiGPU方面一段时间，并且只关注复制倾斜数据（在UVA平台中，如何处理GPU到GPU的传输基本上是您不需要的实现细节知道），只有三个做你想做什么需要的东西：

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

这里是一个具体的例子，它基于您发布的代码，它假定列主要顺序执行两个音高分配之间的数据子集复制。请注意，为了简洁起见，我将大部分寻址机制封装在一个简单的类中，该类可以在主机和设备上使用。分配两个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

来源

2014-01-26 16:42:38 talonmies

感谢您的回复。非常清楚。它正是我想要的。 – Phong

CUDA：从GPU发送数据到GPU

回答

相关问题