Nuevo algoritmo de escalado e inicialización con NVIDIA Collective Communications Library 2.23

La biblioteca de comunicaciones colectivas de NVIDIA (NCCL) implementa primitivas de comunicación multi-GPU y multinodo optimizadas para redes y GPU NVIDIA. 

NCCL es una pieza central de software para el entrenamiento de aprendizaje profundo multi-GPU. Maneja cualquier tipo de comunicación entre GPU, ya sea a través de PCI, NVLink o redes. Utiliza detección de topología avanzada, gráficos de comunicación optimizados y modelos de ajuste para obtener el mejor rendimiento desde el primer momento en las plataformas de GPU NVIDIA.

En esta publicación, analizamos las nuevas funciones y correcciones lanzadas en NCCL 2.23. Consulta el repositorio de GitHub de NVIDIA/nccl . 

Aspectos destacados y características del lanzamiento 

NVIDIA Magnum IO NCCL es una biblioteca diseñada para optimizar la comunicación entre GPU y multinodo, algo fundamental para la computación paralela eficiente en aplicaciones de IA y computación de alto rendimiento (HPC). El valor de esta versión radica en sus nuevas características: 

• Nuevo algoritmo PAT para ReduceScatter y AllGather: Introducimos el algoritmo Parallel Aggregated Trees (PAT), basado en Brucks, para AllGather y ReduceScatter, consiguiendo escalamiento logarítmico. 
• Inicialización acelerada: rendimiento de inicialización mejorado, incluida la capacidad de utilizar redes en banda para la comunicación de arranque. 
• ncclCommInitRankScalable: Una nueva API de inicialización, para usar múltiples ncclUniqueIds para acelerar la inicialización a gran escala. 
• Registro de buffer de usuario intranodo: aproveche los buffers de usuario registrados para operaciones intranodo. 
• Nueva API del complemento de generación de perfiles: ganchos de API para medir el rendimiento de NCCL en detalle. 

Las siguientes secciones profundizan en los detalles de las nuevas funciones: 

• Escala logarítmica PAT para ReduceScatter y AllGather   
• Nueva ncclCommInitRankScalableAPI 
• Operaciones de arranque aceleradas 
• Registro de búfer de usuario intranodo 
• Nueva API del complemento de creación de perfiles 
• Corrección de errores y funciones menores

Escala logarítmica PAT para ReduceScatter y AllGather 

El algoritmo PAT es una variación del algoritmo Bruck, que presenta una cantidad logarítmica de pasos de red para tamaños pequeños a escala, aumentando progresivamente la cantidad de transferencias de red a medida que aumentan los tamaños, para mantener las necesidades de almacenamiento en búfer al mínimo. Se aplica tanto a AllGather como a ReduceScatter. Puede esperar que los tamaños de mensajes pequeños a medianos tengan un mejor rendimiento con PAT, y esta mejora aumentará a medida que su carga de trabajo se amplíe. 

Este algoritmo ejecuta un árbol binomial desplazado para cada rango. Su ventaja en comparación con algoritmos similares como la duplicación recursiva es que funciona en cualquier número de rangos y no requiere una potencia de dos. 

Inicialmente, PAT solo admite una GPU por nodo. El caso de una GPU por nodo ReduceScatter y AllGather es importante para el entrenamiento de modelos de lenguaje grandes (LLM) , donde el paralelismo de tuberías y el paralelismo de tensores se encuentran en dimensiones ortogonales al paralelismo de datos. La dimensión de paralelismo de tensores generalmente está alineada con la conectividad NVLink intranodo, lo que significa que otras dimensiones solo tendrán una GPU por nodo. 

Busque nuestro próximo artículo que describe los detalles del algoritmo.

Nueva ncclCommInitRankScalableAPI 

Esta característica agrega una nueva función de inicialización, ncclCommInitRankScalable, para permitir el uso de múltiples identificadores únicos durante la creación del comunicador. Esta incorporación evita los patrones de comunicación de todos a uno durante la inicialización y proporciona un rendimiento de inicialización más escalable. 

En la creación del comunicador, NCCL necesita obtener las direcciones de todos los rangos del comunicador (paso de arranque). Para ello, NCCL se basa en un identificador único conocido por todos los rangos. Durante el paso de arranque de la inicialización del comunicador, cada rango intercambia su dirección con el identificador único conocido, lo que crea un patrón de comunicación de todos a uno y un importante cuello de botella a escala. 

Con ncclCommInitRankScalable, el usuario ahora tiene la libertad de proporcionar más de una identificación única para usar durante el arranque. Para lograr la mayor ganancia, NCCL distribuirá la carga entre múltiples identificaciones únicas, lo que permitirá un tiempo de arranque constante a escala, si la cantidad de identificaciones únicas proporcionadas se ajusta al tamaño del comunicador. 

