Jetson Nano（Maxwell架构）启用L1/Texture缓存对NVCC（CUDA10.2）的性能影响

实验背景 #

为对比纹理内存与全局内存在双线性插值中的表现，实现了两个CUDA内核：

全局内存双线性插值核 #

__global__ void GlobalBilinear(typeImg * inPtr, typeOutImg* outPtr, int width, int height)
{
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;

    __shared__ float val[1024];

    if(ix < width && iy < height)
    {
        val[threadIdx.x + blockDim.x*threadIdx.y] 
                =   (1.0f - alpha) * (1.0f - beta) * uint2float(inPtr[ix + width * iy]) + 
                    alpha * (1.0f - beta) * uint2float(inPtr[ix + 1 + width * iy])      +
                    (1.0f - alpha) * (beta)*uint2float(inPtr[ix + width * (iy + 1)])     +
                    alpha * (beta)*uint2float(inPtr[ix + 1 + width * (iy + 1)]);
    }
}

纹理内存双线性插值核 #

__global__ void MemoryBilinear2DTexture(cudaTextureObject_t texObj, typeOutImg* outPtr, int width, int height)
{
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;

    __shared__ float val[1024];
    if(ix < width && iy < height)
    {
        val[threadIdx.x + blockDim.x*threadIdx.y] = (tex2D<typeOutImg>(texObj, ix + 1.0f, iy + 1.0f));
    }
}

测试结果 #

对1000×1000的8位图像进行50次运行计时：

• 默认编译（L1缓存禁用，-Xptxas -dlcm=cg）：
  • 全局内存核耗时：37.87073 ms
  • 纹理内存核耗时：25.75578 ms
• 手动启用L1缓存（-Xptxas -dlcm=ca）：
  • 全局内存核耗时：48.90469 ms
  • 纹理内存核耗时：33.65562 ms

L1启用时的nvprof分析结果 #

**Kernel: GlobalBilinear(void*, void*, int, long)**
global_hit_rate         Global Hit Rate in unified l1/tex      72.48%      72.48%      72.48%
tex_cache_transactions  Unified Cache Transactions      504000      504000      504000
**Kernel: MemoryBilinear2DTexture(__int64, void*, int, long)**
global_hit_rate         Global Hit Rate in unified l1/tex       0.00%       0.00%       0.00%
tex_cache_transactions  Unified Cache Transactions      250000      250000      250000

问题解答 #

1. 为何启用L1缓存会带来性能下降？ #

从nvprof数据可以拆解出核心原因：

• 全局内存核的双线性实现中，每个线程要读取4个相邻全局内存地址。启用L1后，这些访问会占用L1缓存空间，相邻线程的缓存行容易出现覆盖、冲突，导致额外的缓存读写开销；同时L1与纹理缓存共享硬件资源，启用L1会挤占纹理缓存的可用空间，直接降低纹理核的访问效率。
• 默认禁用L1时，全局内存访问直接走L2缓存。对于这种规则的连续+相邻访问模式，L2的预取机制和带宽表现更优——L1容量小，容易触发缓存失效，反而增加了往返L2的次数，整体延迟上升。

2. 为何NVCC默认禁用全局内存访问的L1缓存？ #

NVCC的默认设置是基于通用场景的最优权衡：

• 多数通用CUDA内核的全局内存访问要么是规则连续访问（此时L2预取机制已能高效工作，L1增益有限），要么是随机访问（L1容量小，命中率低，反而增加缓存开销）。
• 在统一缓存架构中，L1与纹理缓存共享资源。禁用L1可让更多资源留给纹理缓存，这对图形、图像处理这类CUDA常见应用场景更友好。
• 启用L1可能引发缓存冲突、额外延迟等“意外”性能损耗，默认禁用能让多数通用内核的表现更稳定。

内容的提问来源于stack exchange，提问作者rooody