ARTICLE · INTELLIGENCE

战地情报 · 详情页

来自尧图项目组的一线实战观察与深度解析

4.6 CUDA 常量内存

4.6 CUDA 常量内存

1. 常量内存是什么?

CUDA 常量内存(Constant Memory)是 GPU 上专门用于保存只读数据的一类内存空间。

它具有以下特点:

  • Kernel 中的线程只能读取,不能写入。
  • 通常由主机端在 Kernel 启动前初始化。
  • 所有线程都可以访问同一份数据。
  • 容量较小,通常适合保存几十 KB 以内的常量数据。
  • 具有缓存机制。
  • 当一个 Warp 中的线程读取同一个地址时,适合进行广播。

可以把它理解成:

主机端设置一次 ↓ GPU 中保存一份只读数据 ↓ 多个线程反复读取

2. 为什么需要常量内存?

假设一个 Kernel 中的所有线程都需要读取相同的参数:

线程 0 ─┐ 线程 1 ─┼── 读取相同的常量参数 线程 2 ─┤ 线程 3 ─┘

如果这些参数放在普通全局内存中,线程可能产生多次独立的内存访问。

如果放入常量内存,并且一个 Warp 中的线程读取同一个地址,GPU 可以将一次读取结果广播给多个线程:

常量内存 ↓ 一次读取 Warp 内多个线程同时获得数据

因此,常量内存适合:

  • 所有线程都读取相同数据。
  • 数据量较小。
  • Kernel 执行期间数据不会改变。
  • 数据会被多个线程重复使用。

3. 如何声明常量内存?

使用__constant__声明设备端常量变量:

__constant__floatcoefficients[64];

这里:

  • coefficients位于 CUDA 常量内存。
  • Kernel 可以读取它。
  • Kernel 不能修改它。
  • 它不是普通的主机变量。

Kernel 中可以直接使用:

__global__voidapply_filter(constfloat*input,float*output,intn){inti=blockIdx.x*blockDim.x+threadIdx.x;if(i<n){output[i]=input[i]*coefficients[0];}}

4. 如何初始化常量内存?

主机端不能直接对设备常量变量赋值:

coefficients[0]=1.0f;// 错误:这是设备端符号

通常使用cudaMemcpyToSymbol将主机数据复制到常量内存:

floath_coefficients[64];for(inti=0;i<64;++i){h_coefficients[i]=1.0f;}cudaMemcpyToSymbol(coefficients,h_coefficients,sizeof(h_coefficients));

典型流程是:

主机准备常量数据 ↓ cudaMemcpyToSymbol ↓ 常量内存初始化 ↓ 启动 Kernel ↓ Kernel 只读访问

也可以只更新一部分数据:

cudaMemcpyToSymbol(coefficients,h_coefficients,count*sizeof(float),offset*sizeof(float));

其中:

  • count * sizeof(float)表示复制字节数。
  • offset * sizeof(float)表示常量数组中的字节偏移。

5. 常量内存的访问模式

5.1 Warp 读取同一个地址

这是常量内存最适合的访问模式:

floatvalue=coefficients[0];

一个 Warp 中的所有线程都读取coefficients[0]:

线程 0 → coefficients[0] 线程 1 → coefficients[0] 线程 2 → coefficients[0] ...

GPU 可以将读取结果广播给整个 Warp。

5.2 Warp 读取不同地址

如果一个 Warp 中的线程读取不同地址:

floatvalue=coefficients[threadIdx.x];

访问变成:

线程 0 → coefficients[0] 线程 1 → coefficients[1] 线程 2 → coefficients[2] ...

此时常量内存可能需要处理多个不同地址的读取请求,访问效率可能下降。

因此,常量内存并不是“线程越多越快”,关键是 Warp 内线程的访问地址是否一致。

5.3 访问模式对比

Warp 内访问模式常量内存适配程度
所有线程读取同一个地址很适合,适合广播
少量不同地址,且数据在缓存中可能可以接受
每个线程读取完全不同的随机地址通常不适合
频繁写入数据不适合,常量内存只读

6. 常量内存的典型应用

6.1 固定系数

例如,图像滤波器的固定系数:

__constant__floatfilter[9];

Kernel 中:

floatsum=0.0f;for(inti=0;i<9;++i){sum+=input_value[i]*filter[i];}

如果所有线程按照相同顺序读取滤波器系数,常量内存通常比较合适。

6.2 查找表

例如三角函数、颜色转换或量化操作使用的查找表:

__constant__intlookup_table[256];

Kernel 中:

intvalue=lookup_table[input[i]];

如果不同线程访问的索引比较集中,常量缓存可能带来收益。

6.3 小型模型参数

某些小型计算需要所有线程反复读取相同的权重或配置参数:

__constant__floatweights[128];

这类数据满足:

  • 数据规模小;
  • Kernel 期间不修改;
  • 多个线程重复读取。

6.4 几何和物理计算参数

例如:

  • 变换矩阵;
  • 光照参数;
  • 仿真常量;
  • 坐标变换系数;
  • 固定边界条件。

7. 常量内存与其他内存的区别

