0

我有一个与复制结构相关的问题，该结构包含从主机指向设备的 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从 device.xml 中的该结构访问。例如，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;

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

4

2 回答 2

1

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

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

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

有关具有嵌入式指针的结构数组，请参阅此贴文。

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

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

于 2013-11-13T14:33:47.460 回答
1

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

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

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

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

必须首先从内存中检索 a[10] 处的指针，这会导致您的 warp 被搁置很长时间（在 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());
}
于 2013-11-13T17:07:33.677 回答