cuda：对设备内存使用全局变量答案

【问题标题】：cuda: using global variables for device memorycuda：对设备内存使用全局变量
【发布时间】：2012-02-03 09:49:52
【问题描述】：

我有一个关于如何在 cuda 代码中使用正确变量的问题。我的程序有很多数组，需要在不同的函数中访问，我想避免传递它们并想使用全局变量和 2D mallocpitch 数组，而不是扁平的 1D 数组。所以，我在想这样的事情：

__device__ double * dataPtr ;
__device__ size_t dataPitch;
....
int main()
{
 double * dataPtrLoc; size_t dataPitchLoc;
cudaMallocPitch( (void**) &dataPtrLoc, &dataPitchLoc, width*sizeof(double), height);
cudaMemcpyToSymbol(dataPtr, &dataPtrLoc, sizeof(dataPtrLoc));
cudaMemcpyToSymbol(dataPitch, &dataPitchLoc, sizeof(dataPitchLoc));
...
}

这看起来是获取全球 2D 设备数据的好方法吗？可以给点建议吗？

编辑：我制作了这个程序，它编译并运行良好：

#include <stdio.h>
__device__ int *d_gridPtr;
__device__ size_t d_gridPitch;

__device__ int valij(int ii, int jj)
{
  int* row = (int*)((char*)d_gridPtr + ii * d_gridPitch);
  return (row[jj]);
}

__global__ void printval()
{
  int val0, val1, val2, val3;
  val0= valij(0,0);
  val1= valij(0,1);
  val2= valij(1,0);
  val3= valij(1,1);
  printf("%d %d %d %d \n", val0, val1, val2, val3);
}

int main()
{
  size_t d_gridPitchLoc;
  int * d_gridPtrLoc;  
  cudaMallocPitch((void**)&d_gridPtrLoc, &d_gridPitchLoc, 2 * sizeof(int), 2);
  cudaMemcpyToSymbol(d_gridPtr, & d_gridPtrLoc, sizeof(d_gridPtrLoc));
  cudaMemcpyToSymbol(d_gridPitch, &d_gridPitchLoc, sizeof(float));

  int h_mem[2*2]={0,1,100,4};  
  size_t hostpitch = 2* sizeof(int);
  cudaMemcpy2D(d_gridPtrLoc,d_gridPitchLoc,h_mem,hostpitch,2*sizeof(int),2,cudaMemcpyHostToDevice );

  printval<<<1,1>>> ();
  cudaDeviceReset();  
}

【问题讨论】：

该代码不起作用 - 复制到 dataPtr 是错误的。但是对于这类事情来说，不断的记忆会更有意义。
无论数据是常量还是不相关，您只是将指向全局内存数据的指针存储在常量内存中，而不是数据本身。无论如何，这就是在 Fermi GPU 中实现内核参数的方式。
我的数据不是固定的。我记得读到常量内存不一定是常量，只是缓存的全局内存（虽然不确定）。那么，你仍然认为，持续记忆会起作用吗？
那么，您是否建议使用 __ device __ __ constant __ double *dataPtr 并将内存符号复制到其中？
在 CUDA 中将 2D 数据保持为 2D 可以提高性能，特别是在内存访问中存在空间局部性的情况下（合并等对于编译器在两个维度上都很明显）。此外，如果您的数组适合 constant 设备内存，则声明如下：__constant__ __device__ int dMyContantArray[1024];

标签： variables memory cuda global

【解决方案1】：

如果 warp 或 block 的所有线程同时访问相同的只读全局内存地址（例如数组索引），则考虑将只读全局数据存储在 __constant__ 内存数组中。如果写入数据，则不能使用__constant__。

如果您的数组是只读的并且您的访问模式具有很强的 2D 局部性（在扭曲和/或块内），请考虑改用纹理。

【讨论】：