Fermi (microarquitectura)
Fermi es el nombre en clave de una microarquitectura de unidad de procesamiento de gráficos (GPU) desarrollada por Nvidia, lanzada por primera vez al por menor en abril de 2010, como sucesora de la microarquitectura de Tesla. Fue la microarquitectura principal utilizada en las series GeForce 400 y GeForce 500. Le siguió Kepler y se usó 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 estaciones de trabajo, Fermi encontró uso en la serie Quadro x000, los modelos Quadro NVS, así como en los módulos informáticos Nvidia Tesla. Todas las GPU Fermi de escritorio se fabricaron en 40nm, las GPU Fermi móviles en 40nm y 28nm. Fermi es la microarquitectura más antigua de NVIDIA que recibió soporte para la API de renderizado de Microsoft Direct3D 12 feature_level 11.
| Fermi | ||
|---|---|---|
| Información | ||
| Tipo | Microarquitectura | |
| Desarrollador | Nvidia | |
| Fecha de lanzamiento | Abril de 2010 | |
La arquitectura lleva el nombre de Enrico Fermi, un físico italiano.
Descripción general editar
Las unidades de procesamiento gráfico (GPU) de Fermi cuentan con 3.000 millones de transistores y se muestra un esquema en la figura 1.
- Streaming Multiprocessor (SM): compuesto por 32 núcleos CUDA (consulte las secciones Streaming Multiprocessor y CUDA core).
- Programador global GigaThread: distribuye bloques de subprocesos a programadores de subprocesos SM y administra los cambios de contexto entre subprocesos durante la ejecución (consulte la sección Programación Warp).
- Interfaz de host: conecta la GPU a la CPU a través de un bus PCI-Express v2 (tasa de transferencia máxima de 8 GB/s).
- DRAM: admite hasta 6 GB de memoria DRAM GDDR5 gracias a la capacidad de direccionamiento de 64 bits (consulte la sección Arquitectura de memoria).
- Frecuencia de reloj: 1.5 GHz (no publicado por NVIDIA, pero estimado por Insight 64).
- Máximo rendimiento: 1,5 TFlops.
- Reloj de memoria global: 2 GHz.
- Ancho de banda DRAM: 192 GB/s.
Multiprocesador de transmisión (SM - Streaming Multiprocessor) editar
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).
Unidades de carga/almacenamiento editar
Permita que las direcciones de origen y destino se calculen para 16 subprocesos por reloj. Cargue y almacene los datos desde/en caché o DRAM.
Unidades de funciones especiales (SFU - Special Functions Units) editar
Ejecutar instrucciones trascendentales como seno, coseno, recíproco y raíz cuadrada. Cada SFU ejecuta una instrucción por subproceso, por reloj; una deformación se ejecuta durante ocho relojes. La canalización de SFU se desacopla 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.
Núcleo CUDA editar
Unidad lógica aritmética entera (ALU - Arithmetic Logic Unit): Admite una precisión total de 32 bits para todas las instrucciones, de acuerdo con los requisitos del lenguaje de programación estándar. También está optimizado para admitir de manera eficiente operaciones de precisión extendida y de 64 bits.
Unidad de coma flotante (FPU - Floating Point Unit) editar
Implementa el nuevo estándar de coma flotante IEEE 754-2008, que proporciona la instrucción de suma y multiplicación fusionada (FMA) para aritmética de precisión simple y doble. Se pueden realizar hasta 16 operaciones de suma y multiplicación fusionadas de doble precisión por SM, por reloj.[1]
Suma y multiplicación fusionada editar
La suma y multiplicación fusionada (FMA) realiza multiplicaciones y sumas (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.
Planificación warp editar
La arquitectura Fermi utiliza un planificador de subprocesos distribuidos de dos niveles.
Cada SM puede emitir instrucciones que consuman dos de las cuatro columnas de ejecución verdes que se muestran en la Fig. 1 esquemática. 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.
Tenga en cuenta que las operaciones de coma flotante de 64 bits consumen las dos primeras columnas de ejecución. Esto implica que un SM puede emitir hasta 32 operaciones de punto flotante de precisión simple (32 bits) o 16 operaciones de punto flotante de precisión doble (64 bits) a la vez.
Motor GigaThread editar
El motor GigaThread programa bloques de subprocesos para varios SM.
Planificador Warp dual editar
En el nivel de SM, cada programador de warp distribuye warps de 32 subprocesos a sus unidades de ejecución. Los hilos se programan en grupos de 32 hilos llamados warps. Cada SM cuenta con dos programadores warp y dos unidades de despacho de instrucciones, lo que permite emitir y ejecutar dos warps simultáneamente. El programador 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 se pueden emitir de forma dual; dos instrucciones enteras, dos instrucciones flotantes o una combinación de instrucciones enteras, de punto flotante, de carga, de almacenamiento y SFU se pueden emitir simultáneamente. Las instrucciones de precisión doble no admiten el envío doble con ninguna otra operación.
Rendimiento editar
La potencia de procesamiento de precisión simple teórica 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 dos 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 sus 32 núcleos CUDA completamente utilizados.[2] 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 en GF100/110. Sin embargo, en la práctica, esta potencia de doble precisión solo está disponible en las tarjetas profesionales Quadro y Tesla, mientras que las tarjetas GeForce de consumo tienen un límite de 1/8.[3]
Memoria editar
Caché L1 por SM y caché L2 unificado que da servicio a todas las operaciones (carga, almacenamiento y textura).
Registros editar
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 núcleo CUDA es 63. La cantidad 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 según la cantidad de subprocesos. Los registros tienen un ancho de banda muy alto: unos 8.000 GB/s.
L1+Memoria compartida editar
Memoria en chip que se puede usar 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 los datos en el chip y reduce en gran medida el tráfico fuera del chip. Los subprocesos en el mismo bloque de subprocesos pueden acceder a la memoria compartida. Proporciona acceso de latencia baja (10-20 ciclos) y 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 usa la idea del bloc de notas local[4]
Memoria local editar
La memoria local se entiende como una ubicación de memoria utilizada para contener registros "derramados". El derrame 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) Los arreglos que el compilador no puede determinar están indexados con cantidades constantes; (2) Grandes estructuras o matrices que consumirían demasiado espacio de registro; Cualquier variable que el compilador decida volcar a la memoria local cuando un kernel usa más registros de los que están disponibles en el SM.
Caché L2 editar
Caché L2 unificado de 768 KB, compartido entre los 16 SM, que da servicio a todas las cargas y almacenamiento desde/hacia la memoria global, incluidas las copias hacia/desde el host de la CPU y también las solicitudes de textura. El subsistema de caché L2 también implementa operaciones atómicas, que se utilizan para administrar el acceso a los datos que deben compartirse entre bloques de subprocesos o incluso núcleos.
Memoria global editar
Accesible por todos los subprocesos, así como por el host (CPU). Alta latencia (400-800 ciclos).
Descompresión/compresión de video editar
Consulte Nvidia NVDEC (anteriormente llamado NVCUVID), así como Nvidia PureVideo.
La tecnología Nvidia NVENC aún no estaba disponible, pero se introdujo en el sucesor, Kepler.
Chips Fermi editar
- GF100
- GF104
- GF106
- GF108
- GF110
- GF114
- GF116
- GF117
- GF119
Véase también editar
Referencias editar
- ↑ «NVIDIA's Next Generation CUDA Compute Architecture: Fermi». 2009. Consultado el 7 de diciembre de 2015.
- ↑ Glaskowsky, Peter N. (septiembre de 2009). «NVIDIA's Fermi: The First Complete GPU Computing Architecture». Consultado el 6 de diciembre de 2015. «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».
- ↑ Smith, Ryan (26 de marzo de 2010). «NVIDIA's GeForce GTX 480 and GTX 470: 6 Months Late, Was It Worth the Wait?». AnandTech. p. 6. Consultado el 6 de diciembre de 2015. «El rendimiento de FP64 de la serie GTX 400 tiene un tope de 1/8 (12,5 %) de su rendimiento de FP32, a diferencia de lo que el hardware puede hacer de forma nativa de 1/2 (50 %) de FP32».
- ↑ Patterson, David (30 de septiembre de 2009). «The Top 10 Innovations in the New NVIDIA Fermi Architecture, and the Top 3 Next Challenges». Parallel Computing Research Laboratory & NVIDIA. Consultado el 3 de octubre de 2013.
General editar
- N. Brookwood, "NVIDIA Solves the GPU Computing Puzzle."
- P.N. Glaskowsky, "NVIDIA’s Fermi: The First Complete GPU Computing Architecture."
- N. Whitehead, A. Fit-Florea, "Precision & Performance: Floating Point and IEEE 754 Compliance for NVIDIA GPUs.", 2011.
- Oberman, S.F.; Siu, M.Y. (2005). «A High-Performance Area-Efficient Multifunction Interpolator». 17th IEEE Symposium on Computer Arithmetic (ARITH'05). pp. 272-279. ISBN 0-7695-2366-8. S2CID 14975421. doi:10.1109/arith.2005.7.
- R. Farber, "CUDA Application Design and Development," Morgan Kaufmann, 2011.
- NVIDIA Application Note "Tuning CUDA applications for Fermi".
Enlaces externos editar
Datos: Q3068993
