c++ - 将包含 2d 指针的结构复制到设备

Question

我有一个问题是关于从主机向设备复制包含二维指针的结构，我的代码如下

struct mymatrix
    {
        matrix m;
        int x;
    };
size_t pitch;

mymatrix m_h[5];
for(int i=0; i<5;i++){
    m_h[i].m = (float**) malloc(4 * sizeof(float*));  
       for (int idx = 0; idx < 4; ++idx)
           {
               m_h[i].m[idx] = (float*)malloc(4 * sizeof(float));
           }
       }
mymatrix *m_hh = (mymatrix*)malloc(5*sizeof(mymatrix));
memcpy(m_hh,m_h,5*sizeof(mymatrix));

for(int i=0 ; i<5 ;i++) 
{
     cudaMallocPitch((void**)&(m_hh[i].m),&pitch,4*sizeof(float),4);
     cudaMemcpy2D(m_hh[i].m, pitch, m_h[i].m, 4*sizeof(float), 4*sizeof(float),4,cudaMemcpyHostToDevice);
}
mymatrix *m_d;
cudaMalloc((void**)&m_d,5*sizeof(mymatrix));
cudaMemcpy(m_d,m_hh,5*sizeof(mymatrix),cudaMemcpyHostToDevice);
distance_calculation_begins<<<1,16>>>(m_d,pitch);

问题

使用此代码我无法访问结构的二维指针元素，但我可以从设备中的该结构访问 x。例如例如，如果我初始化

，我会收到带有指针 mymatrix* m 的 m_d

m[0].m[0][0] = 5;

并打印这个值，例如

cuPrintf("The value is %f",m[0].m[0][0]);

在设备中，我没有得到任何输出。意味着我无法使用 2D 指针，但如果我尝试访问

 m[0].x = 5; 

然后我就可以打印这个了。我认为我的初始化是正确的，但我无法找出问题所在。任何人的帮助将不胜感激。

Answer 1

除了@RobertCrovella 在您的代码中指出的问题外，还请注意:

• 您只能使用将 m_h 复制到 m_hh 的 memcpy 获得结构的浅表拷贝。
• 您假设 pitch 在对 cudaMemcpy2D() 的所有调用中都是相同的(您覆盖 pitch 并在最后仅使用最新的拷贝)。我认为目前这可能是安全的假设，但将来可能会改变。
• 您正在使用 cudaMemcpyHostToDevice() 和 cudaMemcpyHostToDevice 复制到主机上的 m_hh，而不是设备上。

使用许多小缓冲区和指针表在 CUDA 中效率不高。小的分配和解除分配最终可能会花费大量时间。此外，使用指针表会导致额外的内存事务，因为必须先从内存中检索指针，然后才能将它们用作索引的基础。因此，如果您考虑这样的构造:

a[10][20][30] = 3

必须首先从内存中检索 a[10] 处的指针，导致你的扭曲被搁置很长时间(在费米上最多大约 600 个周期)。然后，同样的事情发生在第二个指针上，又增加了 600 个周期。此外，这些请求不太可能合并，从而导致更多的内存事务。

正如 Robert 所提到的，解决方案是扁平化您的内存结构。我为此提供了一个示例，您可以将其用作程序的基础。如您所见，代码总体上要简单得多。确实变得有点复杂的部分是索引计算。此外，此方法假定您的矩阵都具有相同的大小。

我也添加了错误检查。如果您在代码中添加了错误检查，那么您至少会发现一些错误而无需任何额外的努力。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

typedef float* mymatrix;

const int n_matrixes(5);
const int w(4);
const int h(4);


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void test(mymatrix m_d, size_t pitch_floats)
{
  // Print the value at [2][3][4].
  printf("%f ", m_d[3 + (2 * h + 4) * pitch_floats]);
}


int main()
{
  mymatrix m_h;
  gpuErrchk(cudaMallocHost(&m_h, n_matrixes * w * sizeof(float) * h));
  // Set the value at [2][3][4].
  m_h[2 * (w * h) + 3 + 4 * w] = 5.0f;

  // Create a device copy of the matrix.
  mymatrix m_d;
  size_t pitch;
  gpuErrchk(cudaMallocPitch((void**)&m_d, &pitch, w * sizeof(float), n_matrixes * h));
  gpuErrchk(cudaMemcpy2D(m_d, pitch, m_h, w * sizeof(float), w * sizeof(float), n_matrixes * h, cudaMemcpyHostToDevice));

  test<<<1,1>>>(m_d, pitch / sizeof(float));

  gpuErrchk(cudaPeekAtLastError());
  gpuErrchk(cudaDeviceSynchronize());
}