丢失在 CUDA 设备指针中

Posted

技术标签:

【中文标题】丢失在 CUDA 设备指针中【英文标题】：Lost in CUDA device pointers 【发布时间】：2014-03-26 16:24:31 【问题描述】：

作为我论文工作的一部分，我正在从事一个 CUDA 项目（修改别人的代码、添加功能等）。作为 CUDA 的新手，这对我来说是一个真正的挑战。我正在使用 计算能力 1.3 卡，4 x Tesla C1060。遗憾的是，我遇到了平台的一些限制。

我需要将几个新结构传递给设备，我认为这些结构已正确复制。但是，当尝试在我的内核调用中将指针传递给设备上的结构时，我达到了 256 字节的限制（如question 中所述）。

我的代码是这样的：

// main.cu
static void RunGPU(HostThreadState *hstate)

  SimState *HostMem = &(hstate->host_sim_state);
  SimState DeviceMem;

  TetrahedronStructGPU *h_root = &(hstate->root);
  TetrahedronStructGPU *d_root;
  TriangleFacesGPU *h_faces = &(hstate->faces);
  TriangleFacesGPU *d_faces;

  GPUThreadStates tstates;

  unsigned int n_threads = hstate->n_tblks * NUM_THREADS_PER_BLOCK;
  unsigned int n_tetras  = hstate->n_tetras; // 9600
  unsigned int n_faces   = hstate->n_faces;  // 38400

  InitGPUStates(HostMem, h_root, h_faces, &DeviceMem, &tstates, hstate->sim, 
                d_root, d_faces, n_threads, n_tetras, n_faces );
  cudaThreadSynchronize();

  ...

  kernel<<<dimGrid, dimBlock, k_smem_sz>>>(DeviceMem, tstates, /*OK, these 2*/
                                           d_root, d_faces);
                           // Limit of 256 bytes adding d_root and/or d_faces
  cudaThreadSynchronize();

  ...

InitGPUStates 函数在另一个源文件中：

// kernel.cu
int InitGPUStates(SimState* HostMem, TetrahedronStructGPU* h_root,
                  TriangleFacesGPU* h_faces,
                  SimState* DeviceMem, GPUThreadStates *tstates,
                  SimulationStruct* sim, 
                  TetrahedronStructGPU* d_root, TriangleFacesGPU* d_faces,
                  int n_threads, int n_tetras, int n_faces)

  unsigned int size;

  // Allocate and copy RootTetrahedron (d_root) on device
  size = n_tetras * sizeof(TetrahedronStructGPU); // Too big
  checkCudaErrors(cudaMalloc((void**)&d_root, size));
  checkCudaErrors(cudaMemcpy(d_root, h_root, size, cudaMemcpyHostToDevice));

  // Allocate and copy Faces (d_faces) on device
  size = n_faces * sizeof(TriangleFacesGPU); // Too big
  checkCudaErrors(cudaMalloc((void**)&d_faces, size));
  checkCudaErrors(cudaMemcpy(d_faces, h_faces, size, cudaMemcpyHostToDevice));     

  ...

我知道我只需要传递指向设备内存位置的指针。如何获取设备中的地址？这种指针传递是否正确完成？

两个新结构是：

// header.h
typedef struct 
  int idx;
  int vertices[4];
  float Nx, Ny, Nz, d;
 TriangleFacesGPU;

typedef struct 
  int idx, region;
  int vertices[4], faces[4], adjTetras[4];
  float n, mua, mus, g;
 TetrahedronStructGPU;

// other structures
typedef struct 
  BOOLEAN *is_active;
  BOOLEAN *dead;
  BOOLEAN *FstBackReflectionFlag;
  int *NextTetrahedron;
  UINT32 *NumForwardScatters;
  UINT32 *NumBackwardScatters;
  UINT32 *NumBackwardsSpecularReflections;
  UINT32 *NumBiases;
  UINT32 *p_layer;
  GFLOAT *p_x, *p_y, *p_z;
  GFLOAT *p_ux, *p_uy, *p_uz;
  GFLOAT *p_w;
  GFLOAT *Rspecular;
  GFLOAT *LocationFstBias;
  GFLOAT *OpticalPath;
  GFLOAT *MaxDepth;
  GFLOAT *MaxLikelihoodRatioIncrease;
  GFLOAT *LikelihoodRatioIncreaseFstBias;
  GFLOAT *LikelihoodRatio;
  GFLOAT *LikelihoodRatioAfterFstBias;
  GFLOAT *s, *sleft;
  TetrahedronStructGPU *tetrahedron;
  TriangleFacesGPU *faces;
 GPUThreadStates;

typedef struct 
  UINT32 *n_p_left;
  UINT64 *x;
  UINT32 *a;
  UINT64 *Rd_ra;
  UINT64 *A_rz;
  UINT64 *Tt_ra;
 SimState;

kernel的定义是

__global__ void kernel(SimState d_state, GPUThreadStates tstates,
                       TetrahedronStructGPU *d_root,
                       TriangleFacesGPU *d_faces);

我将努力将SimState d_state 更改为指针传递SimState *d_state。以及GPUThreadStates tstates 到GPUThreadStates *tstates。

