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 内快速共享 |
| 数据初始化 | 通常用cudaMemcpyToSymbol | cudaMemcpy或其他方式 | 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 执行期间只读取。
如果需要更新参数，通常应当：
- 等待正在使用该常量数据的 Kernel 完成；
- 使用
cudaMemcpyToSymbol更新； - 再启动下一个 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 常量内存是一种适合小型只读数据的设备内存空间。
它的核心优势是：
- 数据只读。
- 多个线程可以访问同一份数据。
- 具有缓存机制。
- Warp 内线程读取同一地址时可以进行广播。
选择常量内存时，可以依次判断：
数据是否只读？ ↓ 数据规模是否较小？ ↓ Kernel 执行期间是否保持不变？ ↓ Warp 内线程是否经常读取相同地址？ ↓ 如果是，可以考虑常量内存最重要的原则是：
常量内存适合“小规模、只读、重复使用、Warp 内访问地址相同”的数据。