Esta nueva API requiere varios rangos para crear un ID único. Para obtener el mejor rendimiento, recomendamos distribuir los ID únicos de la forma más homogénea posible entre los rangos. 

Operaciones de arranque aceleradas

En la versión 2.23, mejoramos el rendimiento general del código de inicialización. Eliminamos algunos de los colectivos de arranque necesarios, así como también realizamos ajustes de rendimiento en el paso de arranque. 

Ahora puede utilizar la red rápida (IB/RoCE/…) para la comunicación fuera de banda a fin de acelerar los dos pasos lineales de la inicialización, bootstrap y allgather. Esta función está deshabilitada de manera predeterminada para evitar el uso de dispositivos configurados incorrectamente (el uso de ncclNetdispositivos ocurre antes de la detección de topología). Puede habilitarla con NCCL_OOB_NET_ENABLE=1. 

Además, puede especificar qué interfaz se debe utilizar con NCCL_OOB_NET_IFNAME. De manera predeterminada, NCCL utilizará el primer ncclNetdispositivo que encuentre en esa red.

Registro de búfer de usuario intranodo 

NCCL nunca requiere que usted, como usuario, registre y mantenga ningún búfer persistente para funcionar. Esta es una característica excelente para facilitar la usabilidad, pero conlleva desventajas en el rendimiento. Sin acceso directo, debe haber más flujo de control y almacenamiento en búfer cuando NCCL transfiere datos. Esto consume más recursos de la GPU y genera mayores gastos generales para mover la misma cantidad de datos en comparación con los búferes registrados y mapeados explícitamente.

Siempre que sea posible, se recomienda a los desarrolladores de NCCL que registren sus buffers ncclCommRegisterpara permitir que NCCL utilice todas las optimizaciones disponibles. El equipo de NCCL siempre está trabajando para agregar más casos de uso para buffers de usuario registrados. La versión 2.23 implementa soporte de registro de buffer de usuario (UB) intranodo para transportes P2P NvLink y PCIe.

El principal beneficio del registro de UB intranodo es evitar copias adicionales entre pares. Esto reduce la presión sobre el subsistema de memoria, mejora el rendimiento de la comunicación NCCL y también mejora la superposición de computación y comunicación. Se admiten todos los colectivos NCCL y las operaciones basadas en sendrecv, excepto ncclReducey ncclReduceScatter(no se beneficiarían). 

Hay dos formas de habilitar el registro de UB dentro de un nodo. La primera es registrar los buffers de ncclCommRegisterforma explícita, y los buffers se registrarán solo cuando se llamen a los colectivos NCCL correspondientes. La segunda es capturar operaciones NCCL a través de CUDA Graphs, y todos los buffers de usuario se registrarán automáticamente durante la captura del gráfico. Para obtener más pautas y requisitos, consulte la documentación de NCCL . 

Además de la comunicación entre nodos a través de NVLink y PCIe, la función funciona en sistemas NVLink multinodo (MNNVL) dentro de cada dominio NVLink.

Nueva API del complemento de creación de perfiles 

A medida que aumenta la escala de los clústeres de GPU, las anomalías de rendimiento se vuelven más difíciles de detectar y de determinar su causa. Se necesitan herramientas de diagnóstico y monitoreo específicas del dominio para recopilar y analizar datos de telemetría con una sobrecarga mínima para los trabajos en ejecución. La interfaz del complemento de perfilador NCCL se ha diseñado para abordar estas inquietudes. El diseño de la interfaz también facilita su adopción por parte de los perfiladores de marcos de trabajo de DL, como PyTorch Kineto. 

La nueva NCCL_PROFILER_PLUGINvariable de entorno controla la carga y la inicialización del complemento del generador de perfiles de la misma manera que se cargan e inicializan otros complementos de NCCL. Una vez cargado, el complemento del generador de perfiles puede habilitar la generación de perfiles de eventos de NCCL configurando la máscara de activación de eventos que NCCL expone al generador de perfiles durante la inicialización. La máscara de activación de eventos es un entero de 32 bits donde cada bit representa un evento del generador de perfiles de NCCL. Actualmente, NCCL admite los siguientes eventos: 

• ncclProfileGroup(bit-0): Evento de grupo 
• ncclProfileColl(bit-1): Evento colectivo 
• ncclProfileP2p(bit-2): Evento punto a punto 
• ncclProfileProxyOp(bit-3): Evento de canal de progreso del proxy 
• ncclProfileProxyStep(bit-4): Evento de paso de progreso del proxy 
• ncclProfileProxyCtrl(bit-5): Evento de estado interno de progreso del proxy 

NCCL expresa eventos en forma jerárquica. Por ejemplo, los colectivos se pueden agrupar y las operaciones de proxy ayudan a la GPU con la transferencia punto a punto de fragmentos de datos individuales a través de los canales de comunicación de red disponibles. Por lo tanto, NCCL presenta los eventos correspondientes al generador de perfiles, preservando esta relación. A continuación se muestra un diagrama de la jerarquía de eventos de NCCL:

