Límites de registros y memoria compartida

Cada SM posee una cantidad fija de registros y memoria compartida. Cada bloque residente debe tomar su parte de estos recursos.

Su kernel utiliza cierta cantidad de registros por hilo. Multiplíquela por los hilos del bloque para obtener el coste de registros de un bloque.

Si cada hilo necesita muchos registros, caben menos hilos y, por tanto, permanecen residentes menos bloques. Los kernels que usan muchos registros pierden ocupación.

Compile con esta opción y nvcc mostrará los registros por hilo, lo que le permitirá detectar cuándo un kernel consume demasiados registros.

Puede imponer un límite mediante un launch bound para que el compilador utilice menos registros y permita que más warps permanezcan residentes.

Cada bloque puede solicitar memoria compartida. El SM solo admite tantos bloques como permita su memoria compartida, a menudo entre 48 y 100 KB.

Solicite un tile __shared__ grande y solo cabrán uno o dos bloques por SM. Los tiles grandes intercambian ocupación por reutilización de datos.

El SM calcula los bloques permitidos por los registros, la memoria compartida y el límite de warps. El menor de estos valores determina la ocupación.

Cuando un hilo necesita más registros de los permitidos, los adicionales se derraman a la memoria local, que es lenta. Los derrames pueden perjudicar más que una ocupación baja.

Reducir los registros o la memoria compartida aumenta la ocupación, pero una reducción demasiado agresiva provoca derrames. El punto óptimo está en el equilibrio.

Lea siempre el uso real de registros y memoria compartida que indica el compilador antes de ajustar el kernel. Adivinar suele llevar a elegir el parámetro equivocado.

Recuerde cómo decide el SM cuántos bloques pueden permanecer residentes.

Ha visto que los registros y la memoria compartida son presupuestos fijos del SM, y que el límite más restrictivo limita la ocupación. Preste atención a los derrames. 🧮