【问题讨论】：

kernel的声明是什么？例如，您似乎将tstates 按值 传递给kernel。如果sizeof(GPUThreadStates) 很大，您可以通过指针而不是值传递该结构来释放一些喘息空间。问题是，d_root 和 d_faces 是 已经 指针。因此，如果仅添加这两个指针就超出了参数空间，则需要缩小要传递的其他内容的大小，例如 DeviceMem (sizeof(SimState)) 和 tstates (sizeof(GPUThreadStates)) .这也会影响引用这些实体的内核代码。 @RobertCrovella 你是对的。我不确定我是否正确传递了指针。内核定义 __global__ void MCMLKernel(SimState d_state, GPUThreadStates tstates, TetrahedronStructGPU *d_root, TriangleFacesGPU *d_faces) 以及 d_state 和 tstates 都是按值传递的，不是吗？ 是的，它们似乎是，尽管您实际上没有显示GPUThreadStates 和SimState 的定义。如果它们的大小很大，阻止您添加d_root（指针）和d_faces（指针），那么您将不得不关注这些。 @RobertCrovella 再次感谢。我正在研究这些，GPUThreadStates 和 SimState 也很大。我在上面添加了这些定义。 @RobertCrovella 我发布了我所做的修改，作为更好格式化的答案。我有错误code=11(cudaErrorInvalidValue) "cudaMalloc((void**)&amp;DeviceMem-&gt;n_photons_left, size)" 。我将衷心感谢您的帮助！谢谢！ 【参考方案1】：

您似乎还没有初始化 DeviceMem 结构，它应该保存稍后应该使用 cudaMalloc 初始化的指针。

你应该这样做：

SimState* DeviceMem;

cudaMalloc(&DeviceMem, sizeof(SimState))

也是（或为该指针分配内存的任何其他方式）。

【讨论】：

我自己分配SimState结构中的每个元素都没有关系？我没有包含该代码，只有前三个元素 DeviceMem-&gt;n_p_left、DeviceMem-&gt;a 和 DeviceMem-&gt;x。但是，我也对以下元素执行相同操作（Rd_ra、A_rz 和 Tt_ra）。 由于我在 kernel 调用中通过引用（而不是以前的值）传递了 DeviceMem，因此我可能需要分配所提到的完整结构。但是，不能百分百确定。问题是 GPU 上的调试能力非常有限（也许我还没有找到正确的方法）。我发现很难在设备上跟踪我的数据。 好吧，严格来说这不是“通过引用传递”。您正在通过指针传递。如果您通过引用传递，则您的对象将已经存在，并且您一直在使用对它的引用作为函数参数，例如void f(int&amp; i) - 这里i，函数参数，作为引用传递。但这只能在 C++ 中实现，而不是在纯 C 中。 @kronos @t_carn 在question。你如何让你的内核调用doThings？您将指针传递给整个结构 Matrix 还是单独传递给 Matrix.elements？ 据我所知，他们通过常规的malloc分配Matrix结构的数组，然后将这些结构的数据成员传递给CUDA函数。【参考方案2】：

终于解决了 256 字节的问题。但是，真的还是迷路了

我修改后的代码是这样的：

// main.cu
static void RunGPU(HostThreadState *hstate)

  SimState *HostMem = &(hstate->host_sim_state);

  // new pointers to pass
  SimState *DeviceMem = (SimState*)malloc(sizeof(SimState));
  GPUThreadStates *tstates = (GPUThreadStates*)malloc(sizeof(GPUThreadStates));

  TetrahedronStructGPU *h_root = hstate->root; //root, pointer in HostThreadState
  TetrahedronStructGPU *d_root;
  TriangleFacesGPU *h_faces = hstate->faces; //faces, pointer in HostThreadState
  TriangleFacesGPU *d_faces;

  unsigned int n_threads = hstate->n_tblks * NUM_THREADS_PER_BLOCK;
  unsigned int n_tetras  = hstate->n_tetras; // 9600
  unsigned int n_faces   = hstate->n_faces;  // 38400

  InitGPUStates(HostMem, h_root, h_faces, DeviceMem, tstates, hstate->sim, 
                d_root, d_faces, n_threads, n_tetras, n_faces );
  cudaThreadSynchronize();

  ...

  kernel<<<dimGrid, dimBlock, k_smem_sz>>>(DeviceMem, tstates,
                                           d_root, d_faces);
                                         // No limit reached!
  cudaThreadSynchronize();

  ...

在InitGPUStates 函数中的更改如下。特别注意DeviceMem的副本（我尝试了很多形式都没有成功）。某些表格（带括号，例如cudaMalloc((void **)&amp;(*DeviceMem).n_p_left, size)）不会给我任何错误。我假设没有错误意味着没有数据复制到设备。在当前形式中，错误是code=11(cudaErrorInvalidValue) "cudaMalloc((void**)&amp;DeviceMem-&gt;n_photons_left, size)"。

