Obligar a CUDA a usar el registro para una variable

Question

Tengo muchos registros sin usar en mi kernel. Me gustaría decirle a CUDA que use algunos registros para guardar algunos datos, en lugar de hacer una lectura global de datos cada vez que lo necesito. (No soy capaz de utilizar mem compartido.)

__global__ void simple(float *gData) { 
float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

compilar w /: NVCC -arch sm_20 --ptxas-options = -v simple.cu, y consigo
0 bytes pila marco, 0 tiendas de derrames de bytes, 0 bytes cargas derrame
usado 2 registros, 40 bytes cmem [0]

__global__ void simple(float *gData) { 
register float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

registro declaración no hace nada.
0 bytes pila marco, 0 bytes tiendas de derrames, 0 cargas bytes derrame
Usado 2 registros, 40 bytes cmem [0]

__global__ void simple(float *gData) { 
volatile float rData[1024]; 
for(int i=0; i<1024; i++) { 
    rData[i]=gData[i]; 
    } 
// work on the data here 
} 

volátil declaración crea ordenador pila:
4096 bytes pila marco, 0 bytes tiendas de derrames, 0 bytes cargas derrame
usados ​​resistentes 21, 40 bytes cmem [0]

1) ¿hay una manera fácil de decir al compilador que utilice el espacio de registro para una variable?
2) ¿Dónde está 'stack frame': register, global mem, local mem, ...? ¿Qué es un marco de pila? (Desde cuando la GPU tiene una pila Una pila virtual??)
3) El archivo simple.ptx es básicamente vacío: (NVCC -arch sm_20 -ptx simple.cu)

.loc 2 14 2 
ret; 

Cualquier idea de dónde puede encontrar el código máquina/compilado real?

Answer 1

• matrices dinámicamente indexado no se pueden almacenar en los registros, porque el archivo de registro de la GPU no es direccionable de forma dinámica.
• Las variables escalares se almacenan automáticamente en los registros por el compilador.
• estáticamente indexados (es decir, donde el índice se puede determinar en tiempo de compilación), pequeñas arrays (por ejemplo, menos de 16 flotadores) puede ser almacenados en los registros por el compilador.

SM Las GPU 2.0 (Fermi) solo admiten hasta 63 registros por subproceso. Si se excede esto, los valores de registro se derramarán/llenarán desde la memoria local (fuera del chip), compatible con la jerarquía del caché. Las GPU SM 3.5 amplían esto a hasta 255 registros por subproceso.

En general, como menciona Jared, el uso de demasiados registros por subproceso no es deseable porque reduce la ocupación y, por lo tanto, reduce la capacidad de ocultación de latencia en el kernel. Las GPU prosperan en el paralelismo y lo hacen cubriendo la latencia de la memoria con el trabajo de otros hilos.

Por lo tanto, probablemente no deberías optimizar las matrices en los registros. En su lugar, asegúrese de que su memoria tenga acceso a esas matrices entre subprocesos lo más secuencialmente posible para que maximice la fusión (es decir, minimice las transacciones de memoria).

El ejemplo que proporciona puede ser un caso para la memoria compartida si:

1. muchos hilos en el bloque de utilizar los mismos datos, o
2. El per-hilo tamaño de la matriz es lo suficientemente pequeño para asignar espacio suficiente para todos los hilos en múltiples bloques de hilos (1024 flotantes por hilo es mucho).

Como njuffa mencionó, la razón por la cual su kernel solo usa 2 registros es porque no hace nada útil con los datos en el kernel, y el código muerto fue eliminado por el compilador.

Answer 2

Como ya se señaló, los registros (y el "espacio de parámetros" PTX) no se pueden indexar de forma dinámica. Para hacer eso, el compilador debería emitir código como para un bloque switch...case para convertir el índice dinámico en inmediato. No estoy seguro de que alguna vez lo haga automáticamente. Puede ayudarlo utilizando una estructura de tupla de tamaño fijo y un switch...case. Es probable que la metaprogramación C/C++ sea el arma de elección para mantener un código como este manejable.

Además, para CUDA 4.0, use el interruptor de línea de comando -Xopencc=-O3 para tener todo lo que no sean escalares simples (como estructuras de datos) asignados a los registros (consulte this post). Para CUDA> 4.0, debe deshabilitar la compatibilidad de depuración (no existe la opción de línea de comando -G; la optimización solo ocurre cuando la depuración está deshabilitada).

El nivel de PTX permite muchos más registros virtuales que el hardware. Esos son mapeados a registros de hardware en tiempo de carga. El límite de registro que especifique le permite establecer un límite superior en los recursos de hardware utilizados por el binario generado. Sirve como una heurística para que el compilador decida cuándo depositar (vea a continuación) los registros al compilar para PTX, de modo que se puedan satisfacer ciertas necesidades de concurrencia (consulte "límites de lanzamiento", "ocupación" y "ejecución simultánea del kernel" en la documentación de CUDA - también puede disfrutar this most interesting presentation).

Para las GPU Fermi hay como máximo 64 registros de hardware. El 64 (o el último - cuando se utiliza menos que el máximo del hardware) es utilizado por el ABI como el puntero de pila y por lo tanto para "registrar derrame" (significa liberar registros almacenando temporalmente sus valores en la pila y sucede cuando más registros son necesarios que disponibles) por lo que es intocable.