CUDA deshabilitar la caché L1 solo para una variable

¿Hay alguna forma en los dispositivos CUDA 2.0 para deshabilitar la caché L1 solo para una variable específica? Sé que se puede deshabilitar la memoria caché L1 en tiempo de compilación agregando la marca -Xptxas -dlcm=cg a nvcc para todas las operaciones de memoria. Sin embargo, quiero desactivar la memoria caché solo para lecturas de memoria en una variable global específica para que todo el resto de la memoria se lea en la caché L1.CUDA deshabilitar la caché L1 solo para una variable

Basado en una búsqueda que he realizado en la web, una posible solución es a través del código de ensamblaje PTX.

Fuente

2012-09-23 zeus2

como se mencionó anteriormente se puede utilizar en línea PTX, aquí es un ejemplo:

__device__ __inline__ double ld_gbl_cg(const double *addr) { 
    double return_value; 
    asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr)); 
    return return_value; 
}

Usted fácilmente puede variar mediante el canje de esta .f64 para .f32 (float) o .s32 (int), etc., el constraint of return_value "= d" para "= f" (float) o "= r" (int) etc. Tenga en cuenta que la última restricción antes (addr) - "l" - denota el direccionamiento de 64 bits, si está utilizando el direccionamiento de 32 bits, debe ser "r".

Fuente

2012-09-23 20:39:48 Reguj

¡Gracias! Eso está funcionando bien! – zeus2

@Reguj, ¿esto no es proporcionado por los encabezados de NVIDIA en cualquier lugar? – einpoklum

[este] (https://nvlabs.github.io/cub/classcub_1_1_cache_modified_input_iterator.html#details) puede ser de interés –

Inline PTX se puede utilizar para cargar y almacenar la variable. Las instrucciones ld.cg y st.cg solo almacenan datos en caché en L2. Los operadores de caché se describen en la sección 8.7.8.1 Operadores de caché del documento PTX ISA 2.3. Las instrucciones o intereses son ld y st. PTX en línea se describe en Using Inline PTX Assembly in CUDA.

Fuente

2012-09-23 16:56:22

Si declara que la variable es volatile, solo se almacenará en caché en la memoria caché L2 en las GPU de Fermi. Tenga en cuenta que algunas optimizaciones del compilador, como la eliminación de cargas repetidas, no se realizan en variables volátiles porque el compilador supone que pueden ser escritas por otro hilo.

Fuente

2012-09-23 17:59:02 Heatsink

No creo que el modelo de programación haga ninguna representación sobre la caché de las variables volátiles. – ArchaeaSoftware

@Archaea La arquitectura de Fermi hace que el almacenamiento en caché de datos volátiles no sea factible debido a la ausencia de coherencia del caché. Habiendo encontrado errores en la documentación de CUDA en el pasado, no considero que la documentación del modelo de memoria de CUDA sea confiable. – Heatsink

Intenté la solución con declinación variable volátil y no funcionó. Parece que la variable se almacena en caché de nuevo. – zeus2

CUDA deshabilitar la caché L1 solo para una variable

Respuesta

Cuestiones relacionadas