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){intiblockIdx.x*blockDim.xthreadIdx.x;if(in){output[i]input[i]*coefficients[0];}}4. 如何初始化常量内存主机端不能直接对设备常量变量赋值coefficients[0]1.0f;// 错误这是设备端符号通常使用cudaMemcpyToSymbol将主机数据复制到常量内存floath_coefficients[64];for(inti0;i64;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 读取同一个地址这是常量内存最适合的访问模式floatvaluecoefficients[0];一个 Warp 中的所有线程都读取coefficients[0]线程 0 → coefficients[0] 线程 1 → coefficients[0] 线程 2 → coefficients[0] ...GPU 可以将读取结果广播给整个 Warp。5.2 Warp 读取不同地址如果一个 Warp 中的线程读取不同地址floatvaluecoefficients[threadIdx.x];访问变成线程 0 → coefficients[0] 线程 1 → coefficients[1] 线程 2 → coefficients[2] ...此时常量内存可能需要处理多个不同地址的读取请求访问效率可能下降。因此常量内存并不是“线程越多越快”关键是 Warp 内线程的访问地址是否一致。5.3 访问模式对比Warp 内访问模式常量内存适配程度所有线程读取同一个地址很适合适合广播少量不同地址且数据在缓存中可能可以接受每个线程读取完全不同的随机地址通常不适合频繁写入数据不适合常量内存只读6. 常量内存的典型应用6.1 固定系数例如图像滤波器的固定系数__constant__floatfilter[9];Kernel 中floatsum0.0f;for(inti0;i9;i){suminput_value[i]*filter[i];}如果所有线程按照相同顺序读取滤波器系数常量内存通常比较合适。6.2 查找表例如三角函数、颜色转换或量化操作使用的查找表__constant__intlookup_table[256];Kernel 中intvaluelookup_table[input[i]];如果不同线程访问的索引比较集中常量缓存可能带来收益。6.3 小型模型参数某些小型计算需要所有线程反复读取相同的权重或配置参数__constant__floatweights[128];这类数据满足数据规模小Kernel 期间不修改多个线程重复读取。6.4 几何和物理计算参数例如变换矩阵光照参数仿真常量坐标变换系数固定边界条件。7. 常量内存与其他内存的区别对比维度常量内存全局内存共享内存Kernel 中能否写入不能可以可以访问范围设备上多个线程只读设备上多个线程同一个 Block容量较小较大较小主要优势相同地址广播和缓存容量大、用途广Block 内快速共享数据初始化通常用cudaMemcpyToSymbolcudaMemcpy或其他方式Kernel 内加载适合场景小型只读公共数据输入、输出和大数据数据复用和线程协作简单判断所有线程读取相同的小数据优先考虑常量内存。Block 内线程需要读写共享数据考虑共享内存。数据量大或需要写回使用全局内存。8. 常量内存与const的区别这两个概念容易混淆。constconst是 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. 一个完整示例下面示例使用常量内存保存缩放系数#includecuda_runtime.h#includecstdio__constant__floatscale_factor;__global__voidscale(constfloat*input,float*output,intn){intiblockIdx.x*blockDim.xthreadIdx.x;if(in){output[i]input[i]*scale_factor;}}intmain(){constintn1024;constsize_t bytesn*sizeof(float);floath_scale2.0f;float*d_inputnullptr;float*d_outputnullptr;cudaMalloc(d_input,bytes);cudaMalloc(d_output,bytes);cudaMemcpyToSymbol(scale_factor,h_scale,sizeof(float));// 这里可以继续把输入数据复制到 d_inputconstintblock_size256;constintgrid_size(nblock_size-1)/block_size;scalegrid_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 需要修改数据。每个线程读取不同的随机地址。数据会频繁更新。数据只被单个线程使用。访问模式无法利用常量缓存和广播。例如valuetable[thread_id*large_stride];如果不同线程经常读取不同地址常量内存的优势可能不明显。此时应通过基准测试比较全局内存、只读缓存或其他方案。12. 总结CUDA 常量内存是一种适合小型只读数据的设备内存空间。它的核心优势是数据只读。多个线程可以访问同一份数据。具有缓存机制。Warp 内线程读取同一地址时可以进行广播。选择常量内存时可以依次判断数据是否只读 ↓ 数据规模是否较小 ↓ Kernel 执行期间是否保持不变 ↓ Warp 内线程是否经常读取相同地址 ↓ 如果是可以考虑常量内存最重要的原则是常量内存适合“小规模、只读、重复使用、Warp 内访问地址相同”的数据。
