cuda - problemas - memoria cache pdf

Acceso a memoria global y caché L1 en Kepler

Question

Al perfilar mis núcleos en Visual Profiler en hardware Kepler, me he dado cuenta de que el generador de perfiles muestra que las cargas globales y las tiendas están en caché en L1. Estoy confundido porque la guía de programación y el manual de ajuste de Kepler indican que:

El almacenamiento en caché L1 en las GPU de Kepler está reservado solo para los accesos de memoria local, como los derrames de registros y los datos de la pila. Las cargas globales se almacenan en caché solo en L2 (o en el caché de datos de solo lectura).

No hay ningún derrame de registros (Profiler muestra el almacenamiento en caché de L1 incluso para el kernel primitivo de "agregar" 2 líneas) y no estoy seguro de qué significa "datos de pila" aquí.

El Whitepaper GK110 muestra que los accesos globales pasarán por la caché L1 en todos los casos menos uno: se carga a través de caché de solo lectura (__ldg). ¿Significa que, si bien los accesos globales pasan por el hardware L1, en realidad no están en la memoria caché? ¿Significa también que si he derramado los datos de los registros almacenados en caché en L1, estos datos pueden ser desalojados como resultado del acceso gmem?

ACTUALIZACIÓN : me he dado cuenta de que podría estar malinterpretando la información que el generador de perfiles me está dando, así que aquí está el código del núcleo y los resultados del generador de perfiles (lo he intentado tanto en Titan como en K40 con los mismos resultados).

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

Salida Visual Profiler:

Los números L1 tienen mucho sentido dado que la memoria caché L1 está habilitada para los accesos gmem. Para las cargas que tenemos:

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

ACTUALIZACIÓN 2 : código SASS y PTX agregado. El código SASS es muy simple y contiene lecturas de memoria constante y cargas / almacena desde / a la memoria global (instrucciones 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

En las arquitecturas de Fermi y Kepler, todas las operaciones genéricas, globales, locales y de memoria compartida son manejadas por la memoria caché L1. Los accesos a memoria compartida no requieren una búsqueda de etiqueta y no invalidan una línea de caché. Todos los accesos de memoria locales y globales requieren una búsqueda de etiqueta. Los almacenes y lecturas de memoria global no almacenados en caché invalidarán una línea de caché. En la capacidad de cálculo 3.0 y 3.5, todas las lecturas de memoria global con excepción de LDG en CC 3.5 se descartarán. La instrucción LDG pasa por la memoria caché.