小编Ale*_*nev的帖子

在Kepler硬件上的Visual Profiler中分析我的内核时,我注意到分析器显示全局加载和存储缓存在L1中.我很困惑,因为编程指南和开普勒调音手册指出:

Kepler GPU中的L1缓存仅保留用于本地存储器访问,例如寄存器溢出和堆栈数据.全局加载仅缓存在L2中(或在只读数据缓存中).

没有寄存器溢出(探查器显示L1缓存,即使是原始的,2行'添加'内核),我不知道'堆栈数据'在这里意味着什么.

GK110白皮书显示除了一种情况外,全局访问将通过L1缓存:通过只读缓存(__ldg)加载.这是否意味着当全局访问通过L1硬件时,它们实际上并未缓存？这是否也意味着如果我在L1中缓存了溢出的寄存器数据,那么这些数据可能会因为访问gmem而被驱逐？

更新:我意识到我可能误读了分析器给我的信息,所以这里是内核代码以及分析器结果(我在Titan和K40上都试过了相同的结果).

template<typename T>
__global__ void addKernel(T *c, const T *a, const T *b)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

...
// Kernel call
float* x;
float* y;
float* d;
// ...
addKernel<<<1024, 1024>>>(d, x, y);
cudaError_t cudaStatus = cudaDeviceSynchronize();
assert(cudaSuccess == cudaStatus);

Visual Profiler输出:

在为gmem访问启用L1缓存的情况下,L1数字非常有意义.对于我们的负载:

65536*128 == 2*4*1024*1024

更新2:添加了SASS和PTX代码.SASS代码非常简单,包含来自常量存储器的读取以及来自/到全局存储器的加载/存储(LD/ST指令).

Function : _Z9addKernelIfEvPT_PKS0_S3_
.headerflags    @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
                                                             /* 0x088cb0a0a08c1000 */
/*0008*/                MOV R1, c[0x0][0x44];                /* …