arrays - En un kernel CUDA, ¿cómo almaceno una matriz en la "memoria de subproceso local"?

memory gpu-local-memory

Question

Estoy tratando de desarrollar un pequeño programa con CUDA, pero como era LENTO, hice algunas pruebas y busqué un poco en Google. Descubrí que, si bien las variables individuales están almacenadas de forma predeterminada en la memoria del subproceso local, las matrices generalmente no lo están. Supongo que por eso lleva tanto tiempo ejecutarlo. Ahora me pregunto: ya que la memoria de subprocesos local debe ser de al menos 16KB y como mis arreglos tienen una longitud de 52 caracteres, ¿hay alguna forma (sintaxis, por favor :)) para almacenarlos en la memoria local?

¿No debería ser algo como:

__global__ my_kernel(int a)
{
  __local__ unsigned char p[50];
}

Answer 1

Arreglos, memorias locales y registros.

Hay una idea errónea aquí con respecto a la definición de "memoria local". La "memoria local" en CUDA es en realidad memoria global (y en realidad debería llamarse "memoria global local de subproceso") con direccionamiento intercalado (lo que hace que la iteración en una matriz en paralelo sea un poco más rápida que tener los datos de cada subproceso bloqueados). Si desea que las cosas sean realmente rápidas, desea utilizar los registros de memoria compartida o, mejor aún, los registros (especialmente en los dispositivos más recientes donde obtiene hasta 255 registros por subproceso). Explicar la jerarquía de memoria CUDA completa queda fuera del alcance de esta publicación. En cambio, centrémonos en hacer que los cálculos de matriz pequeña sean más rápidos.

Las matrices pequeñas, al igual que las variables, pueden almacenarse completamente en registros. Sin embargo, en el hardware NVIDIA actual, es difícil colocar matrices en registros. ¿Por qué? Porque los registros necesitan un tratamiento muy cuidadoso. Si no lo hace exactamente bien, sus datos terminarán en la memoria local (que, de nuevo, es realmente una memoria global, que es la memoria más lenta que tiene). La Guía de programación de CUDA, sección 5.3.2 le informa cuándo se utiliza la memoria local:

Memoria local

Los accesos a la memoria local solo se producen para algunas variables automáticas como se menciona en los Calificadores de tipo de variable. Las variables automáticas que el compilador puede colocar en la memoria local son:

1. Arrays para los que no puede determinar que están indexados con cantidades constantes,
2. Grandes estructuras o matrices que consumirían demasiado espacio de registro,
3. Cualquier variable si el kernel usa más registros que los disponibles (esto también se conoce como registro de derrames).

¿Cómo funciona la asignación de registros?

Tenga en cuenta que la asignación de registros es un proceso extremadamente complicado por lo que no puede (y no debe) interferir con él. En su lugar, el compilador convertirá el código CUDA en código PTX (un tipo de código de bytes) que asume una máquina con infinitos registros. Puede escribir PTX en línea pero no hará mucho para registrar la asignación. El código PTX es un código independiente del dispositivo y es solo la primera etapa. En una segunda etapa, PTX se compilará en el código de ensamblaje del dispositivo, llamado SASS. El código SASS tiene las asignaciones de registro reales. El compilador SASS y su optimizador también serán la máxima autoridad sobre si una variable estará en los registros o en la memoria local. Todo lo que puede hacer es tratar de entender lo que hace el compilador SASS en ciertos casos y usarlo para su beneficio. La vista de correlación de código en Nsight puede ayudarlo con eso (ver más abajo). Sin embargo, dado que el compilador y el optimizador siguen cambiando, no hay garantías en cuanto a lo que habrá o no en los registros.

Registros insuficientes

El Apéndice G, sección 1 le dice cuántos registros puede tener un hilo. Busque "Número máximo de registros de 32 bits por subproceso". Para interpretar esa tabla, debe conocer su capacidad de cómputo (ver más abajo). No olvide que los registros se utilizan para todo tipo de cosas, y no solo se correlacionan con variables individuales. Los registros en todos los dispositivos hasta CC 3.5 son de 32 bits cada uno. Si el compilador es lo suficientemente inteligente (y el compilador CUDA sigue cambiando), puede, por ejemplo, empaquetar varios bytes en el mismo registro. La vista de correlación del código de Nsight (ver "Análisis de accesos de memoria" más abajo) también revela eso.

Indización constante vs. dinámica

Si bien la restricción de espacio es un obstáculo obvio para las matrices en el registro, lo que se supervisa fácilmente es el hecho de que, en el hardware actual (capacidad de cálculo 3.xy posterior), el compilador coloca cualquier matriz en la memoria local a la que se accede La indexación dinámica. Un índice dinámico es un índice que el compilador no puede entender. Las matrices a las que se accede con índices dinámicos no se pueden colocar en registros porque los registros deben ser determinados por el compilador y, por lo tanto, el registro real que se está utilizando no debe depender de un valor determinado en el tiempo de ejecución. Por ejemplo, dada una matriz arr , arr[k] es una indexación constante si y solo si k es una constante, o solo depende de las constantes. Si k , de alguna manera, depende de algún valor no constante, el compilador no puede calcular el valor de k y se obtiene una indexación dinámica . En los bucles en los que k comienza y termina en números constantes (pequeños), el compilador (lo más probable) puede desenrollar su bucle y aún así puede lograr una indexación constante.

Ejemplo

