Reducción del número de registros utilizados en Kernel CUDA

Question

Tengo un núcleo que usa 17 registros, reducirlo a 16 me daría un 100% de ocupación. Mi pregunta es: ¿existen métodos que puedan usarse para reducir el número o los registros utilizados, excluyendo la reescritura completa de mis algoritmos de una manera diferente? Siempre he asumido que el compilador es mucho más inteligente que yo, así que, por ejemplo, a menudo uso variables adicionales solo por claridad. ¿Estoy equivocado en este pensamiento?

Tenga en cuenta: Yo sé de los --max_registers (o lo que sea la sintaxis es) bandera, pero el uso de la memoria local sería más perjudicial que una ocupación del 25% inferior (que debería probar esto)

Answer 1

Es realmente difícil de decir, el compilador nvcc no es muy inteligente en mi opinión.
Puede probar cosas obvias, por ejemplo, usando short en lugar de int, pasando y usando variables por referencia (por ejemplo, & variable), desenrollando bucles, usando plantillas (como en C++). Si tiene divisiones, funciones trascendentales, aplicadas en secuencia, intente hacerlas como un bucle. Intenta eliminar los condicionales, posiblemente reemplazándolos con cálculos redundantes.

Si publica un código, tal vez obtendrá respuestas específicas.

Answer 2

La ocupación puede ser un poco engañosa y el 100% de ocupación no debe ser su objetivo principal. Si puede obtener accesos totalmente fusionados a la memoria global, entonces, en una GPU de gama alta, una ocupación del 50% será suficiente para ocultar la latencia a la memoria global (para flotadores, incluso menor para dobles). Consulte la presentación Advanced CUDA C de GTC el año pasado para obtener más información sobre este tema.

En su caso, debe medir el rendimiento con y sin maxrregcount establecido en 16. La latencia de la memoria local debe ocultarse como resultado de tener suficientes hilos, suponiendo que no tiene acceso aleatorio a las matrices locales (lo que dar como resultado accesos no fusionados).

Para responder a su pregunta específica sobre la reducción de registros, publique el código para obtener respuestas más detalladas. Comprender cómo funcionan los compiladores en general puede ayudar, pero recuerde que nvcc es un compilador de optimización con un gran espacio de parámetros, por lo que minimizar el recuento de registros debe equilibrarse con el rendimiento general.

Answer 3

El recuento de instrucciones aumenta cuando la reducción del uso del registro tiene una explicación simple. El compilador podría estar usando registros para almacenar los resultados de algunas operaciones que se usan más de una vez a través de su código para evitar volver a calcular esos valores, cuando se obliga a usar menos registros, el compilador decide recalcular los valores que se almacenarían en los registros de otra manera.

Answer 4

En general, no es un buen enfoque para minimizar la presión de registro. El compilador hace un buen trabajo optimizando el rendimiento general proyectado del núcleo, y tiene en cuenta muchos factores, incluido el registro.

¿Cómo funciona cuando se registra la reducción causada menor velocidad

Lo más probable es que el compilador tuvo que derramar datos del registro insuficientes en la memoria "local", que es esencialmente la misma que la memoria global, y por lo tanto muy lento

Para fines de optimización, recomendaría usar palabras clave como const, volátil y demás cuando sea necesario, para ayudar al compilador en la fase de optimización.

De todos modos, no son estos pequeños problemas como los registros que a menudo hacen que los kernels CUDA corran lento.Recomiendo optimizar el trabajo con la memoria global, el patrón de acceso, el almacenamiento en memoria caché en la memoria de textura si es posible, las transacciones a través de la PCIe.

Answer 5

Utilizando la memoria caché compartida como puede conducir menos registrar el uso y evitar que se derrame registro a la memoria local ...

piensa que el kernel calcula algunos valores y estos valores calculados son utilizados por todos los hilos,

__global__ void kernel(...) { 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    int id0 = blockDim.x * blockIdx.x; 

    int reg = id0 * ...; 
    int reg0 = reg * a/x + y; 


    ... 

    int val = reg + reg0 + 2 * idx; 

    output[idx] = val > 10; 
} 

Por lo tanto, en lugar de mantener reg y reg0 como registros y hacer que se derramen en la memoria local (memoria global), podemos usar memoria compartida.

__global__ void kernel(...) { 
    __shared__ int cache[10]; 

    int idx = threadIdx.x + blockDim.x * blockIdx.x; 

    if (threadIdx.x == 0) { 
     int id0 = blockDim.x * blockIdx.x; 

     cache[0] = id0 * ...; 
     cache[1] = cache[0] * a/x + y; 
    } 
    __syncthreads(); 


    ... 

    int val = cache[0] + cache[1] + 2 * idx; 

    output[idx] = val > 10; 
} 

Tome un vistazo a este paper para más información ..