La biblioteca de plantillas C++ CUB es fundamental para algoritmos primitivos de GPU de alto rendimiento. Sin embargo, su API tradicional de «dos fases», que separa la estimación de memoria de la asignación, puede resultar engorrosa. Este modelo de programación, aunque ofrece flexibilidad, a menudo genera código repetitivo.
Este artículo detalla la transición de esta API a la nueva API de llamada única de CUB, presentada en CUDA 13.1, que simplifica el desarrollo al gestionar la memoria de manera interna sin comprometer el rendimiento.
¿Qué es CUB?
Si necesitas ejecutar un algoritmo estándar (como escaneo, histograma o clasificación) en una GPU, CUB es probablemente la manera más rápida de hacerlo. Como componente principal de las NVIDIA CUDA Core Compute Libraries (CCCL), CUB está diseñado para abstraer la complejidad del manejo manual de hilos CUDA sin sacrificar rendimiento.
A diferencia de bibliotecas como Thrust, que ofrecen una interfaz de alto nivel, CUB proporciona un conjunto de primitivas “del lado del dispositivo”. Esto permite a los desarrolladores integrar algoritmos altamente optimizados directamente en sus propios núcleos personalizados. Para aprender a utilizar CUB, consulta el curso de DLI de NVIDIA Fundamentals of Accelerated Computing with Modern CUDA C++.
La API de dos fases de CUB
CUB es ampliamente recomendado para aprovechar al máximo las capacidades computacionales de las GPUs de NVIDIA. Sin embargo, su uso puede presentar complejidades que resultan no triviales. Este apartado ofrece una perspectiva sobre estos mecanismos subyacentes.
Se asume un flujo de ejecución sencillo, donde una única llamada a una función es suficiente para ejecutar el algoritmo y obtener los resultados inmediatamente. Se espera que los efectos secundarios de la función, como la modificación de una variable o el retorno de un resultado, sean visibles de inmediato.
El modelo de ejecución de CUB se aparta de este patrón familiar. Invocar una primitiva de CUB es un proceso de dos pasos: primero, calcular el tamaño necesario de memoria del dispositivo, y segundo, asignar explícitamente y ejecutar el núcleo.
El siguiente es un llamado convencional a CUB:
// PRIMER LLAMADO: determinar tamaño de almacenamiento temporal cub::DeviceScan::ExclusiveSum(nullptr, temp_storage_bytes, d_input, d_output, num_items); // Asignar el almacenamiento temporal requerido cudaMalloc(&d_temp_storage, temp_storage_bytes); // SEGUNDO LLAMADO: ejecutar el escaneo cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, d_input, d_output, num_items);
La interfaz de CUB presenta un desafío práctico. Las primitivas deben ser invocadas dos veces: primero para determinar la cantidad de memoria temporal necesaria y luego para ejecutar el algoritmo con el almacenamiento asignado.
Una desventaja significativa de la API de dos fases tradicional es la falta de claridad sobre qué argumentos deben permanecer consistentes entre los pasos de estimación y ejecución. Tomando el fragmento anterior como referencia, no está claro cuáles parámetros influyen en el estado interno y pueden cambiar entre las llamadas, ya que las firmas de función para ambas fases son idénticas. Por ejemplo, los argumentos d_input y d_output solo se utilizan en la segunda llamada.
La nueva API de llamada única de CUB
Dada la amplia utilización de wrappers en muchos códigos de producción, se ha reconocido la necesidad de extender CUB con la nueva API de llamada única:
// LLAMADA ÚNICA: asignación y ejecución en un solo paso cub::DeviceScan::ExclusiveSum(d_input, d_output, num_items);
Este ejemplo muestra que no se requiere asignación de memoria explícita. Sin embargo, el proceso de asignación aún ocurre internamente. La Figura 1 demuestra que la interfaz de llamada única, que incluye la estimación de almacenamiento temporal, la asignación de memoria y la invocación del algoritmo, no introduce sobrecarga en comparación con la API de dos fases.
La figura 1 compara el tiempo de ejecución en GPU del llamado original de ExclusiveSum de dos fases con el nuevo llamado de una fase. En el eje x se representan múltiples tamaños de entrada, mientras que en el eje y se muestra el tiempo de ejecución normalizado para cada tipo de invocación. Se pueden extraer dos conclusiones principales de estos datos de rendimiento:
- La nueva API no introduce sobrecarga
- La asignación de memoria se mantiene bajo la nueva API; simplemente ocurre internamente
El segundo punto se puede verificar al inspeccionar la implementación de la nueva API. La asignación asíncrona está incrustada dentro de la primitiva del dispositivo:
cub::DeviceScan::ExclusiveSum(d_input, d_output, num_items, env = {}) { . . . d_temp_storage = mr.allocate(stream, bytes); mr.deallocate(stream, d_temp_storage, bytes); . . . }
Las APIs de dos fases no han sido eliminadas; estas siguen siendo llamadas válidas de las APIs existentes de CUB. Más bien, se han añadido las llamadas de una fase a las APIs existentes. Se espera que la mayoría de los usuarios utilicen estas nuevas llamadas.
El entorno y los recursos de memoria
Más allá de resolver los problemas mencionados, la nueva API de llamada única de CUB también amplía las capacidades de configuración de ejecución de la primitiva invocada. Introduce un argumento de entorno, que puede personalizar la asignación de memoria utilizando recursos de memoria o simplemente proporcionar un flujo para ejecutar (como la API de dos fases).
Los recursos de memoria son una nueva utilidad de memoria para asignar y liberar memoria. El argumento de entorno en las APIs de llamada única puede opcionalmente contener un recurso de memoria. Cuando no se proporciona un recurso de memoria, la API utilizará un recurso de memoria predeterminado proporcionado por CCCL. Por el contrario, puedes optar por pasar uno de los recursos de memoria no predeterminados proporcionados como parte de la base de código, o incluso pasar tu propio recurso de memoria personalizado.
Por ejemplo:
// Usar tipo de recurso de memoria proporcionado por CCCL cuda::device_memory_pool mr{cuda::devices[0]}; cub::DeviceScan::ExclusiveSum(d_input, d_output, num_items, mr);
// Crear y usar tu recurso de memoria personalizado my_memory_resource my_mr{cuda::experimental::devices[0]}; // Usarlo con CUB cub::DeviceScan::ExclusiveSum(d_input, d_output, num_items, my_mr);
Con la nueva API, el manejo del flujo de ejecución CUDA no se elimina, sino que se encapsula dentro de la nueva variable env. Por supuesto, también se puede pasar explícitamente como antes, incluso si se elimina el manejo de asignación temporal. CUB ahora también proporciona cuda::stream_ref que es seguro en cuanto a tipos y su uso debe ser preferido. También puedes pasar cuda::stream, que posee el flujo de ejecución subyacente.
Combinando opciones de ejecución
La API de llamada única permite más que solo pasar un recurso de memoria o un flujo como último argumento. En el futuro, el argumento de entorno será el lugar para todos los ajustes relacionados con la ejecución, incluyendo requisitos deterministas, garantías, ajustes definidos por el usuario y mucho más.
Con la introducción de la API de llamada única, CUB ha desbloqueado un conjunto masivo de características de configuración de ejecución. Con la multitud de nuevas características de ejecución, la pregunta es: ¿cuál es la mejor manera de combinarlas todas?
La solución radica en el nuevo argumento env. Al aprovechar cuda::std::execution, CUB proporciona un punto de control central que actúa como un flexible “panel de control” para tu algoritmo. En lugar de argumentos de función rígidamente definidos, el entorno te permite crear una mezcla combinatoria de cualquier característica que necesites. Ya sea que desees emparejar un flujo personalizado con un grupo de memoria específico, o combinar requisitos deterministas estrictos con una política de ajuste personalizada, el argumento env lo maneja todo en un solo objeto seguro en cuanto a tipos.
cuda::stream custom_stream{cuda::device_ref{0}}; auto memory_prop = cuda::std::execution::prop{cuda::mr::get_memory_resource, cuda::device_default_memory_pool(cuda::device_ref{0})}; auto env = cuda::std::execution::env{custom_stream.get(), memory_prop}; DeviceScan::ExclusiveSum(d_input, d_output, num_items, env);
CUB actualmente proporciona los siguientes algoritmos que soportan la interfaz de entorno, con más por venir:
- cub::DeviceReduce::Reduce
- cub::DeviceReduce::Sum
- cub::DeviceReduce::Min/Max/ArgMin/ArgMax
- cub::DeviceScan::ExclusiveSum
- cub::DeviceScan::ExclusiveScan
Para el progreso actualizado de las sobrecargas basadas en el nuevo entorno, consulta el issue de seguimiento de primitivas de CUB en el repositorio de GitHub de NVIDIA/cccl.
Comienza con CUB
Al reemplazar el patrón verbose de dos fases con una interfaz de llamada única más ágil, CUB ofrece una API moderna que elimina el código repetitivo sin añadir sobrecarga. Al aprovechar el argumento extensible env, obtienes un panel de control unificado para mezclar sin problemas recursos de memoria, flujos y otras facilidades. Se te anima a adoptar este nuevo estándar para simplificar tu base de código y aprovechar al máximo el poder computacional de tu GPU. Descarga CUDA 13.1 o posterior y comienza a utilizar estas APIs de llamada única.
