Fermi es el nombre en clave de una microarquitectura de unidad de procesamiento de gráficos (GPU) desarrollada por Nvidia , lanzada al comercio minorista por primera vez en abril de 2010, como sucesora de la microarquitectura de Tesla . Fue la microarquitectura principal utilizada en las series GeForce 400 y GeForce 500 . Todas las GPU Fermi de escritorio se fabricaron en 40 nm, las GPU Fermi móviles en 40 nm y 28 nm [ cita requerida ] . Fermi es la microarquitectura más antigua de NVIDIA que recibió soporte para la API de renderizado Direct3D 12 de Microsoft, nivel de característica 11.
A Fermi le siguió Kepler , y se utilizó junto con Kepler en las series GeForce 600 , GeForce 700 y GeForce 800 , en las dos últimas solo en GPU móviles .
En el mercado de las estaciones de trabajo, Fermi encontró uso en la serie Quadro x000, los modelos Quadro NVS y en los módulos informáticos Nvidia Tesla .
La arquitectura lleva el nombre de Enrico Fermi , un físico italiano.
Las unidades de procesamiento gráfico ( GPU ) de Fermi cuentan con 3.000 millones de transistores y en la figura 1 se muestra un esquema.
Cada SM cuenta con 32 núcleos CUDA de precisión simple, 16 unidades de carga/almacenamiento, cuatro unidades de funciones especiales (SFU), un bloque de 64 KB de memoria en chip de alta velocidad (consulte la subsección L1+Memoria compartida) y una interfaz para la caché L2 ( consulte la subsección Caché L2).
Permita que las direcciones de origen y destino se calculen para 16 subprocesos por reloj. Cargue y almacene los datos desde/hacia caché o DRAM .
Ejecute instrucciones trascendentales como pecado, coseno, recíproco y raíz cuadrada. Cada SFU ejecuta una instrucción por subproceso, por reloj; una deformación se ejecuta en ocho relojes. La tubería SFU está desacoplada de la unidad de despacho, lo que permite que la unidad de despacho emita a otras unidades de ejecución mientras la SFU está ocupada.
Admite precisión total de 32 bits para todas las instrucciones, de acuerdo con los requisitos del lenguaje de programación estándar. [ ¿cual? ] También está optimizado para admitir de manera eficiente 64 bits en modelos de estaciones de trabajo y servidores, pero está artificialmente paralizado para las versiones de consumo.
Implementa el nuevo estándar de punto flotante IEEE 754-2008, que proporciona la instrucción de suma múltiple fusionada (FMA) para aritmética de precisión simple y doble. Se pueden realizar hasta 16 operaciones de suma múltiple fusionadas de doble precisión por SM, por reloj. [2]
La suma múltiple fusionada (FMA) realiza la multiplicación y la suma (es decir, A*B+C) con un único paso de redondeo final, sin pérdida de precisión en la suma. FMA es más preciso que realizar las operaciones por separado.
La arquitectura Fermi utiliza un programador de subprocesos distribuido de dos niveles .
Cada SM puede emitir instrucciones que consumen dos de las cuatro columnas de ejecución verdes que se muestran en la figura esquemática 1. Por ejemplo, el SM puede mezclar 16 operaciones de los 16 núcleos de la primera columna con 16 operaciones de los 16 núcleos de la segunda columna, o 16 operaciones. de las unidades de carga/almacenamiento con cuatro de SFU, o cualquier otra combinación que especifique el programa.
Las operaciones de punto flotante de 64 bits requieren las dos primeras columnas de ejecución, por lo que se ejecutan a la mitad de la velocidad de las operaciones de 32 bits.
A nivel SM, cada programador warp distribuye warps de 32 subprocesos a sus unidades de ejecución. Cada SM cuenta con dos programadores warp y dos unidades de envío de instrucciones, lo que permite emitir y ejecutar dos warps simultáneamente. El programador de warp 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 dos veces; Se pueden emitir simultáneamente dos instrucciones de número entero, dos instrucciones flotantes o una combinación de instrucciones de número entero, punto flotante, carga, almacenamiento y SFU. Las instrucciones de doble precisión no admiten el envío dual con ninguna otra operación. [ cita necesaria ]
La potencia teórica de procesamiento de precisión simple de una GPU Fermi en GFLOPS se calcula como 2 (operaciones por instrucción FMA por núcleo CUDA por ciclo) × número de núcleos CUDA × velocidad de reloj del sombreador (en GHz). Tenga en cuenta que la generación anterior de Tesla podía emitir MAD+MUL a núcleos CUDA y SFU en paralelo, pero Fermi perdió esta capacidad ya que solo puede emitir 32 instrucciones por ciclo por SM, lo que mantiene solo sus 32 núcleos CUDA en pleno uso. [3] Por lo tanto, no es posible aprovechar las SFU para alcanzar más de 2 operaciones por núcleo CUDA por ciclo.
La potencia teórica de procesamiento de doble precisión de una GPU Fermi es la mitad del rendimiento de precisión simple de GF100/110. Sin embargo, en la práctica, esta potencia de doble precisión sólo está disponible en tarjetas Quadro y Tesla profesionales , mientras que las tarjetas GeForce de consumo tienen un límite de 1/8. [4]
Caché L1 por SM y caché L2 unificado que da servicio a todas las operaciones (carga, almacenamiento y textura).
Cada SM tiene 32K de registros de 32 bits. Cada hilo tiene acceso a sus propios registros y no a los de otros hilos. El número máximo de registros que puede utilizar un kernel CUDA es 63. El número de registros disponibles se degrada gradualmente de 63 a 21 a medida que la carga de trabajo (y, por tanto, los requisitos de recursos) aumenta con el número de subprocesos. Los registros tienen un ancho de banda muy alto: unos 8.000 GB/s.
Memoria en chip que se puede utilizar para almacenar en caché datos para subprocesos individuales (desbordamiento de registros/caché L1) y/o para compartir datos entre varios subprocesos (memoria compartida). Esta memoria de 64 KB se puede configurar 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 subprocesos dentro del mismo bloque de subprocesos cooperen, facilita la reutilización extensiva de datos en el chip y reduce en gran medida el tráfico fuera del chip. Los subprocesos del mismo bloque de subprocesos pueden acceder a la memoria compartida. Proporciona acceso de baja latencia (10-20 ciclos) y ancho de banda muy alto (1600 GB/s) para 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 un bloc de notas local [5]
La memoria local es una ubicación de memoria utilizada para contener registros "derramados". El desbordamiento de registros ocurre cuando un bloque de subprocesos requiere más almacenamiento de registros del que está disponible en un SM. La memoria local se usa 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 por lo siguiente: (1) Las matrices que el compilador no puede determinar están indexadas con cantidades constantes; (2) Estructuras o matrices grandes que consumirían demasiado espacio de registro; Cualquier variable que el compilador decida transferir a la memoria local cuando un kernel usa más registros de los disponibles en el SM.
Caché L2 unificado de 768 KB, compartido entre los 16 SM, que brinda servicios de carga y almacenamiento desde/hacia la memoria global, incluidas copias hacia/desde el host de la CPU y también solicitudes de textura. 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.
Todos los subprocesos pueden acceder a la memoria global (VRAM) directamente, así como al sistema host a través del bus PCIe. Tiene una alta latencia de 400-800 ciclos. [ cita necesaria ]
Consulte Nvidia NVDEC (anteriormente llamada NVCUVID) y Nvidia PureVideo .
La tecnología Nvidia NVENC aún no estaba disponible, pero se introdujo en su sucesor, Kepler .
Se pueden enviar un total de 32 instrucciones de uno o dos warps en cada ciclo a dos de los cuatro bloques de ejecución dentro de un Fermi SM.
El rendimiento FP64 de la serie GTX 400 tiene un límite de 1/8 (12,5 %) de su rendimiento FP32, a diferencia de lo que el hardware puede hacer de forma nativa de 1/2 (50 %) FP32.
