Fermi (microarquitectura) - Fermi (microarchitecture)
Fecha de lanzamiento | Abril de 2010 |
---|---|
Proceso de fabricación | 40 nm y 28 nm |
Historia | |
Predecesor | Tesla 2.0 |
Sucesor | Kepler |
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 mercado minorista en abril de 2010, como sucesora de la microarquitectura de Tesla . Fue la microarquitectura principal utilizada en las series GeForce 400 y GeForce 500 . Fue seguido por Kepler , y se utiliza junto con Kepler en la serie GeForce 600 , 700 de la serie GeForce y GeForce serie 800 , en los dos últimos sólo en móviles GPU. En el mercado de las estaciones de trabajo, Fermi encontró uso en la serie Quadro x000, los modelos Quadro NVS, así como en los módulos de computación Nvidia Tesla . Todas las GPU Fermi de escritorio se fabricaron en 40 nm, las GPU Fermi móviles en 40 nm y 28 nm. Fermi es la microarquitectura más antigua de NVIDIA que recibió soporte para la API de renderizado de Microsoft Direct3D 12 feature_level 11.
La arquitectura lleva el nombre de Enrico Fermi , un físico italiano.
Visión general
Las unidades de procesamiento gráfico ( GPU ) de Fermi cuentan con 3000 millones de transistores y en la figura 1 se muestra un esquema.
- 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 los programadores de subprocesos SM y administra los cambios de contexto entre subprocesos durante la ejecución (consulte la sección Programación de deformaciones).
- 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).
- Rendimiento máximo: 1,5 TFlops.
- Reloj de memoria global: 2 GHz.
- Ancho de banda DRAM : 192 GB / s.
Streaming multiprocesador
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 L1 + Memoria compartida) y una interfaz para la caché L2 ( consulte la subsección Caché L2).
Unidades de carga / almacenamiento
Permita que las direcciones de origen y destino se calculen para 16 subprocesos por reloj. Cargue y almacene los datos desde / hacia la caché o DRAM .
Unidades de funciones especiales (SFU)
Ejecuta instrucciones trascendentales como pecado, coseno, recíproco y raíz cuadrada. Cada SFU ejecuta una instrucción por hilo, por reloj; una deformación se ejecuta en ocho relojes. La tubería SFU está desacoplada de la unidad de envío, lo que permite que la unidad de envío emita a otras unidades de ejecución mientras la SFU está ocupada.
Núcleo CUDA
Unidad Aritmética Lógica Entera (ALU) : Admite 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 ampliadas y de 64 bits.
Unidad de punto flotante (FPU)
Implementa el nuevo estándar de punto flotante IEEE 754-2008, que proporciona la instrucción fusionada de multiplicación-suma (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.
Multiplicar-agregar fusionado
Fused multiply-add (FMA) realiza la multiplicación y la suma (es decir, A * B + C) con un solo paso de redondeo final, sin pérdida de precisión en la suma. FMA es más precisa que realizar las operaciones por separado.
Programación de deformaciones
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 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 punto 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 doble precisión (64 bits) a la vez.
Motor GigaThread
El motor GigaThread programa bloques de subprocesos para varios SM
Programador de doble deformación
En el nivel SM, cada programador warp distribuye warps de 32 hilos a sus unidades de ejecución. Los subprocesos se programan en grupos de 32 subprocesos denominados deformaciones. Cada SM cuenta con dos programadores warp y dos unidades de envío de instrucciones, lo que permite que se emitan y ejecuten dos warps al mismo tiempo. 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 en dos ocasiones; Se pueden emitir simultáneamente dos instrucciones enteras, dos instrucciones flotantes o una combinación de instrucciones enteras, de coma flotante, de carga, de almacenamiento y SFU. Las instrucciones de doble precisión no admiten el envío dual con ninguna otra operación.
Rendimiento
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 podrí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. 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 teórica 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 tarjetas profesionales Quadro y Tesla , mientras que las tarjetas GeForce de consumo tienen un límite de 1/8.
Memoria
Caché L1 por SM y caché L2 unificado que da servicio a todas las operaciones (carga, almacenamiento y textura).
Registros
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 fácilmente de 63 a 21 a medida que la carga de trabajo (y, por tanto, los requisitos de recursos) aumenta según el número de subprocesos. Los registros tienen un ancho de banda muy alto: alrededor de 8.000 GB / s.
Memoria en chip que se puede utilizar para almacenar en caché datos para subprocesos individuales (registro de derrame / 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 extensa de datos en el chip y reduce en gran medida el tráfico fuera del chip. La memoria compartida es accesible por los subprocesos en el mismo bloque de subprocesos. Proporciona acceso de baja latencia (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 video, etc.). David Patterson dice que esta memoria compartida usa la idea de scratchpad local
Memoria local
La memoria local se entiende como una ubicación de memoria que se utiliza para almacenar 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) Las matrices que el compilador no puede determinar están indexadas con cantidades constantes; (2) Grandes estructuras o matrices que consumirían demasiado espacio de registro; Cualquier variable que el compilador decida pasar a la memoria local cuando un kernel usa más registros de los que están disponibles en el SM.
Caché L2
Caché L2 unificado de 768 KB, compartido 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 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
Accesible por todos los subprocesos, así como por el host (CPU). Alta latencia (400-800 ciclos).
Compresión / descompresión de video
Consulte Nvidia NVDEC (anteriormente llamado NVCUVID) y Nvidia PureVideo .
La tecnología Nvidia NVENC aún no estaba disponible, pero se introdujo en el sucesor, Kepler .
Chips de Fermi
- GF100
- GF104
- GF106
- GF108
- GF110
- GF114
- GF116
- GF118
- GF119
- GF117
Ver también
- Lista de unidades de procesamiento de gráficos Nvidia
- CUDA
- Interfaz de enlace escalable (SLI)
- Qualcomm Adreno
Referencias
General
- N. Brookwood, "NVIDIA resuelve el rompecabezas de la informática de la GPU".
- PN Glaskowsky, "Fermi de NVIDIA: la primera arquitectura informática completa de GPU".
- N. Whitehead, A. Fit-Florea, "Precisión y rendimiento: punto flotante y cumplimiento de IEEE 754 para GPU NVIDIA". , 2011.
- Oberman, SF; Siu, MI (2005). "Un interpolador multifunción de área eficiente de alto rendimiento". 17º Simposio IEEE sobre Aritmética Informática (ARITH'05) . págs. 272-279. doi : 10.1109 / arith.2005.7 . ISBN 0-7695-2366-8. S2CID 14975421 .
- R. Farber, "Diseño y desarrollo de aplicaciones CUDA", Morgan Kaufmann, 2011.
- Nota de aplicación de NVIDIA "Ajuste de aplicaciones CUDA para Fermi".