Por ejemplo, la clasificación de una pequeña matriz se puede hacer en los registros, pero debe usar redes de clasificación o enfoques "cableados" de manera similar, y no puede usar un algoritmo estándar porque la mayoría de los algoritmos utilizan la indexación dinámica.

Con una probabilidad bastante alta, en el siguiente ejemplo de código, el compilador mantiene toda la matriz aBytes en registros porque no es demasiado grande y los bucles se pueden desenrollar completamente (porque el bucle se repite en un rango constante). El compilador (muy probablemente) sabe a qué registro se está accediendo en cada paso y, por lo tanto, puede mantenerlo completamente en los registros. Tenga en cuenta que no hay garantías. Lo mejor que puede hacer es verificarlo caso por caso utilizando las herramientas de desarrollo de CUDA, como se describe a continuación.

__global__
void
testSortingNetwork4(const char * aInput, char * aResult)
{
    const int NBytes = 4;

    char aBytes[NBytes];

    // copy input to local array
    for (int i = 0; i < NBytes; ++i)
    {
        aBytes[i] = aInput[i];
    }

    // sort using sorting network
    CompareAndSwap(aBytes, 0, 2); CompareAndSwap(aBytes, 1, 3); 
    CompareAndSwap(aBytes, 0, 1); CompareAndSwap(aBytes, 2, 3); 
    CompareAndSwap(aBytes, 1, 2); 


    // copy back to result array
    for (int i = 0; i < NBytes; ++i)
    {
        aResult[i] = aBytes[i];
    }
}

Analizando accesos de memoria.

Una vez que haya terminado, generalmente desea verificar si los datos están realmente almacenados en registros o si fueron a la memoria local. Lo primero que puede hacer es decirle a su compilador que le dé estadísticas de memoria usando el --ptxas-options=-v . Una forma más detallada de analizar los accesos de memoria es usando Nsight .

Nsight tiene muchas características interesantes. Nsight para Visual Studio tiene un generador de perfiles incorporado y una vista de correlación de código CUDA <-> SASS. La característica se explica aquí . Tenga en cuenta que las versiones de Nsight para diferentes IDE probablemente se desarrollen de forma independiente y, por lo tanto, sus características pueden variar entre las diferentes implementaciones.

Si sigue las instrucciones en el enlace anterior (asegúrese de agregar las banderas correspondientes al compilar), puede encontrar el botón "Transacciones de memoria CUDA" en la parte inferior del menú inferior. En esa vista, desea encontrar que no hay ninguna transacción de memoria proveniente de las líneas que solo están trabajando en la matriz correspondiente (por ejemplo, las líneas CompareAndSwap en mi ejemplo de código). Porque si no informa el acceso a la memoria para esas líneas, (muy probablemente) pudo mantener todo el cómputo en registros y podría haber ganado una aceleración de miles, si no decenas de miles, del porcentaje (También puede querer Compruebe la ganancia de velocidad real, salga de esto!).

Averiguar la capacidad de cálculo

Para saber cuántos registros tiene, debe conocer la capacidad de cálculo de su dispositivo. La forma estándar de obtener dicha información del dispositivo es ejecutar el ejemplo de consulta del dispositivo. Para CUDA 5.5 en Windows 64bit, que se encuentra por defecto en C: / ProgramData / NVIDIA Corporation / CUDA Samples / v5.5 / Bin / win64 / Release / deviceQuery.exe (en Windows, la ventana de la consola se cerrará de inmediato, es posible que desee abrir cmd primero y ejecutarlo desde allí). Tiene una ubicación similar en Linux y MAC.

Si tiene Nsight para Visual Studio, solo vaya a Nsight -> Windows -> Información del sistema.

No optimices temprano

Estoy compartiendo esto hoy porque encontré este problema muy recientemente. Sin embargo, como se mencionó en este hilo , forzar que los datos estén en los registros definitivamente no es el primer paso que desea tomar. Primero, asegúrese de que realmente entiende lo que está pasando, luego aborde el problema paso a paso. Ver el código de ensamblaje es ciertamente un buen paso, pero generalmente no debería ser el primero. Si es nuevo en CUDA, la Guía de Mejores Prácticas de CUDA lo ayudará a descubrir algunos de esos pasos.

Answer 2

Estás mezclando espacio de memoria local y registro.

Las variables individuales y las matrices de tamaño constante se guardan automáticamente en el espacio de registro en el chip, casi sin costos de lectura y escritura.

Si excede la cantidad de registros por multiprocesador, se almacenarán en la memoria local.

La memoria local reside en el espacio de la memoria global y tiene el mismo ancho de banda lento para las operaciones de lectura y escritura.

#DEFINE P_SIZE = 50

__global__ void kernel()
{
    unsigned char p[P_SIZE];
}

Answer 3

La palabra clave que busca es __shared__ . Las matrices grandes no cabrán en el espacio de memoria compartida, pero el compilador debería usar la memoria compartida para una matriz pequeña de tamaño fijo como en este caso. Puede utilizar la palabra clave __shared__ para asegurarse de que esto suceda. Verá un error en tiempo de compilación si excede la cantidad máxima de memoria compartida para un bloque.

Answer 4

Todo lo que necesitas es esto:

__global__ my_kernel(int a)
{
    unsigned char p[50];
    ........
}

El compilador automáticamente derramará esto para enhebrar la memoria local si es necesario. Pero tenga en cuenta que la memoria local se almacena en SDRAM fuera de la GPU y es tan lenta como la memoria global. Entonces, si espera que esto produzca una mejora en el rendimiento, es posible que esté en una decepción ...