目录摘要常量内存常量缓存特性分配方式使用常量内存实现一维模板只读缓存分配方式实验总结摘要本文将介绍常量内存和只读缓存然后对比其差异常量内存常量内存是一种专用的全局内存区域用于存放只读数据。对于GPU端的内核代码常量内存是完全只读的GPU线程只能读取不能写入。主机端CPU可以在内核启动之前通过cudaMemcpyToSymbol等API向常量内存写入数据或者在内核执行完毕后读取数据。常量内存是位于片外而不是片上片上的是常量缓存关于常量缓存我们是查不到的其实对我们也没啥可编程的作用这个是硬件的作用不像共享内存一样开放给我们进行编程这样知道其大小就有用我们只能知道常量内存的大小cout常量内存大小prop.totalConstMemendl;由此可得常量内存大小是64KB常量缓存特性常量内存是存放数据的地方接下来讲的特性是属于常量缓存的常量缓存具有单周期广播性质也就是当请求的地址是同一个的时候常量缓存能够把数据广播给32个线程第 1 步Warp Scheduler 发射指令Warp Scheduler 从就绪的 Warp 中挑中它发射一条LDCLoad Constant指令给LSU。第 2 步LSU 把请求路由给常量缓存LSU 收到指令后收集这 32 个线程的地址。但它不做合并常量内存走的是常量缓存通路不走 L1而是直接把这 32 个地址发送给常量缓存控制器。第 3 步常量缓存控制器判断是否广播常量缓存控制器检查这 32 个地址如果所有地址都相同→ 触发广播。控制器直接从常量缓存的 SRAM 中取出这份数据通过专用的广播总线在一个时钟周期内发送给所有 32 个线程。LSU 负责把数据写入各线程的目标寄存器。如果地址不同→ 广播失效。控制器只能串行化处理这些请求一个接一个地从常量缓存或更底层的 DRAM读取数据。此时性能暴跌甚至可能比普通全局内存还慢。第 4 步如果常量缓存缺失如果请求的数据不在常量缓存的SRAM 里控制器会把请求转发给L2 缓存L2 再未命中就从片外 DRAM 的常量内存区域那 64KB搬来整个 128 字节 Cache Line。数据进入常量缓存后再按第 3 步的逻辑返回给 LSU。注意一旦这32个线程要的是不同的地址就会变串行如果是全局显存LSU拿到数据之后不是一个周期就能写到32个线程的寄存器的需要多个周期分配方式常量内存必须在所有函数外部以__constant__关键字声明。它的大小是静态固定的且必须在主机端通过专用 API 来初始化或更新数据。① 声明常量内存变量// 全局作用域文件作用域内可见 __constant__ float weights[256];② 从主机端拷贝数据到常量内存必须使用cudaMemcpyToSymbol不能用普通的cudaMemcpy。float h_weights[256]; // ... 初始化 h_weights ... cudaMemcpyToSymbol(weights, h_weights, sizeof(h_weights));为什么不能用cudaMemcpy因为__constant__变量不是普通的指针它没有可以被cudaMemcpy识别的“设备地址”。cudaMemcpyToSymbol内部会自动处理这种特殊的地址映射。③ 从主机端读取常量内存中的数据使用cudaMemcpyFromSymbol同样不能直接用cudaMemcpy。float h_copy[256]; cudaMemcpyFromSymbol(h_copy, weights, sizeof(weights));④ 如果常量变量在多个.cu文件中共享需要在声明时使用extern关键字一个文件定义其他文件声明// file1.cu: 定义 __constant__ float shared_data[1024]; // file2.cu: 声明 extern __constant__ float shared_data[1024];注意常量内存的生命周期和程序一致程序结束常量内存才结束初始化一次后面都可以一直用了内存类型生命周期声明方式常量内存(__constant__)整个应用程序最长文件作用域全局声明全局内存(cudaMalloc)由程序员控制从分配cudaMalloc到释放cudaFree动态分配共享内存(__shared__)线程块Block的生命周期核函数内声明寄存器 / 局部内存线程的生命周期核函数内声明局部变量使用常量内存实现一维模板相信人工智能发达的今天大家都听过卷积或者上过视觉的或者数值分析等课程如果没听过的去查一下很简单就是一个卷积核模板一个矩阵卷积就是把模板依次滑过矩阵在对应位置相乘再相加这是数值分析里的一维高阶差分求导模板那这就是一维高阶求导d_in是你的数据然后我们经过模板之后每个位置都用相邻的8个元素包括本身就是9个元素然后进行差分求导求出来的数据就是d_out[]每个点都有导数通过以上的背景我们知道c0c1c2c3是常量是模板那么可以使用常量内存而且此时的输入数据d_in可以使用共享内存因为每个数据都要被使用8次所以最好的办法就是一次加载多次使用那就是共享内存#includecuda_runtime.h #include ../freshman.hpp using namespace std; #define TEMPLATE_SIZE 9 #define TEMP_RADIO_SIZE (TEMPLATE_SIZE/2) #define BDIM 32 __constant__ float coef[TEMP_RADIO_SIZE]; void convolution(float *in,float *out,float* template_,const unsigned int array_size) { for(int iTEMP_RADIO_SIZE;iarray_size-TEMP_RADIO_SIZE;i) { for(int j1;jTEMP_RADIO_SIZE;j) { out[i]template_[j-1]*(in[ij]-in[i-j]); } //printf(%d:CPU :%lf\n,i,out[i]); } } __global__ void stencil_1d(float * in,float * out){ //block边缘会越界为了不越界并且能够计算所以填充前面四个后面四个一共8个位置 __shared__ float smem[BDIM2*TEMP_RADIO_SIZE]; //计算全局索引 int idxthreadIdx.xblockDim.x*blockIdx.x; //跳过前面四个进行加载数据到共享内存 int sidxthreadIdx.xTEMP_RADIO_SIZE; smem[sidx]in[idx]; //此时处理前面四个和后面四个 if (threadIdx.xTEMP_RADIO_SIZE){ //如果此时的block是中间的那么前面和后面都有数据直接用 if(idxTEMP_RADIO_SIZE) smem[sidx-TEMP_RADIO_SIZE]in[idx-TEMP_RADIO_SIZE]; if(idxgridDim.x*blockDim.x-BDIM) smem[sidxBDIM]in[idxBDIM]; } //同步保证所有数据都加载到共享内存 __syncthreads(); //全局边缘位置不计算 if (idxTEMP_RADIO_SIZE||idxgridDim.x*blockDim.x-TEMP_RADIO_SIZE) return; float temp.0f; //计算差分 #pragma unroll for(int i1;iTEMP_RADIO_SIZE;i) { tempcoef[i-1]*(smem[sidxi]-smem[sidx-i]); } out[idx]temp; //printf(%d:GPU :%lf,\n,idx,temp); } //read only __global__ void stencil_1d_readonly(float * in,float * out,const float* __restrict__ dcoef) { __shared__ float smem[BDIM2*TEMP_RADIO_SIZE]; int idxthreadIdx.xblockDim.x*blockIdx.x; int sidxthreadIdx.xTEMP_RADIO_SIZE; smem[sidx]in[idx]; if (threadIdx.xTEMP_RADIO_SIZE) { if(idxTEMP_RADIO_SIZE) smem[sidx-TEMP_RADIO_SIZE]in[idx-TEMP_RADIO_SIZE]; if(idxgridDim.x*blockDim.x-BDIM) smem[sidxBDIM]in[idxBDIM]; } __syncthreads(); if (idxTEMP_RADIO_SIZE||idxgridDim.x*blockDim.x-TEMP_RADIO_SIZE) return; float temp.0f; #pragma unroll for(int i1;iTEMP_RADIO_SIZE;i) { tempdcoef[i-1]*(smem[sidxi]-smem[sidx-i]); } out[idx]temp; //printf(%d:GPU :%lf,\n,idx,temp); } int main(int argc,char** argv){ int dimxBDIM; unsigned int nxy116; int nBytesnxy*sizeof(float); //Malloc float* in_host(float*)malloc(nBytes); float* out_gpu(float*)malloc(nBytes); float* out_cpu(float*)malloc(nBytes); memset(out_cpu,0,nBytes); //initialData(in_host,nxy); //cudaMalloc float *in_devNULL; float *out_devNULL; initialData(in_host,nxy); float templ_[]{-1.0,-2.0,2.0,1.0}; CHECK(cudaMemcpyToSymbol(coef,templ_,TEMP_RADIO_SIZE*sizeof(float))); CHECK(cudaMalloc((void**)in_dev,nBytes)); CHECK(cudaMalloc((void**)out_dev,nBytes)); CHECK(cudaMemcpy(in_dev,in_host,nBytes,cudaMemcpyHostToDevice)); CHECK(cudaMemset(out_dev,0,nBytes)); // cpu compute double iStartefficiency(); convolution(in_host,out_cpu,templ_,nxy); double iElapsefficiency()-iStart; //printf(CPU Execution Time elapsed %f sec\n,iElaps); // stencil 1d dim3 block(dimx); dim3 grid((nxy-1)/block.x1); double stencil_1d_startefficiency(); stencil_1dgrid,block(in_dev,out_dev); CHECK(cudaDeviceSynchronize()); double stencil_1d_iElapsefficiency()-stencil_1d_start; printf(stencil_1d Time elapsed %f sec\n,stencil_1d_iElaps); CHECK(cudaMemcpy(out_gpu,out_dev,nBytes,cudaMemcpyDeviceToHost)); check(out_cpu,out_gpu,nxy); CHECK(cudaMemset(out_dev,0,nBytes)); // stencil 1d read only float * dcoef_ro; CHECK(cudaMalloc((void**)dcoef_ro,TEMP_RADIO_SIZE * sizeof(float))); CHECK(cudaMemcpy(dcoef_ro,templ_,TEMP_RADIO_SIZE * sizeof(float),cudaMemcpyHostToDevice)); double stencil_1d_readonly_startefficiency(); stencil_1d_readonlygrid,block(in_dev,out_dev,dcoef_ro); CHECK(cudaDeviceSynchronize()); double stencil_1d_readonly_iElapsefficiency()-stencil_1d_readonly_start; printf(stencil_1d_readonly Time elapsed %f sec\n,stencil_1d_readonly_iElaps); CHECK(cudaMemcpy(out_gpu,out_dev,nBytes,cudaMemcpyDeviceToHost)); check(out_cpu,out_gpu,nxy); cudaFree(dcoef_ro); cudaFree(in_dev); cudaFree(out_dev); free(out_gpu); free(out_cpu); free(in_host); cudaDeviceReset(); return 0; }只读缓存只读缓存在现代 Ada Lovelace 架构中并不是一个独立的物理硬件而是L1 数据缓存的一种“VIP 使用模式”。它的作用是当你明确告诉编译器“我这块数据在整个核函数运行期间都不会被修改”时L1 缓存会用更优的策略来缓存这些数据让它们尽量不被频繁读写的数据挤掉。分配方式方式一使用__ldg()内联函数最直接__ldg()是 CUDA 提供的一个内置函数专门用于发起一次只读的全局内存加载。你不需要改变指针的类型声明只需在读取数据时用__ldg()包裹一下即可。__global__ void kernel_readonly(float *output, float *large_table, int n) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx n) { // 用 __ldg() 显式告诉编译器这次加载是只读的请走只读缓存路径 float val __ldg(large_table[idx]); output[idx] val * 2.0f; } }关键点__ldg()作用于每次具体的加载操作你可以只对关键的只读数据使用它其他读写数据仍然走普通路径。方式二使用const __restrict__指针更省心在核函数的参数声明中给只读指针加上const __restrict__修饰。编译器会自动把符合条件的内存访问映射到只读缓存路径。// 注意参数中的 const __restrict__ __global__ void kernel_readonly(float *output, const float* __restrict__ large_table, int n) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx n) { // 不需要 __ldg()编译器自动走只读缓存路径 float val large_table[idx]; output[idx] val * 2.0f; } }关键点const表示数据不可修改__restrict__告诉编译器这个指针是访问这块内存的唯一途径没有别名编译器因此可以放心地把所有对large_table的读取都映射到只读缓存路径。实验//read only __global__ void stencil_1d_readonly(float * in,float * out,const float* __restrict__ dcoef) { __shared__ float smem[BDIM2*TEMP_RADIO_SIZE]; int idxthreadIdx.xblockDim.x*blockIdx.x; int sidxthreadIdx.xTEMP_RADIO_SIZE; smem[sidx]in[idx]; if (threadIdx.xTEMP_RADIO_SIZE) { if(idxTEMP_RADIO_SIZE) smem[sidx-TEMP_RADIO_SIZE]in[idx-TEMP_RADIO_SIZE]; if(idxgridDim.x*blockDim.x-BDIM) smem[sidxBDIM]in[idxBDIM]; } __syncthreads(); if (idxTEMP_RADIO_SIZE||idxgridDim.x*blockDim.x-TEMP_RADIO_SIZE) return; float temp.0f; #pragma unroll for(int i1;iTEMP_RADIO_SIZE;i) { tempdcoef[i-1]*(smem[sidxi]-smem[sidx-i]); } out[idx]temp; //printf(%d:GPU :%lf,\n,idx,temp); }两次的数据非常接近甚至有时候只读的还能跑过常量可能是跟数据量非常小的问题但是只读的波动非常大有时候还很慢总结特性常量内存 (Constant Memory)只读缓存 (Read-Only Cache)物理存储片外 DRAM但拥有一块独立的片上常量缓存Ada 架构为 8KB/SM数据仍在片外 DRAM加速路径是L1 数据缓存的一部分非独立硬件触发方式__constant__声明 cudaMemcpyToSymbol初始化__ldg()内联函数或const __restrict__指针核心优势单周期广播一个 Warp 内所有线程读同一地址时一个时钟周期完成缓存保护数据被标记为只读优先保留在 L1 缓存不被普通读写数据轻易挤掉最佳访问模式全 Warp 统一地址如所有线程读同一个权重每个线程读不同地址但这些地址通常是连续且只读的容量限制64KB整个 GPU受 L1 缓存容量和全局内存容量共同限制远大于 64KB主机端如何写入cudaMemcpyToSymbol专用 APIcudaMemcpy普通拷贝常量内存是专为“小、统一、只读”场景打造的广播利器只读缓存是为“大、各异、只读”场景提供的 L1 保护策略。