Fermi es el nombre en código de una microarquitectura de unidad de procesamiento gráfico (GPU) desarrollada por Nvidia , lanzada por primera vez al mercado en abril de 2010, como sucesora de la microarquitectura Tesla . Fue la microarquitectura principal utilizada en las series GeForce 400 y 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 recibe soporte para la API de renderizado Direct3D 12 feature_level 11 de Microsoft.
A Fermi le siguió Kepler , y se utilizó junto con Kepler en las series GeForce 600 , GeForce 700 y GeForce 800 , en estas dos últimas solo en GPU móviles .
En el mercado de 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 función especial (SFU), un bloque de 64 KB de memoria en chip de alta velocidad (consulte la subsección Memoria compartida L1+) y una interfaz con el caché L2 (consulte la subsección Caché L2).
Permite calcular direcciones de origen y destino para 16 subprocesos por ciclo de reloj. Carga y almacena datos desde/hacia la memoria caché o DRAM .
Ejecutar instrucciones trascendentales como seno, coseno, recíproco y raíz cuadrada. Cada SFU ejecuta una instrucción por subproceso, por ciclo de reloj; una deformación se ejecuta en ocho ciclos de reloj. La tubería de la SFU está desacoplada de la unidad de despacho, lo que permite que la unidad de despacho envíe instrucciones 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. [ ¿cuál? ] También está optimizado para soportar de manera eficiente 64 bits en modelos de estaciones de trabajo y servidores, pero está artificialmente limitado para las versiones para consumidores.
Implementa el nuevo estándar de punto flotante IEEE 754-2008, que proporciona la instrucción de multiplicación-suma fusionada (FMA) para operaciones aritméticas de precisión simple y doble. Se pueden realizar hasta 16 operaciones de multiplicación-suma fusionada de precisión doble por SM, por reloj. [2]
La multiplicación y suma fusionadas (FMA) realizan 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 precisa que realizar las operaciones por separado.
La arquitectura Fermi utiliza un programador de subprocesos distribuido de dos niveles .
Cada SM puede emitir instrucciones que consuman dos de las cuatro columnas de ejecución verdes que se muestran en el esquema de la Fig. 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 las 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.
En el nivel SM, cada programador de warp distribuye warps de 32 subprocesos a sus unidades de ejecución. Cada SM cuenta con dos programadores de warp y dos unidades de despacho de instrucciones, lo que permite que se emitan y ejecuten dos warps simultáneamente. El programador dual de warp 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 se pueden emitir de forma dual; se pueden emitir simultáneamente dos instrucciones de números enteros, dos instrucciones flotantes o una combinación de instrucciones de números enteros, de punto flotante, de carga, de almacenamiento y SFU. Las instrucciones de doble precisión no admiten el despacho dual con ninguna otra operación. [ cita requerida ]
La potencia de procesamiento teórica 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 shader (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 completamente utilizados. [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 de procesamiento de doble precisión teórica de una GPU Fermi es la mitad del rendimiento de precisión simple de una GF100/110. Sin embargo, en la práctica, esta potencia de doble precisión solo está disponible en tarjetas profesionales Quadro y Tesla , mientras que las tarjetas GeForce para consumidores están limitadas a 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 subproceso tiene acceso a sus propios registros y no a los de otros subprocesos. El número máximo de registros que puede utilizar un núcleo 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 lo tanto, los requisitos de recursos) aumenta en número de subprocesos. Los registros tienen un ancho de banda muy alto: alrededor de 8000 GB/s.
Memoria en chip que se puede utilizar para almacenar en caché datos de 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 chip y reduce en gran medida el tráfico fuera del chip. La memoria compartida es accesible para los subprocesos en el mismo bloque de subprocesos. Proporciona acceso de baja latencia (10-20 ciclos) y un 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 video, etc.). David Patterson dice que esta memoria compartida utiliza la idea de un bloc de notas local [5].
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 subproceso requiere más almacenamiento de registros que el 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 núcleo utiliza más registros que los disponibles en el SM.
Caché L2 unificada de 768 KB, compartida entre los 16 SM, que atiende toda la carga y el almacenamiento desde/hacia la memoria global, incluidas las copias hacia/desde el host de la CPU y también las solicitudes de texturas. El subsistema de caché L2 también implementa operaciones atómicas, que se utilizan para gestionar el acceso a los datos que deben compartirse entre bloques de subprocesos o incluso núcleos.
La memoria global (VRAM) es accesible para todos los subprocesos directamente, así como para el sistema host a través del bus PCIe. Tiene una latencia alta de 400 a 800 ciclos. [ cita requerida ]
Consulte Nvidia NVDEC (anteriormente llamada NVCUVID) así como Nvidia PureVideo .
La tecnología NVENC de Nvidia aún no estaba disponible, pero se introdujo en el 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 está limitado a 1/8 (12,5 %) de su rendimiento FP32, a diferencia de lo que el hardware puede hacer de forma nativa con 1/2 (50 %) FP32
