Interpretación de salida de --ptxas-options = -v

Estoy tratando de entender el uso de recursos para cada uno de mis hilos CUDA para un kernel escrito a mano.Interpretación de salida de --ptxas-options = -v

compilé mi archivo kernel.cu a un archivo kernel.o con nvcc -arch=sm_20 -ptxas-options=-v

y me dio el siguiente resultado

ptxas info : Compiling entry function '_Z12searchkernel6octreePidiPdS1_S1_' for 'sm_20' 
ptxas info : Function properties for _Z12searchkernel6octreePidiPdS1_S1_ 
    72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14]

En cuanto a la salida anterior, ¿es correcto decir que

cada hilo CUDA está usando 46 registros?
no hay ningún registro derramando a la memoria local?

También estoy teniendo algunos problemas para comprender la salida.

Mi núcleo llama a un montón de funciones __device__. ¿Es de 72 bytes la suma total de la memoria para las estructuras de pila de las funciones __global__ y __device__?
¿Cuál es la diferencia entre 0 byte spill stores y 0 bytes spill loads
¿Por qué es la información para cmem (que estoy asumiendo que es la memoria constante) se repite dos veces con diferentes figuras? Dentro del kernel no estoy usando ninguna memoria constante . ¿Eso significa que el compilador está, bajo el capó, va a decirle a la GPU que use algo de memoria constante?

Fuente

2012-09-12 curiousexplorer

'usados' 46 registros indican que el compilador ha reservado 46 registros por hilo para el kernel compilado y los demás registros se derramen. Puede encontrar el número de registros derramados restando este número (46) del número total de registro utilizado en el PTX del kernel. – ahmad

@Ahmad: La primera oración es correcta, pero la segunda no. Un núcleo puede usar menos que los registros máximos permisibles por hilo y no tener derrames en la memoria local. – talonmies

Para elaborar la respuesta de talonmies, PTX es una abstracción de alto nivel con registros infinitos. Esto se debe a que se puede compilar para varias generaciones de GPU y la cantidad de registros puede ser diferente. Solo cuando compila hasta el código específico de la máquina, puede ver realmente el uso del registro. En cualquier caso, ptxas (compilando PTX con el código específico de la máquina) le informa la cantidad de derrames. El compilador – Tom

Cada hilo CUDA está utilizando 46 registros? Sí, correcto
¿No se está desviando el registro a la memoria local? Sí, es correcto
¿Tiene 72 bytes la suma total de la memoria para los cuadros de pila de las funciones __global__ y __device__? Sí, correcto
¿Cuál es la diferencia entre las tiendas de derrames de 0 bytes y las cargas de derrames de 0 bytes?
- Pregunta razonable, las cargas podrían ser mayores que las tiendas ya que podría derramar un valor calculado, cargarlo una vez, descartarlo (es decir, almacenar algo más en ese registro) y luego cargarlo nuevamente (es decir, reutilizarlo). Actualización: nota también que el recuento derrame de carga/almacenamiento se basa en el análisis estático como se describe por @njuffa en los comentarios a continuación
¿Por qué es la información para cmem (que estoy asumiendo que es la memoria constante) repite dos veces con diferentes figuras? Dentro del kernel no estoy usando ninguna memoria constante. ¿Eso significa que el compilador, bajo el capó, va a decirle a la GPU que use algo de memoria constante?
- memoria constante se utiliza durante unos propósitos, incluyendo __constant__ variables y argumentos del núcleo, diferentes "bancos" se utilizan, que empieza a ser un poco detallada pero siempre y cuando se utiliza menos de 64 KB para sus __constant__ las variables y menos de 4 KB para los argumentos del kernel estarás bien.

Fuente

2012-09-12 13:31:16 Tom

Tenga en cuenta que las cargas de derrame y las tiendas se cuentan estáticamente, es decir, el número de carga local y las instrucciones de la tienda local multiplicadas por el ancho acceso de cada carga/tienda. Están normalizados en bytes porque el compilador puede vectorizar cargas/almacenes de derrames si tiene suficiente información sobre la alineación y la asignación de registros lo permite. Como los conteos son estáticos, esto no es directamente una medida del tráfico de derrames, ya que los derrames/rellenos pueden ser bucles internos. Las cargas de derrame pueden exceder las reservas de derrames si hay una reutilización de los datos derramados. Esto implicaría que los bytes de carga de derrame> = cubren los bytes del almacén. – njuffa

Gracias @njuffa - excelentes puntos. El compilador no puede conocer recuentos de viajes para bucles (a menos que sea constante en tiempo de compilación). La mejor manera de analizar realmente el costo del derrame/relleno es usar un perfilador como Nsight (o NVVP independiente), que le proporcionará datos basados en la ejecución en lugar de en la compilación. – Tom

De acuerdo con respecto a la creación de perfiles. Las estadísticas de derrames del compilador son medianamente útiles como indicadores de primera línea. Si no hay derrames, no hay nada de qué preocuparse. Si los números son pequeños (por ejemplo, <32 bytes), la memoria caché L1 debería ocuparse de ellos sin afectar el rendimiento (recuerde que los números son por subproceso, ya que la memoria local de subprocesos se usa para derrames). Si las cifras son miles, es probable que haya un impacto negativo en el rendimiento y puede ser momento de un análisis más detallado. – njuffa

Interpretación de salida de --ptxas-options = -v

Respuesta

Cuestiones relacionadas