对比维度常量内存全局内存共享内存
Kernel 中能否写入不能可以可以
访问范围设备上多个线程只读设备上多个线程同一个 Block
容量较小较大较小
主要优势相同地址广播和缓存容量大、用途广Block 内快速共享
数据初始化通常用cudaMemcpyToSymbolcudaMemcpy或其他方式Kernel 内加载
适合场景小型只读公共数据输入、输出和大数据数据复用和线程协作

简单判断:

  • 所有线程读取相同的小数据:优先考虑常量内存。
  • Block 内线程需要读写共享数据:考虑共享内存。
  • 数据量大或需要写回:使用全局内存。

8. 常量内存与const的区别

这两个概念容易混淆。

const

const是 C/C++ 中的类型限定符,表示通过某个变量或指针不能修改数据:

__global__voidkernel(constfloat*input){// 通过 input 只能读取数据}

这并不表示input一定位于 CUDA 常量内存中。

__constant__

__constant__是 CUDA 的内存空间限定符:

__constant__floatcoefficients[64];

它表示数据位于设备常量内存中,并具有常量内存的访问特性。

对比:

写法主要含义
const float*通过该指针只读
__constant__ float[]数据位于 CUDA 常量内存
const __constant__常量内存中的只读数据,通常不需要额外强调

因此:

const主要描述能不能写,__constant__主要描述数据放在哪里。


9. 使用常量内存时的注意事项

9.1 容量有限

常量内存容量有限,不适合保存大型数组、完整图像或大规模模型参数。

如果数据太大,应考虑:

  • 全局内存;
  • 纹理或只读缓存;
  • 共享内存;
  • 分块加载。

9.2 Kernel 执行期间不要修改

常量内存适合在 Kernel 启动前设置,Kernel 执行期间只读取。

如果需要更新参数,通常应当:

  1. 等待正在使用该常量数据的 Kernel 完成;
  2. 使用cudaMemcpyToSymbol更新;
  3. 再启动下一个 Kernel。

9.3 注意符号名称

cudaMemcpyToSymbol使用的是设备端符号:

__constant__floatcoefficients[64];cudaMemcpyToSymbol(coefficients,h_coefficients,sizeof(h_coefficients));

不要把它误写成普通设备指针传输方式。

9.4 注意复制大小

复制大小必须使用字节数:

cudaMemcpyToSymbol(coefficients,h_coefficients,64*sizeof(float));

不能只传入元素数量64,否则只会复制 64 个字节。

9.5 注意访问模式

如果每个线程都读取不同的随机地址,常量内存的广播优势会减弱。选择常量内存前,应先观察一个 Warp 的访问地址。


10. 一个完整示例

下面示例使用常量内存保存缩放系数:

#include<cuda_runtime.h>#include<cstdio>__constant__floatscale_factor;__global__voidscale(constfloat*input,float*output,intn){inti=blockIdx.x*blockDim.x+threadIdx.x;if(i<n){output[i]=input[i]*scale_factor;}}intmain(){constintn=1024;constsize_t bytes=n*sizeof(float);floath_scale=2.0f;float*d_input=nullptr;float*d_output=nullptr;cudaMalloc(&d_input,bytes);cudaMalloc(&d_output,bytes);cudaMemcpyToSymbol(scale_factor,&h_scale,sizeof(float));// 这里可以继续把输入数据复制到 d_inputconstintblock_size=256;constintgrid_size=(n+block_size-1)/block_size;scale<<<grid_size,block_size>>>(d_input,d_output,n);cudaDeviceSynchronize();cudaFree(d_input);cudaFree(d_output);return0;}

这个示例的关键点是:

主机变量 h_scale ↓ cudaMemcpyToSymbol ↓ 设备常量 scale_factor ↓ 所有线程读取 scale_factor

因为一个 Warp 中的线程读取的是同一个scale_factor地址,所以适合常量内存的广播访问。


11. 什么时候不应该使用常量内存?

以下情况通常不适合使用常量内存:

  • 数据量超过常量内存容量。
  • Kernel 需要修改数据。
  • 每个线程读取不同的随机地址。
  • 数据会频繁更新。
  • 数据只被单个线程使用。
  • 访问模式无法利用常量缓存和广播。

例如:

value=table[thread_id*large_stride];

如果不同线程经常读取不同地址,常量内存的优势可能不明显。此时应通过基准测试比较全局内存、只读缓存或其他方案。


12. 总结

CUDA 常量内存是一种适合小型只读数据的设备内存空间。

它的核心优势是:

  1. 数据只读。
  2. 多个线程可以访问同一份数据。
  3. 具有缓存机制。
  4. Warp 内线程读取同一地址时可以进行广播。

选择常量内存时,可以依次判断:

数据是否只读? ↓ 数据规模是否较小? ↓ Kernel 执行期间是否保持不变? ↓ Warp 内线程是否经常读取相同地址? ↓ 如果是,可以考虑常量内存

最重要的原则是:

常量内存适合“小规模、只读、重复使用、Warp 内访问地址相同”的数据。

RELATED READING

延伸阅读

更多一线实战笔记与深度复盘,助您持续精进