复制包含指向设备的二维指针的结构答案

【问题标题】：Copying structure containing 2d pointer to device复制包含指向设备的二维指针的结构
【发布时间】：2013-11-26 01:16:49
【问题描述】：

我有一个与复制结构相关的问题，该结构包含从主机指向设备的 2D 指针，我的代码如下

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);

问题

使用此代码，我无法访问结构的 2D 指针元素，但我可以从设备中的该结构访问 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;

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

【问题讨论】：

可以出示matrix的声明吗？

标签： c++ visual-studio-2010 cuda

【解决方案1】：

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

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

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

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

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

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

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

#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());
}

【讨论】：

【解决方案2】：

您的 matrix m 类/结构成员似乎是某种双指针，具体取决于您在主机上的初始化方式：

    m_h[i].m = (float**) malloc(4 * sizeof(float*));

在主机和设备之间复制带有嵌入式指针的结构数组有些复杂。复制双指针指向的数据结构也很复杂。

有关嵌入指针的结构数组，请参阅this posting。

要复制二维数组（双指针，即**），请参阅this posting。我们不使用cudaMallocPitch/cudaMemcpy2D 来完成此操作。（请注意，cudaMemcpy2D 采用单指针 * 参数，您传递给它的是双指针 ** 参数，例如 m_h[i].m）

建议您展平您的数据，而不是上述方法，以便可以通过单指针引用来引用所有数据，而无需嵌入指针。

【讨论】：