常量内存特性与广播机制在 CUDA 的内存层级中常量内存Constant Memory是一块常被忽视却极具特色的存储空间。它的容量只有64KB相较于全局内存动辄数 GB 的规模显得微不足道但正是这个小而精的设计让它在特定场景下能带来成倍的性能提升。理解常量内存关键在于抓住它的两个核心机制同地址广播与异地址串行。声明与数据传输__constant__与cudaMemcpyToSymbol在设备端常量内存的声明方式与全局内存类似使用__constant__限定符__constant__ float kernel_weights[128]; // 128 个 float共 512 字节 __constant__ int image_dim[3]; // 存储图像维度信息需要强调的是__constant__变量必须声明在文件作用域即所有函数体外而非某个函数的内部。这说明常量内存的生命周期贯穿整个 CUDA 上下文而非某个 kernel 的局部存在。常量内存的宿主机端访问方式与普通全局内存不同——不能直接通过cudaMemcpy复制数据。CUDA 提供了专用的 APIcudaMemcpyToSymbolfloat host_weights[128]; // ... 在宿主机端填充 host_weights ... cudaMemcpyToSymbol(kernel_weights, host_weights, sizeof(float) * 128, 0, cudaMemcpyHostToDevice); // 参数依次为设备端符号、宿主机指针、复制字节数、偏移量、复制方向同理读取常量内存中的数据使用cudaMemcpyFromSymbol方向参数改为cudaMemcpyDeviceToHost。这种独立的 API 设计反映了常量内存在编译链接层面的特殊性——它更接近一个内嵌于指令寻址空间的常量池而不是一块普通的可寻址显存。64KB 是常量内存的硬性上限。这个限制源于 NVIDIA 的硬件设计常量内存在 GPU 上配有独立的一级缓存Constant Cache缓存行宽度为 64 字节与共享内存一致。当数据量超过 64KB 时只能将多余部分存放在全局内存中由编译器在访问时通过额外的全局内存指令处理。同地址广播一个内存读取服务所有线程常量内存的性能秘密在于它的缓存架构。当 warp 中的 32 个线程同时访问同一个常量内存地址时硬件只需发起一次内存读取操作然后通过广播机制将数据分发给所有线程。这个过程完全由硬件完成不需要任何额外的指令参与。更具体的数字是NVIDIA 的常量缓存Constant Cache以 64 字节为缓存行单位带宽设计为一次读取可向 32 个线程同时广播 4 字节数据。因此一个 warp 内所有线程访问同一地址时整个操作只耗时一条内存指令的时钟周期通常约 4 个周期与寄存器访问相当。相比之下如果这些数据放在全局内存中即使 L2 缓存命中也需要至少两个线程束周期处理内存事务。与之形成鲜明对比的是全局内存的访问模式无论线程访问的是同一地址还是不同地址硬件都会将内存访问合并成尽量少的事务。若 32 个线程访问同一地址全局内存在 L2 缓存命中时也能较快响应通常约 200-300 个周期但常数缓存的广播访问在延迟上几乎可以忽略不计。将这个概念与共享内存对比更有助于建立直觉共享内存是按槽位访问——每个线程独立读取自己的槽位硬件无广播能力只有 bank 冲突时会产生串行化。常量内存是按广播访问——所有线程可以同时获得同一份数据不涉及任何冲突。因此当所有线程需要读取同一个值如卷积核系数、归一化阈值、循环上限时常量内存是比共享内存更优的选择它天然避免了共享内存的 bank 冲突问题。不过这种广播的优雅是有限度条件的——它要求所有线程访问的是完全相同的地址。一旦线程访问的地址不同瓶颈就会出现。异地址串行退化的访问模式当 warp 内不同线程访问不同地址的常量内存时性能逻辑则完全反转。常量缓存的访问流程是硬件在每个时钟周期只能处理一个内存请求并将结果广播给请求该地址的线程如果存在多个不同的地址请求硬件会将这些请求串行化——每周期只服务一个地址直至处理完所有请求。NVIDIA 官方文档给出的描述是如果 warp 中的线程访问了 n 个不同的常量内存地址则该 warp 需要n 个独立的时钟周期来完成这次访问需要可能更多视分支与数据分布。这意味着当 32 个线程全部访问不同地址时常量内存的读取速度会退化为单个内存事务的 32 倍延迟——远超全局内存合并访问的性能。一种直观的类比广播是全体整齐划一串行则是排成一列的过窄门。每次只允许一个访客通过其余线程只能等待。与全局内存合并访问相比这种串行模式在吞吐量上往往处于劣势。这个特性直接推导出一个设计准则常量内存的变量应尽量保证在分支内对齐访问。如果 kernel 依赖线程 ID 或索引动态计算常量内存下标如weights[threadIdx.x]性能会急剧下降。常见的模式是所有线程读取同一个constants[i]其中i由外部循环变量控制而非由线程 ID 直接映射。以下面两段代码作为对比// 低效模式threadIdx 索引常量内存 → 32 个地址 → 32 次串行访问 float w kernel_weights[threadIdx.x % 128]; // 高效模式所有线程访问同一地址 __shared__ float shared_w; if (threadIdx.x 0) shared_w kernel_weights[blockIdx.x % 128]; __syncthreads(); float w shared_w;第二种方式利用共享内存将常量值快速分发到所有线程消除了常量缓存的串行化惩罚。小结与选择策略常量内存的取舍可以浓缩成一句话以空间换广播以同步失速度。64KB 的容量是明确的硬约束同地址广播是实现高吞吐的首选路径适合所有线程共享系数、参数、配置等场景而异地址访问则会让常量内存变成性能陷阱必须用共享内存或全局内存替代。开发时先将常量数据按全程同址访问设计再根据实际分支结构评估是否改用其他存储空间——这样就能避免踩入大索引差异化访问的性能误区。理解了常量内存的这套广播-串行权衡之后下一个自然的问题是当我们需要更大的存储空间、更灵活的缓存策略时常量内存力不能及的场景该交给谁这正是纹理内存登场的地方。纹理内存同样拥有专用的缓存层次但它以二维局部性为核心设计哲学与常量内存的单值广播走了一条截然不同的道路。常量内存使用案例滤波器系数理解了常量内存的广播与串行机制后最自然的疑问是什么场景能真正吃到广播的红利答案是——当一个 warp 内的所有线程读取同一个地址时。这在图像处理中恰好有一个极其典型的场景卷积滤波器的系数读取。为什么滤波器系数适合放在常量内存以 3×3 的高斯模糊为例卷积核包含 9 个浮点系数。在一张 1920×1080 的图像上朴素实现为每个像素执行 9 次全局内存读取每次读一个系数总读取量为1920×1080×9≈18661920 \times 1080 \times 9 \approx 18661920×1080×9≈1866万次。这些系数对所有像素完全相同且在整个核函数执行期间只读不写——完美契合常量内存的两大前提。更关键的是访问模式。当 warp 中的 32 个线程处理相邻像素时它们在同一个循环迭代中访问的是同一个卷积核系数。假设某次循环正在计算K1,1K_{1,1}K1,1若系数存于全局内存32 个线程发出 32 次读取请求虽然 L2 缓存可能合并但首次访问必然产生多次内存事务若系数存于常量内存32 个线程读取同一地址硬件只需1 次内存读取结果通过广播总线同时送达 32 个线程。这正是第 1 节中同地址广播机制的实战价值——当过滤器系数的读取次数等于像素总数乘以核大小数百万次时广播带来的收益会被放大到极致。完整代码示例3×3 高斯模糊下面给出一个可直接编译运行的示例。它定义了一个 3×3 高斯核存入常量内存并利用广播机制加速卷积运算。#include cstdio #include cmath // 卷积核大小 #define KERNEL_RADIUS 1 #define KERNEL_SIZE (2 * KERNEL_RADIUS 1) // 3 // 图像尺寸为简化使用一维表示二维图像 #define WIDTH 64 #define HEIGHT 64 #define IMG_SIZE (WIDTH * HEIGHT) // 常量内存声明文件作用域 // 3x3 高斯核系数仅 9 个 float远远小于 64KB 限制 __constant__ float c_kernel[KERNEL_SIZE * KERNEL_SIZE]; // 设备端卷积核函数 __global__ void gaussianBlur(const float* input, float* output, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width || y height) return; float sum 0.0f; // 双重循环遍历 3x3 邻域 // 关键点同一 warp 的线程在相同迭代中读取 c_kernel 的同一地址 // → 触发常量内存的同地址广播机制 for (int dy -KERNEL_RADIUS; dy KERNEL_RADIUS; dy) { for (int dx -KERNEL_RADIUS; dx KERNEL_RADIUS; dx) { int nx x dx; int ny y dy; // 边界处理越界像素用 0 填充 float pixel 0.0f; if (nx 0 nx width ny 0 ny height) { pixel input[ny * width nx]; } // 计算线性索引并累加 int kIdx (dy KERNEL_RADIUS) * KERNEL_SIZE (dx KERNEL_RADIUS); sum pixel * c_kernel[kIdx]; } } output[y * width x] sum; } // 主机端代码 int main() { // 3x3 高斯核归一化前 float h_kernel[KERNEL_SIZE * KERNEL_SIZE] { 1, 2, 1, 2, 4, 2, 1, 2, 1 }; // 归一化使所有权重之和为 1 float sum 0.0f; for (int i 0; i KERNEL_SIZE * KERNEL_SIZE; i) sum h_kernel[i]; for (int i 0; i KERNEL_SIZE * KERNEL_SIZE; i) h_kernel[i] / sum; // 使用 cudaMemcpyToSymbol 将系数从主机复制到设备常量内存 cudaMemcpyToSymbol(c_kernel, h_kernel, sizeof(h_kernel)); // 注意cudaMemcpyToSymbol 的默认方向是 HostToDevice此处无需显式指定 // 分配图像内存 float *d_input, *d_output; cudaMalloc(d_input, IMG_SIZE * sizeof(float)); cudaMalloc(d_output, IMG_SIZE * sizeof(float)); // 构造简单的测试图像中心区域为 1.0其余为 0.0 float* h_input (float*)malloc(IMG_SIZE * sizeof(float)); for (int i 0; i IMG_SIZE; i) h_input[i] 0.0f; for (int y HEIGHT/4; y 3*HEIGHT/4; y) for (int x WIDTH/4; x 3*WIDTH/4; x) h_input[y * WIDTH x] 1.0f; cudaMemcpy(d_input, h_input, IMG_SIZE * sizeof(float), cudaMemcpyHostToDevice); // 启动核函数block 16x16grid 按需计算 dim3 block(16, 16); dim3 grid((WIDTH block.x - 1) / block.x, (HEIGHT block.y - 1) / block.y); gaussianBlurgrid, block(d_input, d_output, WIDTH, HEIGHT); // 核对错误 cudaError_t err cudaDeviceSynchronize(); if (err ! cudaSuccess) { printf(CUDA error: %s\n, cudaGetErrorString(err)); return -1; } // 拷贝结果回主机并验证核心区域 float* h_output (float*)malloc(IMG_SIZE * sizeof(float)); cudaMemcpy(h_output, d_output, IMG_SIZE * sizeof(float), cudaMemcpyDeviceToHost); // 验证中心区域 (WIDTH/2, HEIGHT/2) 的模糊结果 int cx WIDTH / 2, cy HEIGHT / 2; printf(Center pixel after blur: %f\n, h_output[cy * WIDTH cx]); // 预期输出中心区域被周围 0 值像素稀释结果在 0.3~0.6 之间 // 清理 free(h_input); free(h_output); cudaFree(d_input); cudaFree(d_output); return 0; }编译指令nvcc -o blur blur.cu。运行后中心像素的输出值约为0.44由中心 4、四个正交邻域 2×4/个、四个对角邻域 1×4/个加权叠加后被 16 归一化得到。代码要点拆解__constant__声明必须位于文件作用域。第 1 节提到常量内存的声明要求这里再次验证c_kernel被定义为全局数组因为常量内存在 CUDA 中不支持动态分配且所有核函数共享同一份常量内存。9 个 float 仅占 36 字节远低于 64KB 上限。cudaMemcpyToSymbol是唯一的人口。它负责将主机的h_kernel拷贝到设备端的c_kernel。注意不需要在参数中写cudaMemcpyHostToDevice因为cudaMemcpyToSymbol从主机到设备是默认方向。若要从设备读回常量内存则需显式使用cudaMemcpyDeviceToHost。核函数内层循环是广播的核心。当dx0, dy0时所有线程都在读c_kernel[4]。此时硬件检测到 warp 内 32 个线程请求同一地址触发单次内存读取加广播。整个核函数执行过程中c_kernel[kIdx]被读取的次数等于线程总数 × 9 次但全部命中广播路径全局内存流量几乎为零。用共享内存做对比何时常量内存更优有人会问把卷积核放进共享内存不是一样的吗确实共享内存也有高速缓存特性。但两者有本质差异对比维度常量内存共享内存访问模式同一地址→广播任意地址→按槽位访问存储上限64KB48KB可配至 96KB数据更新核函数启动前一次性赋值核函数内可动态写入硬件开销广播总线单次传输需同步 按槽位冲突管理如果滤波器系数在核函数运行过程中需要动态更新比如自适应滤波共享内存是更好的选择。但对于静态系数的传统卷积常量内存的广播机制让所有线程在同一个循环迭代中只产生一次内存事务——这种零冲突特性是共享内存按槽位分配的访问方式所不具备的不同槽位可能产生 bank conflict。因此在卷积核系数固定且 warp 内线程步调一致时常量内存是更精确的选择。由此可以提炼出一个普适的决策准则当一个 warp 内的所有线程在同一个指令周期内读取同一地址的只读数据时常量内存就是最优解。卷积核系数、查找表LUT、数学函数的预计算参数都属于这类一读多用的数据。而接下来要讨论的纹理内存则把这种特殊缓存路径的思路推进到了更广阔的维度——它不仅提供缓存还内置了硬件插值、归一化坐标等图像处理专属能力。纹理内存与硬件纹理单位如果说常量内存通过广播特性解决了同一地址的高频访问问题那么纹理内存Texture Memory则通过缓存特性解决了另一类截然不同的访问模式——相邻地址的密集读取。这正是图像处理中像素采样的典型形态一个像素的邻域数据往往在下一次计算中立即被再次使用。常量内存管的是所有人读同一个数纹理内存管的是每个人读自己附近的那片数。空间局部性纹理缓存的底层逻辑图像处理的核心计算模式是邻域操作均值模糊取 3×3 邻域边缘检测取 Sobel 核覆盖的区域双线性插值取周围 4 个像素。这类操作有一个天然的访存特征——空间局部性Spatial Locality当前线程访问的像素附近极大概率是相邻线程下一个要访问的目标。全局内存的缓存行为并不理想它的缓存行通常 128 字节虽然能覆盖相邻地址但全局内存缓存更多是为通用数据复用设计缺乏对二维空间结构的感知。纹理内存则专门针对这一模式做了硬件优化——它内部实现了专用的纹理缓存该缓存针对二维甚至三维空间的遍历模式进行了调优。当一个线程读取了坐标(x,y)(x, y)(x,y)处的像素硬件会预取(x1,y)(x1, y)(x1,y)、(x,y1)(x, y1)(x,y1)附近的纹素texel到缓存中。相邻线程访问相邻坐标时命中率极高。以一张 1920×1080 的灰度图像为例处理 3×3 模糊时全局内存方式每个像素需要 9 次独立访问而纹理内存方式下第一次读取(x,y)(x, y)(x,y)时周围的纹素已驻留缓存后续 8 个邻域坐标的访问直接命中纹理缓存实际的内存流量降低了数倍。缓存未命中率从全局内存的数个数量级之高的循环重复访问降至几近于零。// 纹理内存的申请与绑定 texturefloat, 2, cudaReadModeElementType texRef; // 声明 2D 纹理引用 cudaChannelFormatDesc desc cudaCreateChannelDescfloat(); cudaArray* cuArray; cudaMallocArray(cuArray, desc, width, height); // 分配二维数组 // 将图像数据拷贝到 CUDA 数组纹理只能绑定到 CUDA 数组 cudaMemcpyToArray(cuArray, 0, 0, h_data, width * height * sizeof(float), cudaMemcpyHostToDevice); cudaBindTextureToArray(texRef, cuArray); // 完成绑定设备端通过tex2D(texRef, x, y)进行采样由硬件纹理单位Texture Unit完成地址计算和值的返回。纹理硬件Texture Unit是 GPU 上独立于通用计算核心的专用电路它负责坐标换算、缓存管理和数据过滤——这些操作不占用 CUDA 核心的计算资源而是由纹理单元的专用逻辑并行完成。硬件插值让双线性采样近乎免费纹理内存最令人惊艳的特性是硬件插值Hardware Interpolation。在图像缩放、旋转、坐标变换等场景中目标像素往往落在源图像的非整数坐标上。传统的计算方式需要软件实现双线性插值取周围 4 个像素、计算权重、加权求和。这需要约 4 次内存访问和多次浮点运算。纹理硬件将这一过程封装为单条指令。当使用归一化坐标Normalized Coordinates采样时——坐标范围映射到[0,1)[0, 1)[0,1)——硬件会自动处理坐标到纹素地址的换算并通过cudaFilterModeLinear模式直接在硬件层面完成双线性插值texturefloat, 2, cudaReadModeElementType texRef; texRef.addressMode[0] cudaAddressModeWrap; // 坐标环绕模式 texRef.addressMode[1] cudaAddressModeWrap; texRef.filterMode cudaFilterModeLinear; // 硬件双线性插值 // 采样时直接传归一化坐标0.0 ~ 1.0 之间 float val tex2D(texRef, nx, ny); // nx, ny ∈ [0, 1)这组代码中tex2D的返回值就是已经完成双线性插值的结果。软件实现需要至少 4 次全局/共享内存读取加 4 次乘加运算硬件方案只需 1 次纹理读取指令。对于大规模图像变换这种差异意味着数倍的吞吐提升——并且像素越连续、坐标越密集插值硬件的工作效率越高因为它内部有专用的插值单元流水线。2D/3D 局部性的完整收益除了插值纹理内存的缓存设计也超越了简单的一维缓存。2D 纹理缓存针对行、列两个维度的访问都做了预取优化3D 纹理缓存则进一步覆盖深度维度的局部性。这意味着体数据如医学 CT 切片、气候模拟网格的访问也能获得相似的红利。从整个 CUDA 内存体系来看纹理内存的价值定位是清晰而互补的对比维度全局内存共享内存常量内存纹理内存适用访问模式任意读写同 block 内协作warp 内同地址广播相邻地址密集读取板上缓存L2无显式管理常量缓存专用纹理缓存插值支持无无无硬件级插值边界处理无手动无Wrap/Clamp 模式理解了这个坐标就能回答一个常见困惑为什么有了共享内存还要纹理内存共享内存需要程序员手动管理数据搬运和复用而纹理缓存是硬件自动维护的。当你不想显式管理块边界、但又需要利用空间局部性时纹理内存是更轻量的选择——你只管采样缓存策略交给硬件。常量内存靠广播解决了所有线程读同一地址的效率问题纹理内存靠专用的缓存硬件和插值电路解决了相邻线程读相邻地址的效率问题。这两种内存共同补齐了 CUDA 内存体系中读密集、规律性访问的短板——接下来的问题很自然纹理和常量之外还有一类适用于散列访问的只读缓存机制它与纹理内存有何本质区别这正是下一节要展开的内容。只读缓存与__ldg再讨论常量内存的广播机制解决了所有人读同一个数的问题纹理内存的缓存与插值硬件解决了每个人读自己附近那片数的问题。但这两者都不是通用方案常量内存受限于 64KB 容量纹理内存则需要引入纹理坐标、采样器等一套繁琐的 API。那么当数据量超过常量内存上限、访问模式又不适合纹理内存时还有没有第三条路答案是肯定的——只读缓存Read-Only Cache。它是现代 GPUKepler 架构及以后内置的一条独立缓存通路专门服务于只读访问全局内存的场景。只读缓存绕过 L1 的独立通路在 Fermi 时代所有全局内存访问都要经过统一的 L1 缓存——如果你用const __restrict__修饰指针编译器会尝试通过 L1 缓存路径来读取数据但这并不保证只读语义被充分利用。Kepler 架构引入了专门的数据缓存通路即只读缓存它与 L1 数据缓存并行存在两者都有各自独立的能力。在编译器的眼中一个被标记为只读的全局内存访问会走只读缓存——这带来的收益在于只读缓存拥有比 L1 数据缓存更高的带宽只读缓存支持全速率的 128 字节/周期的读取而 L1 数据缓存则有所保留只读缓存的命中延迟与常量和纹理缓存相当几十个周期量级但容量明显大于常量内存从编程接口的角度让代码走只读缓存有两种方式// 方式一显式使用 __ldg 内置函数 __device__ float kernel_explicit(const float* data, int idx) { return __ldg(data idx); // 显式告知编译器这是一次只读访问 } // 方式二使用 const __restrict__ 修饰指针让编译器自动生成 LDG 指令 __global__ void kernel_auto(const float* __restrict__ data, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) out[i] data[i] * 2.0f; // 编译器自动识别为只读访问 }需要说明的是__ldg并非凭空发明的新功能——它只是强制编译器把某个访问编译为 LDGLoad Global指令。而const __restrict__则给了编译器一个证明“这块数据在核函数执行期间不会修改”编译器从而可以进行自动优化。两行代码看起来相似本质不同__ldg是命令__restrict__是承诺。在引入__ldg时NVIDIA 官方给出的性能数据显示当只读访问存在空间局部性即相邻线程访问相邻地址时只读缓存可以带来约 1.7 倍的带宽提升特别是与没有走只读通路的版本相比。这与我们说纹理缓存命中空间局部性的逻辑一脉相承只不过纹理缓存需要坐标而只读缓存只是简单地缓存 32 字节的缓存行天然适合线性数组的遍历。适用差异三大特殊内存如何选现在我们对 CUDA 中三种带缓存的内存通路已经有了全景视野。它们各自的物理基础不同、适用场景各异但核心区别可以归结为一个问题一个 warp 里的 32 个线程访问的地址分布是什么样的为了便于决策我们可以用一张表来总结内存通路最佳访问模式容量限制硬件插值编程方式常量内存所有线程读同一地址广播64KB无__constant__cudaMemcpyToSymbol纹理内存相邻线程读相邻地址局部性可达数十 GB绑定全局内存有双线性/三线性插值纹理对象 tex2D/tex3D只读缓存相邻线程读相邻地址局部性无缓存全局内存无__ldg()或const __restrict__更具体地看常量内存的不可替代性。当 warp 内所有线程必须读取完全相同的地址时如卷积核系数、数学查表广播机制将一次内存访问广播给所有线程吞吐量为一次访问的代价。任何其他方式——包括只读缓存——都无法替代这一点因为它们的缓存粒度是缓存行如 32 字节当一个 warp 读取同一地址时整条缓存行只有一次访问被利用其余存储空间完全浪费。纹理内存的独特优势在于插值。双线性/三线性插值并非一个可以在通用计算中随意简化掉的特性。图像缩放、体渲染、光学仿真等场景中采样坐标往往落在像素网格之外——使用纹理内存硬件直接返回插值结果避免了在代码中层叠 4 次2D甚至 8 次3D的混合计算。若这些操作在通用路径上实现每一层插值都会消耗若干条乘加指令和若干次内存访问。只读缓存则是通用领域的均衡解。数据量超过 64KB、不需要插值、但访问模式有空间局部性——这是最普遍的情况。以矩阵转置为例线程 i 按行步长读取矩阵 A 的若干行每个 warp 内 32 个线程读取的地址彼此相差 4 字节的倍数——正好命中只读缓存的缓存行粒度。这种情况下__ldg是唯一不需要改变数据布局、不需要引入纹理坐标、只需一行代码改动就能获得缓存增益的方案。决策框架从访问模式出发如果以一句话总结三类通路的取舍逻辑可以沿着如下思路来判断先看是否需要插值—— 如果需要硬件插值直接选择纹理内存无其他替代再问数据量是否在 64KB 以内—— 是且 warp 内所有线程访问同一地址选择常量内存剩余情况—— 数据量大、只读、有空间局部性使用__ldg或const __restrict__让编译器走只读缓存这个框架并不替代实测——CUDA 的优化空间往往超出直觉的线性推断。但它的意义在于当我们面对一个新的核函数时不再被用纹理内存总是更好或常量内存万能这类朴素信念所误导而是带着清晰的访问模式图景去选择最匹配的硬件通路。至此从共享内存到常量内存从纹理内存到只读缓存CUDA 自有存储体系的全貌已经完整呈现——而这四条通路之间的协同与博弈将在下一节以完整的决策树形式展开总结。内存类型选择决策表至此我们已经走完了 CUDA 内存体系的最后一公里常量内存用 64KB 的容量换来了同地址广播的高效通路纹理内存借由专用缓存与插值硬件统治了邻域读取场景只读缓存则为超出这两种特殊约束的通用访问提供了第三条路。但有一个问题始终悬而未决——面对一个具体的核函数到底该选哪种内存本节将把这四类内存放在同一张决策表中给出一个可复用的选择框架。五类内存的一次性对照在给出决策流程之前先建立全局视角。下表从容量、延迟、缓存机制、适用访问模式、典型场景五个维度对全局内存、共享内存、常量内存、纹理内存和只读缓存做一次系统性对照维度全局内存共享内存常量内存纹理内存只读缓存容量数 GB48KB~228KB按架构64KB数 GB映射全局数 GB映射全局延迟命中400~600 周期约 20~30 周期约 200~300 周期约 200~300 周期约 200~300 周期缓存通路L2L1 可选无片上 SRAM专用常量缓存专用纹理缓存专用只读缓存广播/合并规则合并访问128B 对齐无冲突即并行同地址广播、异地址串行空间局部性缓存空间局部性缓存可写性可读可写可读可写仅设备端只读主机可写设备端只读映射全局仅只读典型场景大数组、通用计算数据复用、block 内协作滤波器系数、查找表图像采样、插值、2D 局部性通用只读数据、const __restrict__指针这张表揭示了一个容易被忽略的事实常量内存和纹理内存的延迟并非一定优于全局内存命中——它们真正的优势在于缓存命中率和带宽利用效率。全局内存未命中时 400 周期的代价在常量/纹理缓存命中时被压缩到 200~300 周期而共享内存始终以片上 SRAM 的极低延迟独树一帜。因此选择的关键不在于哪块内存快而在于你的访问模式恰好能命中哪条缓存的规则。决策树从访问模式倒推内存类型有了对照表接下来给出可操作的决策流程。核心思路是先分析 warp 的访问模式再匹配内存类型。下图以决策树形式展示了完整的判断路径否是是否是否是否是否分析 warp 的访问模式数据是否只读全局内存或共享内存所有线程读同一地址常量内存广播机制是否需要硬件插值纹理内存纹理缓存 插值硬件数据量 ≤ 64KB 且访问地址固定不变常量内存查找表场景访问是否表现出空间局部性只读缓存 __ldg或 const __restrict__走 L2 的普通全局内存读取沿这棵决策树走一遍可以总结出三条实用准则准则一判定只读是前提。如果数据在核函数执行期间需要修改那么常量内存设备端不可写、纹理内存设备端不可写和只读缓存路径本身只读全部出局只能在全局内存和共享内存之间做选择。准则二广播优先于空间局部性。当 warp 内 32 个线程读取完全相同的地址时如卷积系数、查找表常量内存的广播机制能以一条内存指令同时满足所有线程的需求。此时即使数据量超过 64KB也应考虑拆分策略或改用纹理内存——但永远不要用共享内存去模拟广播那意味着每个线程都要从全局内存搬运一份数据到共享内存反而引入了多余的传输。准则三空间局部性指向纹理或只读缓存。当每个线程读取自己的地址、但这些地址在空间上相邻时选择取决于是否需要插值图像缩放、坐标变换等需要硬件双线性/三线性插值的场景纹理内存是唯一选择纯数据读取如科学计算中的网格数据则优先使用__ldg或const __restrict__指针利用只读缓存获得接近纹理的缓存效果同时免去纹理 API 的样板代码。一个完整的判断示例假设有一个流体力学模拟核函数需要读取一个 256×256 的压强场单精度浮点共 256KB对每个网格点计算其与上下左右四邻域的压力差。按决策树逐层分析数据是否只读是——压强场在整个核函数执行期间不变。所有线程读同一地址否——每个线程读自己周围的 4 个邻居。需要硬件插值否——压力差计算要求精确的原始值插值反而引入误差。数据量 ≤ 64KB否——256KB 远超常量内存上限。有空间局部性是——四邻域访问天然覆盖相邻地址。结论使用__ldg或const float* __restrict__指针读取压强场让只读缓存承担邻域数据的缓存任务如需进一步提升可结合共享内存做分块tiling处理。这个例子直观展示了决策树如何将抽象原则落地为具体选择。选择的核心让数据流匹配硬件规则纵观这张决策表和整篇文章可以提炼出一个贯穿始终的思想CUDA 内存选择的本质是让程序的访问模式去匹配硬件的缓存规则而非让硬件去适配程序的访问模式。常量内存的广播缓存只对同地址高效纹理缓存只对相邻地址高效共享内存只对block 内复用高效——每类内存都有自己擅长的特定访问形态。在动手写任何核函数之前先画出 warp 的访存轨迹再对照决策树选择内存类型往往能避免大量盲目的性能调优。这张决策表不仅适用于本节讨论的常量、纹理和只读缓存也同样为下一节探讨内存访问模式分析奠定了判断基准。
常量内存特性与广播机制在 CUDA 的内存层级中常量内存Constant Memory是一块常被忽视却极具特色的存储空间。它的容量只有64KB相较于全局内存动辄数 GB 的规模显得微不足道但正是这个小而精的设计让它在特定场景下能带来成倍的性能提升。理解常量内存关键在于抓住它的两个核心机制同地址广播与异地址串行。声明与数据传输__constant__与cudaMemcpyToSymbol在设备端常量内存的声明方式与全局内存类似使用__constant__限定符__constant__ float kernel_weights[128]; // 128 个 float共 512 字节 __constant__ int image_dim[3]; // 存储图像维度信息需要强调的是__constant__变量必须声明在文件作用域即所有函数体外而非某个函数的内部。这说明常量内存的生命周期贯穿整个 CUDA 上下文而非某个 kernel 的局部存在。常量内存的宿主机端访问方式与普通全局内存不同——不能直接通过cudaMemcpy复制数据。CUDA 提供了专用的 APIcudaMemcpyToSymbolfloat host_weights[128]; // ... 在宿主机端填充 host_weights ... cudaMemcpyToSymbol(kernel_weights, host_weights, sizeof(float) * 128, 0, cudaMemcpyHostToDevice); // 参数依次为设备端符号、宿主机指针、复制字节数、偏移量、复制方向同理读取常量内存中的数据使用cudaMemcpyFromSymbol方向参数改为cudaMemcpyDeviceToHost。这种独立的 API 设计反映了常量内存在编译链接层面的特殊性——它更接近一个内嵌于指令寻址空间的常量池而不是一块普通的可寻址显存。64KB 是常量内存的硬性上限。这个限制源于 NVIDIA 的硬件设计常量内存在 GPU 上配有独立的一级缓存Constant Cache缓存行宽度为 64 字节与共享内存一致。当数据量超过 64KB 时只能将多余部分存放在全局内存中由编译器在访问时通过额外的全局内存指令处理。同地址广播一个内存读取服务所有线程常量内存的性能秘密在于它的缓存架构。当 warp 中的 32 个线程同时访问同一个常量内存地址时硬件只需发起一次内存读取操作然后通过广播机制将数据分发给所有线程。这个过程完全由硬件完成不需要任何额外的指令参与。更具体的数字是NVIDIA 的常量缓存Constant Cache以 64 字节为缓存行单位带宽设计为一次读取可向 32 个线程同时广播 4 字节数据。因此一个 warp 内所有线程访问同一地址时整个操作只耗时一条内存指令的时钟周期通常约 4 个周期与寄存器访问相当。相比之下如果这些数据放在全局内存中即使 L2 缓存命中也需要至少两个线程束周期处理内存事务。与之形成鲜明对比的是全局内存的访问模式无论线程访问的是同一地址还是不同地址硬件都会将内存访问合并成尽量少的事务。若 32 个线程访问同一地址全局内存在 L2 缓存命中时也能较快响应通常约 200-300 个周期但常数缓存的广播访问在延迟上几乎可以忽略不计。将这个概念与共享内存对比更有助于建立直觉共享内存是按槽位访问——每个线程独立读取自己的槽位硬件无广播能力只有 bank 冲突时会产生串行化。常量内存是按广播访问——所有线程可以同时获得同一份数据不涉及任何冲突。因此当所有线程需要读取同一个值如卷积核系数、归一化阈值、循环上限时常量内存是比共享内存更优的选择它天然避免了共享内存的 bank 冲突问题。不过这种广播的优雅是有限度条件的——它要求所有线程访问的是完全相同的地址。一旦线程访问的地址不同瓶颈就会出现。异地址串行退化的访问模式当 warp 内不同线程访问不同地址的常量内存时性能逻辑则完全反转。常量缓存的访问流程是硬件在每个时钟周期只能处理一个内存请求并将结果广播给请求该地址的线程如果存在多个不同的地址请求硬件会将这些请求串行化——每周期只服务一个地址直至处理完所有请求。NVIDIA 官方文档给出的描述是如果 warp 中的线程访问了 n 个不同的常量内存地址则该 warp 需要n 个独立的时钟周期来完成这次访问需要可能更多视分支与数据分布。这意味着当 32 个线程全部访问不同地址时常量内存的读取速度会退化为单个内存事务的 32 倍延迟——远超全局内存合并访问的性能。一种直观的类比广播是全体整齐划一串行则是排成一列的过窄门。每次只允许一个访客通过其余线程只能等待。与全局内存合并访问相比这种串行模式在吞吐量上往往处于劣势。这个特性直接推导出一个设计准则常量内存的变量应尽量保证在分支内对齐访问。如果 kernel 依赖线程 ID 或索引动态计算常量内存下标如weights[threadIdx.x]性能会急剧下降。常见的模式是所有线程读取同一个constants[i]其中i由外部循环变量控制而非由线程 ID 直接映射。以下面两段代码作为对比// 低效模式threadIdx 索引常量内存 → 32 个地址 → 32 次串行访问 float w kernel_weights[threadIdx.x % 128]; // 高效模式所有线程访问同一地址 __shared__ float shared_w; if (threadIdx.x 0) shared_w kernel_weights[blockIdx.x % 128]; __syncthreads(); float w shared_w;第二种方式利用共享内存将常量值快速分发到所有线程消除了常量缓存的串行化惩罚。小结与选择策略常量内存的取舍可以浓缩成一句话以空间换广播以同步失速度。64KB 的容量是明确的硬约束同地址广播是实现高吞吐的首选路径适合所有线程共享系数、参数、配置等场景而异地址访问则会让常量内存变成性能陷阱必须用共享内存或全局内存替代。开发时先将常量数据按全程同址访问设计再根据实际分支结构评估是否改用其他存储空间——这样就能避免踩入大索引差异化访问的性能误区。理解了常量内存的这套广播-串行权衡之后下一个自然的问题是当我们需要更大的存储空间、更灵活的缓存策略时常量内存力不能及的场景该交给谁这正是纹理内存登场的地方。纹理内存同样拥有专用的缓存层次但它以二维局部性为核心设计哲学与常量内存的单值广播走了一条截然不同的道路。常量内存使用案例滤波器系数理解了常量内存的广播与串行机制后最自然的疑问是什么场景能真正吃到广播的红利答案是——当一个 warp 内的所有线程读取同一个地址时。这在图像处理中恰好有一个极其典型的场景卷积滤波器的系数读取。为什么滤波器系数适合放在常量内存以 3×3 的高斯模糊为例卷积核包含 9 个浮点系数。在一张 1920×1080 的图像上朴素实现为每个像素执行 9 次全局内存读取每次读一个系数总读取量为1920×1080×9≈18661920 \times 1080 \times 9 \approx 18661920×1080×9≈1866万次。这些系数对所有像素完全相同且在整个核函数执行期间只读不写——完美契合常量内存的两大前提。更关键的是访问模式。当 warp 中的 32 个线程处理相邻像素时它们在同一个循环迭代中访问的是同一个卷积核系数。假设某次循环正在计算K1,1K_{1,1}K1,1若系数存于全局内存32 个线程发出 32 次读取请求虽然 L2 缓存可能合并但首次访问必然产生多次内存事务若系数存于常量内存32 个线程读取同一地址硬件只需1 次内存读取结果通过广播总线同时送达 32 个线程。这正是第 1 节中同地址广播机制的实战价值——当过滤器系数的读取次数等于像素总数乘以核大小数百万次时广播带来的收益会被放大到极致。完整代码示例3×3 高斯模糊下面给出一个可直接编译运行的示例。它定义了一个 3×3 高斯核存入常量内存并利用广播机制加速卷积运算。#include cstdio #include cmath // 卷积核大小 #define KERNEL_RADIUS 1 #define KERNEL_SIZE (2 * KERNEL_RADIUS 1) // 3 // 图像尺寸为简化使用一维表示二维图像 #define WIDTH 64 #define HEIGHT 64 #define IMG_SIZE (WIDTH * HEIGHT) // 常量内存声明文件作用域 // 3x3 高斯核系数仅 9 个 float远远小于 64KB 限制 __constant__ float c_kernel[KERNEL_SIZE * KERNEL_SIZE]; // 设备端卷积核函数 __global__ void gaussianBlur(const float* input, float* output, int width, int height) { int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; if (x width || y height) return; float sum 0.0f; // 双重循环遍历 3x3 邻域 // 关键点同一 warp 的线程在相同迭代中读取 c_kernel 的同一地址 // → 触发常量内存的同地址广播机制 for (int dy -KERNEL_RADIUS; dy KERNEL_RADIUS; dy) { for (int dx -KERNEL_RADIUS; dx KERNEL_RADIUS; dx) { int nx x dx; int ny y dy; // 边界处理越界像素用 0 填充 float pixel 0.0f; if (nx 0 nx width ny 0 ny height) { pixel input[ny * width nx]; } // 计算线性索引并累加 int kIdx (dy KERNEL_RADIUS) * KERNEL_SIZE (dx KERNEL_RADIUS); sum pixel * c_kernel[kIdx]; } } output[y * width x] sum; } // 主机端代码 int main() { // 3x3 高斯核归一化前 float h_kernel[KERNEL_SIZE * KERNEL_SIZE] { 1, 2, 1, 2, 4, 2, 1, 2, 1 }; // 归一化使所有权重之和为 1 float sum 0.0f; for (int i 0; i KERNEL_SIZE * KERNEL_SIZE; i) sum h_kernel[i]; for (int i 0; i KERNEL_SIZE * KERNEL_SIZE; i) h_kernel[i] / sum; // 使用 cudaMemcpyToSymbol 将系数从主机复制到设备常量内存 cudaMemcpyToSymbol(c_kernel, h_kernel, sizeof(h_kernel)); // 注意cudaMemcpyToSymbol 的默认方向是 HostToDevice此处无需显式指定 // 分配图像内存 float *d_input, *d_output; cudaMalloc(d_input, IMG_SIZE * sizeof(float)); cudaMalloc(d_output, IMG_SIZE * sizeof(float)); // 构造简单的测试图像中心区域为 1.0其余为 0.0 float* h_input (float*)malloc(IMG_SIZE * sizeof(float)); for (int i 0; i IMG_SIZE; i) h_input[i] 0.0f; for (int y HEIGHT/4; y 3*HEIGHT/4; y) for (int x WIDTH/4; x 3*WIDTH/4; x) h_input[y * WIDTH x] 1.0f; cudaMemcpy(d_input, h_input, IMG_SIZE * sizeof(float), cudaMemcpyHostToDevice); // 启动核函数block 16x16grid 按需计算 dim3 block(16, 16); dim3 grid((WIDTH block.x - 1) / block.x, (HEIGHT block.y - 1) / block.y); gaussianBlurgrid, block(d_input, d_output, WIDTH, HEIGHT); // 核对错误 cudaError_t err cudaDeviceSynchronize(); if (err ! cudaSuccess) { printf(CUDA error: %s\n, cudaGetErrorString(err)); return -1; } // 拷贝结果回主机并验证核心区域 float* h_output (float*)malloc(IMG_SIZE * sizeof(float)); cudaMemcpy(h_output, d_output, IMG_SIZE * sizeof(float), cudaMemcpyDeviceToHost); // 验证中心区域 (WIDTH/2, HEIGHT/2) 的模糊结果 int cx WIDTH / 2, cy HEIGHT / 2; printf(Center pixel after blur: %f\n, h_output[cy * WIDTH cx]); // 预期输出中心区域被周围 0 值像素稀释结果在 0.3~0.6 之间 // 清理 free(h_input); free(h_output); cudaFree(d_input); cudaFree(d_output); return 0; }编译指令nvcc -o blur blur.cu。运行后中心像素的输出值约为0.44由中心 4、四个正交邻域 2×4/个、四个对角邻域 1×4/个加权叠加后被 16 归一化得到。代码要点拆解__constant__声明必须位于文件作用域。第 1 节提到常量内存的声明要求这里再次验证c_kernel被定义为全局数组因为常量内存在 CUDA 中不支持动态分配且所有核函数共享同一份常量内存。9 个 float 仅占 36 字节远低于 64KB 上限。cudaMemcpyToSymbol是唯一的人口。它负责将主机的h_kernel拷贝到设备端的c_kernel。注意不需要在参数中写cudaMemcpyHostToDevice因为cudaMemcpyToSymbol从主机到设备是默认方向。若要从设备读回常量内存则需显式使用cudaMemcpyDeviceToHost。核函数内层循环是广播的核心。当dx0, dy0时所有线程都在读c_kernel[4]。此时硬件检测到 warp 内 32 个线程请求同一地址触发单次内存读取加广播。整个核函数执行过程中c_kernel[kIdx]被读取的次数等于线程总数 × 9 次但全部命中广播路径全局内存流量几乎为零。用共享内存做对比何时常量内存更优有人会问把卷积核放进共享内存不是一样的吗确实共享内存也有高速缓存特性。但两者有本质差异对比维度常量内存共享内存访问模式同一地址→广播任意地址→按槽位访问存储上限64KB48KB可配至 96KB数据更新核函数启动前一次性赋值核函数内可动态写入硬件开销广播总线单次传输需同步 按槽位冲突管理如果滤波器系数在核函数运行过程中需要动态更新比如自适应滤波共享内存是更好的选择。但对于静态系数的传统卷积常量内存的广播机制让所有线程在同一个循环迭代中只产生一次内存事务——这种零冲突特性是共享内存按槽位分配的访问方式所不具备的不同槽位可能产生 bank conflict。因此在卷积核系数固定且 warp 内线程步调一致时常量内存是更精确的选择。由此可以提炼出一个普适的决策准则当一个 warp 内的所有线程在同一个指令周期内读取同一地址的只读数据时常量内存就是最优解。卷积核系数、查找表LUT、数学函数的预计算参数都属于这类一读多用的数据。而接下来要讨论的纹理内存则把这种特殊缓存路径的思路推进到了更广阔的维度——它不仅提供缓存还内置了硬件插值、归一化坐标等图像处理专属能力。纹理内存与硬件纹理单位如果说常量内存通过广播特性解决了同一地址的高频访问问题那么纹理内存Texture Memory则通过缓存特性解决了另一类截然不同的访问模式——相邻地址的密集读取。这正是图像处理中像素采样的典型形态一个像素的邻域数据往往在下一次计算中立即被再次使用。常量内存管的是所有人读同一个数纹理内存管的是每个人读自己附近的那片数。空间局部性纹理缓存的底层逻辑图像处理的核心计算模式是邻域操作均值模糊取 3×3 邻域边缘检测取 Sobel 核覆盖的区域双线性插值取周围 4 个像素。这类操作有一个天然的访存特征——空间局部性Spatial Locality当前线程访问的像素附近极大概率是相邻线程下一个要访问的目标。全局内存的缓存行为并不理想它的缓存行通常 128 字节虽然能覆盖相邻地址但全局内存缓存更多是为通用数据复用设计缺乏对二维空间结构的感知。纹理内存则专门针对这一模式做了硬件优化——它内部实现了专用的纹理缓存该缓存针对二维甚至三维空间的遍历模式进行了调优。当一个线程读取了坐标(x,y)(x, y)(x,y)处的像素硬件会预取(x1,y)(x1, y)(x1,y)、(x,y1)(x, y1)(x,y1)附近的纹素texel到缓存中。相邻线程访问相邻坐标时命中率极高。以一张 1920×1080 的灰度图像为例处理 3×3 模糊时全局内存方式每个像素需要 9 次独立访问而纹理内存方式下第一次读取(x,y)(x, y)(x,y)时周围的纹素已驻留缓存后续 8 个邻域坐标的访问直接命中纹理缓存实际的内存流量降低了数倍。缓存未命中率从全局内存的数个数量级之高的循环重复访问降至几近于零。// 纹理内存的申请与绑定 texturefloat, 2, cudaReadModeElementType texRef; // 声明 2D 纹理引用 cudaChannelFormatDesc desc cudaCreateChannelDescfloat(); cudaArray* cuArray; cudaMallocArray(cuArray, desc, width, height); // 分配二维数组 // 将图像数据拷贝到 CUDA 数组纹理只能绑定到 CUDA 数组 cudaMemcpyToArray(cuArray, 0, 0, h_data, width * height * sizeof(float), cudaMemcpyHostToDevice); cudaBindTextureToArray(texRef, cuArray); // 完成绑定设备端通过tex2D(texRef, x, y)进行采样由硬件纹理单位Texture Unit完成地址计算和值的返回。纹理硬件Texture Unit是 GPU 上独立于通用计算核心的专用电路它负责坐标换算、缓存管理和数据过滤——这些操作不占用 CUDA 核心的计算资源而是由纹理单元的专用逻辑并行完成。硬件插值让双线性采样近乎免费纹理内存最令人惊艳的特性是硬件插值Hardware Interpolation。在图像缩放、旋转、坐标变换等场景中目标像素往往落在源图像的非整数坐标上。传统的计算方式需要软件实现双线性插值取周围 4 个像素、计算权重、加权求和。这需要约 4 次内存访问和多次浮点运算。纹理硬件将这一过程封装为单条指令。当使用归一化坐标Normalized Coordinates采样时——坐标范围映射到[0,1)[0, 1)[0,1)——硬件会自动处理坐标到纹素地址的换算并通过cudaFilterModeLinear模式直接在硬件层面完成双线性插值texturefloat, 2, cudaReadModeElementType texRef; texRef.addressMode[0] cudaAddressModeWrap; // 坐标环绕模式 texRef.addressMode[1] cudaAddressModeWrap; texRef.filterMode cudaFilterModeLinear; // 硬件双线性插值 // 采样时直接传归一化坐标0.0 ~ 1.0 之间 float val tex2D(texRef, nx, ny); // nx, ny ∈ [0, 1)这组代码中tex2D的返回值就是已经完成双线性插值的结果。软件实现需要至少 4 次全局/共享内存读取加 4 次乘加运算硬件方案只需 1 次纹理读取指令。对于大规模图像变换这种差异意味着数倍的吞吐提升——并且像素越连续、坐标越密集插值硬件的工作效率越高因为它内部有专用的插值单元流水线。2D/3D 局部性的完整收益除了插值纹理内存的缓存设计也超越了简单的一维缓存。2D 纹理缓存针对行、列两个维度的访问都做了预取优化3D 纹理缓存则进一步覆盖深度维度的局部性。这意味着体数据如医学 CT 切片、气候模拟网格的访问也能获得相似的红利。从整个 CUDA 内存体系来看纹理内存的价值定位是清晰而互补的对比维度全局内存共享内存常量内存纹理内存适用访问模式任意读写同 block 内协作warp 内同地址广播相邻地址密集读取板上缓存L2无显式管理常量缓存专用纹理缓存插值支持无无无硬件级插值边界处理无手动无Wrap/Clamp 模式理解了这个坐标就能回答一个常见困惑为什么有了共享内存还要纹理内存共享内存需要程序员手动管理数据搬运和复用而纹理缓存是硬件自动维护的。当你不想显式管理块边界、但又需要利用空间局部性时纹理内存是更轻量的选择——你只管采样缓存策略交给硬件。常量内存靠广播解决了所有线程读同一地址的效率问题纹理内存靠专用的缓存硬件和插值电路解决了相邻线程读相邻地址的效率问题。这两种内存共同补齐了 CUDA 内存体系中读密集、规律性访问的短板——接下来的问题很自然纹理和常量之外还有一类适用于散列访问的只读缓存机制它与纹理内存有何本质区别这正是下一节要展开的内容。只读缓存与__ldg再讨论常量内存的广播机制解决了所有人读同一个数的问题纹理内存的缓存与插值硬件解决了每个人读自己附近那片数的问题。但这两者都不是通用方案常量内存受限于 64KB 容量纹理内存则需要引入纹理坐标、采样器等一套繁琐的 API。那么当数据量超过常量内存上限、访问模式又不适合纹理内存时还有没有第三条路答案是肯定的——只读缓存Read-Only Cache。它是现代 GPUKepler 架构及以后内置的一条独立缓存通路专门服务于只读访问全局内存的场景。只读缓存绕过 L1 的独立通路在 Fermi 时代所有全局内存访问都要经过统一的 L1 缓存——如果你用const __restrict__修饰指针编译器会尝试通过 L1 缓存路径来读取数据但这并不保证只读语义被充分利用。Kepler 架构引入了专门的数据缓存通路即只读缓存它与 L1 数据缓存并行存在两者都有各自独立的能力。在编译器的眼中一个被标记为只读的全局内存访问会走只读缓存——这带来的收益在于只读缓存拥有比 L1 数据缓存更高的带宽只读缓存支持全速率的 128 字节/周期的读取而 L1 数据缓存则有所保留只读缓存的命中延迟与常量和纹理缓存相当几十个周期量级但容量明显大于常量内存从编程接口的角度让代码走只读缓存有两种方式// 方式一显式使用 __ldg 内置函数 __device__ float kernel_explicit(const float* data, int idx) { return __ldg(data idx); // 显式告知编译器这是一次只读访问 } // 方式二使用 const __restrict__ 修饰指针让编译器自动生成 LDG 指令 __global__ void kernel_auto(const float* __restrict__ data, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) out[i] data[i] * 2.0f; // 编译器自动识别为只读访问 }需要说明的是__ldg并非凭空发明的新功能——它只是强制编译器把某个访问编译为 LDGLoad Global指令。而const __restrict__则给了编译器一个证明“这块数据在核函数执行期间不会修改”编译器从而可以进行自动优化。两行代码看起来相似本质不同__ldg是命令__restrict__是承诺。在引入__ldg时NVIDIA 官方给出的性能数据显示当只读访问存在空间局部性即相邻线程访问相邻地址时只读缓存可以带来约 1.7 倍的带宽提升特别是与没有走只读通路的版本相比。这与我们说纹理缓存命中空间局部性的逻辑一脉相承只不过纹理缓存需要坐标而只读缓存只是简单地缓存 32 字节的缓存行天然适合线性数组的遍历。适用差异三大特殊内存如何选现在我们对 CUDA 中三种带缓存的内存通路已经有了全景视野。它们各自的物理基础不同、适用场景各异但核心区别可以归结为一个问题一个 warp 里的 32 个线程访问的地址分布是什么样的为了便于决策我们可以用一张表来总结内存通路最佳访问模式容量限制硬件插值编程方式常量内存所有线程读同一地址广播64KB无__constant__cudaMemcpyToSymbol纹理内存相邻线程读相邻地址局部性可达数十 GB绑定全局内存有双线性/三线性插值纹理对象 tex2D/tex3D只读缓存相邻线程读相邻地址局部性无缓存全局内存无__ldg()或const __restrict__更具体地看常量内存的不可替代性。当 warp 内所有线程必须读取完全相同的地址时如卷积核系数、数学查表广播机制将一次内存访问广播给所有线程吞吐量为一次访问的代价。任何其他方式——包括只读缓存——都无法替代这一点因为它们的缓存粒度是缓存行如 32 字节当一个 warp 读取同一地址时整条缓存行只有一次访问被利用其余存储空间完全浪费。纹理内存的独特优势在于插值。双线性/三线性插值并非一个可以在通用计算中随意简化掉的特性。图像缩放、体渲染、光学仿真等场景中采样坐标往往落在像素网格之外——使用纹理内存硬件直接返回插值结果避免了在代码中层叠 4 次2D甚至 8 次3D的混合计算。若这些操作在通用路径上实现每一层插值都会消耗若干条乘加指令和若干次内存访问。只读缓存则是通用领域的均衡解。数据量超过 64KB、不需要插值、但访问模式有空间局部性——这是最普遍的情况。以矩阵转置为例线程 i 按行步长读取矩阵 A 的若干行每个 warp 内 32 个线程读取的地址彼此相差 4 字节的倍数——正好命中只读缓存的缓存行粒度。这种情况下__ldg是唯一不需要改变数据布局、不需要引入纹理坐标、只需一行代码改动就能获得缓存增益的方案。决策框架从访问模式出发如果以一句话总结三类通路的取舍逻辑可以沿着如下思路来判断先看是否需要插值—— 如果需要硬件插值直接选择纹理内存无其他替代再问数据量是否在 64KB 以内—— 是且 warp 内所有线程访问同一地址选择常量内存剩余情况—— 数据量大、只读、有空间局部性使用__ldg或const __restrict__让编译器走只读缓存这个框架并不替代实测——CUDA 的优化空间往往超出直觉的线性推断。但它的意义在于当我们面对一个新的核函数时不再被用纹理内存总是更好或常量内存万能这类朴素信念所误导而是带着清晰的访问模式图景去选择最匹配的硬件通路。至此从共享内存到常量内存从纹理内存到只读缓存CUDA 自有存储体系的全貌已经完整呈现——而这四条通路之间的协同与博弈将在下一节以完整的决策树形式展开总结。内存类型选择决策表至此我们已经走完了 CUDA 内存体系的最后一公里常量内存用 64KB 的容量换来了同地址广播的高效通路纹理内存借由专用缓存与插值硬件统治了邻域读取场景只读缓存则为超出这两种特殊约束的通用访问提供了第三条路。但有一个问题始终悬而未决——面对一个具体的核函数到底该选哪种内存本节将把这四类内存放在同一张决策表中给出一个可复用的选择框架。五类内存的一次性对照在给出决策流程之前先建立全局视角。下表从容量、延迟、缓存机制、适用访问模式、典型场景五个维度对全局内存、共享内存、常量内存、纹理内存和只读缓存做一次系统性对照维度全局内存共享内存常量内存纹理内存只读缓存容量数 GB48KB~228KB按架构64KB数 GB映射全局数 GB映射全局延迟命中400~600 周期约 20~30 周期约 200~300 周期约 200~300 周期约 200~300 周期缓存通路L2L1 可选无片上 SRAM专用常量缓存专用纹理缓存专用只读缓存广播/合并规则合并访问128B 对齐无冲突即并行同地址广播、异地址串行空间局部性缓存空间局部性缓存可写性可读可写可读可写仅设备端只读主机可写设备端只读映射全局仅只读典型场景大数组、通用计算数据复用、block 内协作滤波器系数、查找表图像采样、插值、2D 局部性通用只读数据、const __restrict__指针这张表揭示了一个容易被忽略的事实常量内存和纹理内存的延迟并非一定优于全局内存命中——它们真正的优势在于缓存命中率和带宽利用效率。全局内存未命中时 400 周期的代价在常量/纹理缓存命中时被压缩到 200~300 周期而共享内存始终以片上 SRAM 的极低延迟独树一帜。因此选择的关键不在于哪块内存快而在于你的访问模式恰好能命中哪条缓存的规则。决策树从访问模式倒推内存类型有了对照表接下来给出可操作的决策流程。核心思路是先分析 warp 的访问模式再匹配内存类型。下图以决策树形式展示了完整的判断路径否是是否是否是否是否分析 warp 的访问模式数据是否只读全局内存或共享内存所有线程读同一地址常量内存广播机制是否需要硬件插值纹理内存纹理缓存 插值硬件数据量 ≤ 64KB 且访问地址固定不变常量内存查找表场景访问是否表现出空间局部性只读缓存 __ldg或 const __restrict__走 L2 的普通全局内存读取沿这棵决策树走一遍可以总结出三条实用准则准则一判定只读是前提。如果数据在核函数执行期间需要修改那么常量内存设备端不可写、纹理内存设备端不可写和只读缓存路径本身只读全部出局只能在全局内存和共享内存之间做选择。准则二广播优先于空间局部性。当 warp 内 32 个线程读取完全相同的地址时如卷积系数、查找表常量内存的广播机制能以一条内存指令同时满足所有线程的需求。此时即使数据量超过 64KB也应考虑拆分策略或改用纹理内存——但永远不要用共享内存去模拟广播那意味着每个线程都要从全局内存搬运一份数据到共享内存反而引入了多余的传输。准则三空间局部性指向纹理或只读缓存。当每个线程读取自己的地址、但这些地址在空间上相邻时选择取决于是否需要插值图像缩放、坐标变换等需要硬件双线性/三线性插值的场景纹理内存是唯一选择纯数据读取如科学计算中的网格数据则优先使用__ldg或const __restrict__指针利用只读缓存获得接近纹理的缓存效果同时免去纹理 API 的样板代码。一个完整的判断示例假设有一个流体力学模拟核函数需要读取一个 256×256 的压强场单精度浮点共 256KB对每个网格点计算其与上下左右四邻域的压力差。按决策树逐层分析数据是否只读是——压强场在整个核函数执行期间不变。所有线程读同一地址否——每个线程读自己周围的 4 个邻居。需要硬件插值否——压力差计算要求精确的原始值插值反而引入误差。数据量 ≤ 64KB否——256KB 远超常量内存上限。有空间局部性是——四邻域访问天然覆盖相邻地址。结论使用__ldg或const float* __restrict__指针读取压强场让只读缓存承担邻域数据的缓存任务如需进一步提升可结合共享内存做分块tiling处理。这个例子直观展示了决策树如何将抽象原则落地为具体选择。选择的核心让数据流匹配硬件规则纵观这张决策表和整篇文章可以提炼出一个贯穿始终的思想CUDA 内存选择的本质是让程序的访问模式去匹配硬件的缓存规则而非让硬件去适配程序的访问模式。常量内存的广播缓存只对同地址高效纹理缓存只对相邻地址高效共享内存只对block 内复用高效——每类内存都有自己擅长的特定访问形态。在动手写任何核函数之前先画出 warp 的访存轨迹再对照决策树选择内存类型往往能避免大量盲目的性能调优。这张决策表不仅适用于本节讨论的常量、纹理和只读缓存也同样为下一节探讨内存访问模式分析奠定了判断基准。
平方复杂度无法被常规优化消除｜TaoToken 统一 Key 通道实测)