// kernel.cu
int InitGPUStates(SimState* HostMem, TetrahedronStructGPU* h_root,
                  TriangleFacesGPU* h_faces,
                  SimState* DeviceMem, GPUThreadStates *tstates,
                  SimulationStruct* sim, 
                  TetrahedronStructGPU* d_root, TriangleFacesGPU* d_faces,
                  int n_threads, int n_tetras, int n_faces)

  unsigned int size;

  // Allocate and copy RootTetrahedron (d_root) on device
  size = n_tetras * sizeof(TetrahedronStructGPU); // Too big
  checkCudaErrors(cudaMalloc((void**)&d_root, size));
  checkCudaErrors(cudaMemcpy(d_root, h_root, size, cudaMemcpyHostToDevice));

  // Allocate and copy Faces (d_faces) on device
  size = n_faces * sizeof(TriangleFacesGPU); // Too big
  checkCudaErrors(cudaMalloc((void**)&d_faces, size));
  checkCudaErrors(cudaMemcpy(d_faces, h_faces, size, cudaMemcpyHostToDevice));     

  // HELP NEEDED MAINLY FROM HERE REGARDING POINTER VALUE COPY!
  checkCudaErrors( cudaMalloc((void**)&DeviceMem, sizeof(SimState) ); //Needed?

  size = sizeof(UINT32);
  checkCudaErrors( cudaMalloc(&DeviceMem->n_p_left, size) );
  checkCudaErrors( cudaMemcpy(DeviceMem->n_p_left,
                   HostMem->n_p_left, size, cudaMemcpyHostToDevice) );

  size = n_threads * sizeof(UINT32);
  checkCudaErrors( cudaMalloc(&DeviceMem->a, size) );
  checkCudaErrors( cudaMemcpy(DeviceMem->a, HostMem->a, size,
                                      cudaMemcpyHostToDevice) );
  size = n_threads * sizeof(UINT64);
  checkCudaErrors( cudaMalloc(&DeviceMem->x, size) );
  checkCudaErrors( cudaMemcpy(DeviceMem->x, HostMem->x, size,
                                      cudaMemcpyHostToDevice) );
  ...

我知道我只需要传递指向设备内存位置的指针。如何获取设备中的地址？这种指针传递是否正确完成？

两个新结构是：

// header.h
typedef struct 
  int idx;
  int vertices[4];
  float Nx, Ny, Nz, d;
 TriangleFacesGPU;

typedef struct 
  int idx, region;
  int vertices[4], faces[4], adjTetras[4];
  float n, mua, mus, g;
 TetrahedronStructGPU;

// other structures
typedef struct 
  BOOLEAN *is_active;
  BOOLEAN *dead;
  BOOLEAN *FstBackReflectionFlag;
  int *NextTetrahedron;
  UINT32 *NumForwardScatters;
  UINT32 *NumBackwardScatters;
  UINT32 *NumBackwardsSpecularReflections;
  UINT32 *NumBiases;
  UINT32 *p_layer;
  GFLOAT *p_x, *p_y, *p_z;
  GFLOAT *p_ux, *p_uy, *p_uz;
  GFLOAT *p_w;
  GFLOAT *Rspecular;
  GFLOAT *LocationFstBias;
  GFLOAT *OpticalPath;
  GFLOAT *MaxDepth;
  GFLOAT *MaxLikelihoodRatioIncrease;
  GFLOAT *LikelihoodRatioIncreaseFstBias;
  GFLOAT *LikelihoodRatio;
  GFLOAT *LikelihoodRatioAfterFstBias;
  GFLOAT *s, *sleft;
  TetrahedronStructGPU *tetrahedron;
  TriangleFacesGPU *faces;
 GPUThreadStates;

typedef struct 
  UINT32 *n_p_left;
  UINT64 *x;
  UINT32 *a;
  UINT64 *Rd_ra;
  UINT64 *A_rz;
  UINT64 *Tt_ra;
 SimState;

kernel的定义改为：

__global__ void kernel(SimState *d_state, GPUThreadStates *tstates,
                       TetrahedronStructGPU *d_root,
                       TriangleFacesGPU *d_faces);

【讨论】：

不应该是cudaMalloc((void**)&amp;(DeviceMem-&gt;n_p_left), size)之类的吗？ @Ashalynd 我试过了。仍然得到相同的错误代码=11(cudaErrorInvalidValue) 已经有一个类似的问题：***.com/questions/19404965/…（而且看起来你不需要将指针转换为 void**） 啊，还有别的事。您的 InitGPUStates 函数应该接收指向您的 DeviceMem 的指针，然后将该指针提供给 cudaMalloc。你忘了初始化 DeviceMem 结构:) @Ashalynd 通过初始化 DeviceMem 你的意思是 cudaMalloc 在你的答案中？我知道有一个关于 void** 的问题，首先要更好地重新阅读答案以更好地理解need of void**。谢谢！我也在检查question。

以上是关于丢失在 CUDA 设备指针中的主要内容，如果未能解决你的问题，请参考以下文章

14.linux中设备的访问

OpenHarmony驱动框架HDF中设备管理服务构建过程详解

iOS 14 中设备之间的小部件行为不一致

从 CUDA 中的主机访问设备上的类成员数组指针

设备管理中设备名称聚类分析

TTY体系中设备节点的差别