Fermi (microarquitectura) - Fermi (microarchitecture)

Nvidia Fermi
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

Fig. 1.
Convención de arquitectura NVIDIA Fermi en cifras: naranja - programación y envío; verde - ejecución; azul claro: registros y cachés.
Disparo de la GPU GF100 que se encuentra dentro de las tarjetas GeForce GTX 470

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.

L1 + memoria compartida

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

Referencias

General

enlaces externos