有没有办法动态分配常量内存? CUDA

Is there any way to dynamically allocate constant memory? CUDA

我对将数组复制到常量内存感到困惑。

根据 programming guide 至少有一种方法可以分配常量内存并使用它来存储值数组。这称为静态内存分配:

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));

再次根据programming guide我们可以使用:

__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

看起来好像使用了动态常量内存分配,但我不确定。而且这里也没有使用限定符 __constant__

所以这里有一些问题:

  1. 这个指针是否存储在常量内存中?
  2. (通过这个指针)分配的内存是否也存储在常量内存中?
  3. 这个指针是常量吗?并且不允许使用设备或主机功能更改该指针。但是是否禁止更改数组的值?如果允许改变数组的值,那么是否意味着不使用常量内存来存储这些值?

开发人员可以在文件范围内声明最多 64K 的常量内存。在 SM 1.0 中,工具链使用的常量内存(例如,保存编译时常量)与开发人员可用的常量内存是分开的,而且我认为这从那以后就没有改变。驱动程序在启动驻留在不同编译单元中的内核时动态管理常量内存的不同视图之间的切换。虽然不能动态分配常量内存,但这种模式就足够了,因为 64K 限制不是系统范围的,它只适用于编译单元。

使用问题中引用的第一个模式:静态声明常量数据并在启动引用它的内核之前用 cudaMemcpyToSymbol 更新它。在第二种模式中,只有指针本身的读取才会通过常量内存。使用 读取 指针将由正常的 L1/L2 缓存层次结构提供服务。