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

gpu programming cuda tutorial

Question

Estoy tratando de entender el uso de recursos para cada uno de mis hilos CUDA para un kernel escrito a mano.

kernel.cu mi archivo kernel.cu en un archivo nvcc -arch=sm_20 -ptxas-options=-v con nvcc -arch=sm_20 -ptxas-options=-v

y obtuve la siguiente salida

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]

Mirando el resultado anterior, ¿es correcto decir que

• cada hilo CUDA está usando 46 registros?
• no hay un registro que se derrame a la memoria local?

También estoy teniendo problemas para entender la salida.

• Mi kernel está llamando a un montón de funciones __device__ . ¿Es 72 bytes la suma total de la memoria para los cuadros de pila de las funciones __global__ y __device__ ?

• Cuál es la diferencia entre las 0 byte spill stores 0 bytes spill loads

• ¿Por qué la información para cmem (que cmem es memoria constante) se 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?

Answer 1

• Cada hilo CUDA está usando 46 registros? Sí correcto
• No hay un registro que se derrame a la memoria local? Sí correcto
• ¿Son 72 bytes la suma total de la memoria para los marcos 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: tenga en cuenta también que el recuento de carga / almacenamiento de derrames se basa en el análisis estático según lo descrito por @njuffa en los comentarios a continuación
• ¿Por qué la información para cmem (que asumo es memoria constante) se 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?
  • La memoria constante se usa para algunos propósitos, incluidas __constant__ variables __constant__ y los argumentos del kernel, se usan diferentes "bancos", que comienzan a ser un poco detallados pero siempre que use menos de 64KB para las variables __constant__ y menos de 4KB para los argumentos del kernel estará bien.