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

art*_* me 1 cuda

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

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

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
Run Code Online (Sandbox Code Playgroud)

再次根据编程指南我们可以使用:

__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));
Run Code Online (Sandbox Code Playgroud)

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

所以这里有一些问题:

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

Arc*_*are 6

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

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