在Kepler硬件上的Visual Profiler中分析我的内核时,我注意到分析器显示全局加载和存储缓存在L1中.我很困惑,因为编程指南和开普勒调音手册指出:
Kepler GPU中的L1缓存仅保留用于本地存储器访问,例如寄存器溢出和堆栈数据.全局加载仅缓存在L2中(或在只读数据缓存中).
没有寄存器溢出(探查器显示L1缓存,即使是原始的,2行'添加'内核),我不知道'堆栈数据'在这里意味着什么.
GK110白皮书显示除了一种情况外,全局访问将通过L1缓存:通过只读缓存(__ldg)加载.这是否意味着当全局访问通过L1硬件时,它们实际上并未缓存?这是否也意味着如果我在L1中缓存了溢出的寄存器数据,那么这些数据可能会因为访问gmem而被驱逐?
更新:我意识到我可能误读了分析器给我的信息,所以这里是内核代码以及分析器结果(我在Titan和K40上都试过了相同的结果).
template__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]; /* 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; }
Greg Smith.. 5
在Fermi和Kepler架构上,所有通用,全局,本地和共享内存操作都由L1缓存处理.共享内存访问不需要查找标记,也不会使缓存行无效.所有本地和全局内存访问都需要查找标记.未缓存的全局内存存储和读取将使缓存行无效.在计算能力3.0和3.5上,除了CC 3.5上的LDG之外的所有全局内存读取都将被解除.LDG指令通过纹理缓存.
在Fermi和Kepler架构上,所有通用,全局,本地和共享内存操作都由L1缓存处理.共享内存访问不需要查找标记,也不会使缓存行无效.所有本地和全局内存访问都需要查找标记.未缓存的全局内存存储和读取将使缓存行无效.在计算能力3.0和3.5上,除了CC 3.5上的LDG之外的所有全局内存读取都将被解除.LDG指令通过纹理缓存.