Fermi (microarquitectura)

Permitir que se calculen las direcciones de origen y destino para 16 hilos por ciclo de reloj. Cargar y almacenar los datos desde/hacia la caché o la DRAM .

Ejecuta instrucciones trascendentales como seno, coseno, recíproco y raíz cuadrada. Cada SFU ejecuta una instrucción por hilo y por ciclo de reloj; un warp se ejecuta durante ocho ciclos de reloj. La tubería de la SFU está desacoplada de la unidad de despacho, lo que permite que esta última emita instrucciones a otras unidades de ejecución mientras la SFU está ocupada.

Admite precisión de 32 bits para todas las instrucciones, de acuerdo con los requisitos de los lenguajes de programación estándar. También está optimizado para admitir eficientemente 64 bits en modelos de estaciones de trabajo y servidores, pero su rendimiento se ve limitado artificialmente en las versiones para el consumidor.

Implementa el nuevo estándar de punto flotante IEEE 754-2008, proporcionando la instrucción de multiplicación-suma fusionada (FMA) para aritmética de precisión simple y doble. Se pueden realizar hasta 16 operaciones de multiplicación-suma fusionadas de doble precisión por SM, por ciclo de reloj. ^{[ 1 ]}

A nivel de SM, cada planificador de warps distribuye warps de 32 hilos a sus unidades de ejecución. Cada SM cuenta con dos planificadores de warps y dos unidades de despacho de instrucciones, lo que permite emitir y ejecutar dos warps simultáneamente. El planificador de warps dual selecciona dos warps y emite una instrucción de cada warp a un grupo de 16 núcleos, 16 unidades de carga/almacenamiento o 4 SFU. La mayoría de las instrucciones pueden emitirse simultáneamente; dos instrucciones enteras, dos instrucciones de punto flotante o una combinación de instrucciones enteras, de punto flotante, de carga, de almacenamiento y de SFU pueden emitirse simultáneamente. Las instrucciones de doble precisión no admiten el despacho dual con ninguna otra operación.

Cada SM dispone de 32 KB de registros de 32 bits. Cada hilo tiene acceso a sus propios registros, no a los de los demás. El número máximo de registros que puede utilizar un kernel CUDA es 63. El número de registros disponibles se reduce gradualmente de 63 a 21 a medida que aumenta la carga de trabajo (y, por lo tanto, los requisitos de recursos) con el número de hilos. Los registros tienen un ancho de banda muy elevado: aproximadamente 8000  GB/s.

Memoria en el chip que puede usarse para almacenar en caché datos para hilos individuales (desbordamiento de registros/caché L1) y/o para compartir datos entre varios hilos (memoria compartida). Esta  memoria de 64 KB puede configurarse como 48  KB de memoria compartida con 16  KB de caché L1, o 16  KB de memoria compartida con 48  KB de caché L1. La memoria compartida permite que los hilos dentro del mismo bloque de hilos cooperen, facilita la reutilización extensiva de datos en el chip y reduce considerablemente el tráfico fuera del chip. La memoria compartida es accesible por los hilos en el mismo bloque de hilos. Proporciona acceso de baja latencia (10-20 ciclos) y un ancho de banda muy alto (1600  GB/s) a cantidades moderadas de datos (como resultados intermedios en una serie de cálculos, una fila o columna de datos para operaciones matriciales, una línea de vídeo, etc.). David Patterson dice que esta memoria compartida utiliza la idea de bloc de notas local ^{[ 4 ].}

La memoria local se refiere a una ubicación de memoria utilizada para almacenar registros "desbordados". El desbordamiento de registros ocurre cuando un bloque de hilo requiere más almacenamiento de registros del disponible en un SM. La memoria local se utiliza solo para algunas variables automáticas (que se declaran en el código del dispositivo sin ninguno de los calificadores __device__, __shared__ o __constant__ ). Generalmente, una variable automática reside en un registro excepto en los siguientes casos: (1) Matrices que el compilador no puede determinar que están indexadas con cantidades constantes; (2) Estructuras o matrices grandes que consumirían demasiado espacio de registro; Cualquier variable que el compilador decida desbordar a la memoria local cuando un kernel usa más registros de los disponibles en el SM.

Caché L2 unificada de 768 KB, compartida entre los 16 SM, que gestiona todas las operaciones de carga y almacenamiento desde/hacia la memoria global, incluidas las copias desde/hacia la CPU host, así como las solicitudes de texturas. El subsistema de caché L2 también implementa operaciones atómicas, utilizadas para gestionar el acceso a datos que deben compartirse entre bloques de subprocesos o incluso núcleos.

La memoria global (VRAM) es accesible directamente por todos los subprocesos, así como por el sistema anfitrión a través del bus PCIe. Presenta una alta latencia de entre 400 y 800 ciclos.