Los registros y la memoria compartida dependiendo de la compilación de capacidad de cálculo?

Question

Hola, cuando se compila con nvcc -arch=sm_13 Obtengo:

ptxas info    : Used 29 registers, 28+16 bytes smem, 7200 bytes cmem[0], 8 bytes cmem[1] 

cuando yo uso nvcc -arch=sm_20 Obtengo:

ptxas info    : Used 34 registers, 60 bytes cmem[0], 7200 bytes cmem[2], 4 bytes cmem[16] 

Pensé que todos los parámetros del kernel se pasan a la memoria compartida pero para sm_20, no parece tan...?!Tal vez ellos también se pasan en los registros?La cabeza de mi función tiene el siguiente aspecto:

__global__ void func(double *, double , double, int)

Gracias hasta ahora!

Answer 1

En capacidad de cálculo 2.x dispositivos, los argumentos para granos se almacenan en la memoria constante.El registro de la diferencia es, probablemente, hacia abajo a las diferencias en el código generado para las matemáticas de la biblioteca de funciones entre las versiones.Hay cosas como funciones trascendentes o sqrt en el núcleo?

Answer 2

Como @talonmies estados de memoria compartida, las diferencias son debido a la SM 2.x dispositivos de pasar argumentos del kernel a través de la constante en lugar de la memoria compartida.

Sin embargo, una de las principales diferencias en el registro de uso en cm 2.x dispositivos es el hecho de que mientras que SM 1.x dispositivos han dedicado registros de dirección para cargar y almacenar instrucciones, SM 2.x utiliza registros de propósito general para las direcciones.Esto tiende a aumentar registrar la presión en cm 2.x.Por suerte, el archivo de registro también es 2x más grande en la GF100 (SM 2.0) vsGT200 (SM 1.3).