ncclProfileGroup | +- ncclProfileColl |  | |  +- ncclProfileProxyOp |   | |   +- ncclProfileProxyStep | +- ncclProfileP2p    |    +- ncclProfileProxyOp     |     +- ncclProfileProxyStep  ncclProfileProxyCtrl

Esta representación jerárquica permite que los complementos de generación de perfiles presenten eventos a los usuarios de una forma más significativa y comprensible. 

NCCL también proporciona un complemento de creación de perfiles de ejemplo en el ext-profiler/exampledirectorio que puede usarse como plantilla para desarrollar complementos de creación de perfiles de terceros. 

En total, la interfaz del complemento del generador de perfiles define las siguientes cinco devoluciones de llamadas de función: 

ncclResult_t (*init)(   void** context,   int* eActivationMask);  ncclResult_t (*startEvent)(   void* context,   void** eHandle,   ncclProfilerEventDescr_t* eDescr);  ncclResult_t (*stopEvent)(   void* eHandle);  ncclResult_t (*recordEventState)(   void* eHandle,   ncclProfilerEventState_t eState,   NcclProfilerEventStateArgs_t* eStateArgs);  ncclResult_t (*finalize)(void* context);

La initfunción del generador de perfiles toma un puntero de máscara de activación de eventos y devuelve un objeto de contexto opaco a NCCL. El contexto proporciona aislamiento entre instancias del generador de perfiles, mientras que el generador de perfiles utiliza la máscara de activación de eventos para notificar a NCCL sobre qué eventos deben perfilarse; por ejemplo, estableciendo *eActivationMask = ncclProfileColl | ncclProfileProxyOp. 

La startEvent función del generador de perfiles toma un contexto de generador de perfiles y un descriptor de eventos. El generador de perfiles utiliza la información del descriptor para asignar un nuevo objeto de evento e inicializarlo. Después, el generador de perfiles devuelve un identificador opaco que NCCL puede utilizar para realizar otras operaciones en el evento; por ejemplo, actualizaciones de estado de registros. 

La stopEvent función del generador de perfiles toma un identificador de evento y marca el evento como completado. Luego, el identificador de evento ya no se puede usar (el generador de perfiles podría reciclar internamente el objeto correspondiente para eventos futuros). 

La recordEventState función del generador de perfiles toma un identificador de evento, un estado de evento y (opcionalmente) un objeto de argumento de estado de evento. Esta función permite que el generador de perfiles actualice eventos que pueden pasar por diferentes estados en NCCL. Un ejemplo son los eventos de proxy, donde el proxy necesita coordinarse tanto con la GPU como con la red mientras transfiere datos, pasando de un estado a otro en el proceso. 

La finalize función del generador de perfiles toma el contexto del generador de perfiles y libera todos los recursos asociados a él.

Corrección de errores y funciones menores 

NCCL 2.23 proporciona las siguientes actualizaciones adicionales: 

• La asignación de gráficos asincrónica hace que las llamadas a la asignación de gráficos cudaMallocy cudaMemcpydurante la misma sean asincrónicas. Acelera significativamente la captura de gráficos. 
• El uso de eventos asincrónicos IB fatales para detener las operaciones de red ayuda a detectar errores de enlaces caídos y otros eventos asincrónicos fatales dentro de NCCL. 
• Establezca el nivel P2P en PXB en las CPU AMD cuando utilice más de dos GPU por nodo. 
• Mejorar los initregistros para informar la función NCCL real: informa al usuario si NCCL se está ejecutando ncclCommInitRanko ncclCommSplit. 
• Agregar NCCL_CONF_FILEvariable 
• Aumentar el tiempo de espera predeterminado del IB de 18 a 20 
• Se agregó una nueva verificación para la memoria de pares de NVIDIA. Funciona con los núcleos Linux más recientes. 
• Corrija la regresión de rendimiento anterior al combinar operaciones pequeñas y grandes. 
• Se soluciona el fallo cuando los ID de NUMA son iguales a -1. 
• Se corrige la búsqueda del gráfico de árbol cuando NCCL_CROSS_NICse establece en 1. 

Resumen

NVIDIA NCCL 2.23 presenta nuevas funciones y mejoras para optimizar la comunicación entre GPU y entre nodos, algo fundamental para las aplicaciones de IA y HPC. Las mejoras clave incluyen el nuevo algoritmo PAT, la inicialización acelerada a escala, el registro de búfer de usuario dentro de un nodo y la nueva API del complemento de generación de perfiles. 

Para obtener más información sobre la versión anterior, consulte Eficiencia de memoria, inicialización más rápida y estimación de costos con NVIDIA Collective Communications Library 2.22 .

Obtenga más información sobre Magnum IO y NCCL . NVIDIA Blog. S. J., G. C. T. G., B. W. y G. O. Traducido al español