cuda - Kepler 中的全局内存访问和 L1 缓存

Question

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

L1 caching in Kepler GPUs is reserved only for local memory accesses, such as register spills and stack data. Global loads are cached in L2 only (or in the Read-Only Data Cache).

没有寄存器溢出(即使对于原始的 2 行“添加”内核，分析器也显示 L1 缓存)，我不确定这里的“堆栈数据”是什么意思。

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);

视觉探查器输出:



鉴于为 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];                /* 0x64c03c00089c0006 */
/*0010*/                S2R R0, SR_CTAID.X;                  /* 0x86400000129c0002 */
/*0018*/                MOV32I R5, 0x4;                      /* 0x74000000021fc016 */
/*0020*/                S2R R3, SR_TID.X;                    /* 0x86400000109c000e */
/*0028*/                IMAD R2, R0, c[0x0][0x28], R3;       /* 0x51080c00051c000a */
/*0030*/                IMAD R6.CC, R2, R5, c[0x0][0x148];   /* 0x910c1400291c081a */
/*0038*/                IMAD.HI.X R7, R2, R5, c[0x0][0x14c]; /* 0x93181400299c081e */
                                                             /* 0x08a0a4b0809c80b0 */
/*0048*/                IMAD R8.CC, R2, R5, c[0x0][0x150];   /* 0x910c14002a1c0822 */
/*0050*/                IMAD.HI.X R9, R2, R5, c[0x0][0x154]; /* 0x931814002a9c0826 */
/*0058*/                LD.E R3, [R6];                       /* 0xc4800000001c180c */
/*0060*/                LD.E R0, [R8];                       /* 0xc4800000001c2000 */
/*0068*/                IMAD R4.CC, R2, R5, c[0x0][0x140];   /* 0x910c1400281c0812 */
/*0070*/                IMAD.HI.X R5, R2, R5, c[0x0][0x144]; /* 0x93181400289c0816 */
/*0078*/                FADD R0, R3, R0;                     /* 0xe2c00000001c0c02 */
                                                             /* 0x080000000000b810 */
/*0088*/                ST.E [R4], R0;                       /* 0xe4800000001c1000 */
/*0090*/                EXIT ;                               /* 0x18000000001c003c */
/*0098*/                BRA 0x98;                            /* 0x12007ffffc1c003c */
/*00a0*/                NOP;                                 /* 0x85800000001c3c02 */
/*00a8*/                NOP;                                 /* 0x85800000001c3c02 */
/*00b0*/                NOP;                                 /* 0x85800000001c3c02 */
/*00b8*/                NOP;                                 /* 0x85800000001c3c02 */

PTX:

.visible .entry _Z9addKernelIfEvPT_PKS0_S3_(
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_0,
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_1,
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_2
)
{
.reg .s32 %r<5>;
.reg .f32 %f<4>;
.reg .s64 %rd<11>;

ld.param.u64 %rd1, [_Z9addKernelIfEvPT_PKS0_S3__param_0];
ld.param.u64 %rd2, [_Z9addKernelIfEvPT_PKS0_S3__param_1];
ld.param.u64 %rd3, [_Z9addKernelIfEvPT_PKS0_S3__param_2];
cvta.to.global.u64 %rd4, %rd1;
.loc 1 22 1
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r1, %r2, %r3;
cvta.to.global.u64 %rd5, %rd2;
mul.wide.s32 %rd6, %r4, 4;
add.s64 %rd7, %rd5, %rd6;
cvta.to.global.u64 %rd8, %rd3;
add.s64 %rd9, %rd8, %rd6;
.loc 1 23 1
ld.global.f32 %f1, [%rd9];
ld.global.f32 %f2, [%rd7];
add.f32 %f3, %f2, %f1;
add.s64 %rd10, %rd4, %rd6;
.loc 1 23 1
st.global.f32 [%rd10], %f3;
.loc 1 24 2
ret;
}

Answer 1

在 Fermi 和 Kepler 架构上，所有通用、全局、本地和共享内存操作都由 L1 缓存处理。共享内存访问不需要标记查找并且不会使缓存行无效。所有本地和全局内存访问都需要进行标记查找。未缓存的全局内存存储和读取将使缓存行无效。在计算能力 3.0 和 3.5 上，除 CC 3.5 上的 LDG 外，所有全局内存读取都将不被缓存。 LDG 指令通过纹理缓存。