Las NVIDIA CUDA Core Compute Libraries (CCCL) entregan abstracciones convenientes y eficientes para desarrolladores CUDA en C++ y Python. Sus funcionalidades incluyen:

  • Algoritmos paralelos: algoritmos lanzados desde el host (sort, scan, reduce) que eliminan la necesidad de escribir kernels personalizados para operaciones comunes.
  • Algoritmos cooperativos: algoritmos del lado del dispositivo, como reducciones o scans a nivel de bloque o warp, que simplifican el desarrollo de kernels personalizados.
  • Abstracciones CUDA idiomáticas: abstracciones fundamentales para operaciones específicas de CUDA, incluyendo asignación de memoria, gestión de recursos y funcionalidades de hardware.

Este post presenta un nuevo grupo de funcionalidad dentro de CCCL que entrega abstracciones C++ modernizadas para conceptos fundamentales del modelo de programación CUDA, haciendo el desarrollo en CUDA C++ más seguro y conveniente.

¿Qué es CCCL Runtime?

NVIDIA CCCL Runtime es un nuevo conjunto de APIs idiomáticas en C++ que implementan funcionalidad central de CUDA: gestión de streams, asignación de memoria, lanzamiento de kernels y más.

El familiar CUDA Runtime fue desarrollado originalmente como una capa de conveniencia sobre la CUDA Driver API. El nuevo CCCL Runtime apunta a ser una alternativa con el mismo objetivo, pero con un diseño actualizado alineado con C++ moderno.

Figura 1. Diagrama de stack de las distintas API surfaces de CUDA
Figura 1. Diagrama de stack de las distintas API surfaces de CUDA

CCCL Runtime es una colección de headers dentro de CCCL, como <cuda/stream>, <cuda/buffer> y <cuda/launch>. Aprovecha las funcionalidades modernas de C++ para entregar abstracciones más convenientes y robustas de lo que era posible dentro de las restricciones de compatibilidad de fuente C de la CUDA Runtime API tradicional.

También aprovecharon la oportunidad para incorporar lecciones aprendidas durante más de 20 años de evolución de CUDA en el diseño de la API. Aun con todos esos cambios, CCCL Runtime entrega helpers de compatibilidad que permiten a los desarrolladores adoptarlo incrementalmente sin reescribir el código alrededor que usa la CUDA Runtime API.

A medida que los programas CUDA se vuelven más complejos, con múltiples librerías compartiendo dispositivos, streams y memoria, se hace más urgente la necesidad de APIs que compongan limpiamente y hagan las dependencias explícitas. Ese es el espacio que CCCL Runtime busca llenar.

¿Cómo se ve el código?

Aquí está el clásico ejemplo vectorAdd implementado con las nuevas APIs de CCCL Runtime. Si ha escrito CUDA antes, la estructura general le resultará familiar.

C++
#include <cuda/buffer>
#include <cuda/devices>
#include <cuda/launch>
#include <cuda/memory_pool>
#include <cuda/std/span>
#include <cuda/stream>

struct kernel {
  template <typename Config>
  __device__ void operator()(Config config,
                             cuda::std::span<const int> A,
                             cuda::std::span<const int> B,
                             cuda::std::span<int> C) {
    auto tid = cuda::gpu_thread.rank(cuda::grid, config);
    if (tid < A.size())
      C[tid] = A[tid] + B[tid];
  }
};

int main() {
  // 1. Devices y streams
  cuda::device_ref device = cuda::devices[0];
  cuda::stream stream{device};

  // 2. Allocation de memoria
  auto pool = cuda::device_default_memory_pool(device);

  int num_elements = 1000;
  auto A = cuda::make_buffer<int>(stream, pool, num_elements, 1);
  auto B = cuda::make_buffer<int>(stream, pool, num_elements, 2);
  auto C = cuda::make_buffer<int>(stream, pool, num_elements, cuda::no_init);

  // 3. Kernel launch
  constexpr int threads_per_block = 256;
  auto config = cuda::distribute<threads_per_block>(num_elements);

  cuda::launch(stream, config, kernel{}, A, B, C);

  stream.sync();
  return 0;
}

El ejemplo se puede desglosar en tres secciones principales.

1.) Devices y streams

Consideremos la creación de un stream usando la CUDA Runtime API tradicional:

C++
cudaStream_t stream;
cudaStreamCreate(&stream); // asociado al device "current"

Esto crea un stream, pero el stream queda asociado al device que esté activo en el momento de llamar a cudaStreamCreate. Basándose solo en esa llamada, no se sabe a qué device está asociado el stream.

Contrastemos con la nueva CCCL Runtime API:

C++
cuda::device_ref device = cuda::devices[0];
cuda::stream stream{device};

La primera línea ilustra un principio central del diseño: CCCL Runtime usa tipos dedicados en lugar de identificadores crudos. Un device es un device_ref, no un entero pelado; un stream es un objeto, no un puntero opaco. La tipificación fuerte ayuda a atrapar errores en tiempo de compilación en lugar de cazarlos en runtime.

La segunda línea muestra otro principio: hacer explícitas las dependencias. En ambas APIs un stream se asocia a un device. La diferencia es cómo: aquí el constructor de cuda::stream toma el device como argumento explícito, mientras que con la CUDA Runtime API el stream queda asociado al device que esté activo cuando el stream se crea.

Las dependencias explícitas habilitan razonamiento local: se puede leer una función y entender qué hace sin trackear el estado global. También mejoran la composabilidad: cuando se usan múltiples librerías, ninguna necesita guardar y restaurar estado implícito entre llamadas para evitar interferir con otras.

Una consecuencia relacionada es que CCCL Runtime no expone el default stream. Manejar el significado del default stream requiere trackear el device actual, exactamente el tipo de estado implícito del que se quiere salir. Si bien un default stream de la CUDA Runtime API todavía puede ser envuelto en tipos CCCL Runtime, se desalienta su uso. Como no hay default stream en la API, la noción de "blocking stream" ya no aplica, así que todos los streams CCCL Runtime se crean como non-blocking.

Propiedad de recursos: tipos con dueño y refs

Siguiendo el ejemplo de std::string y std::string_view, muchos objetos CUDA tienen dos tipos en CCCL Runtime: un tipo con dueño y un tipo no-dueño con sufijo _ref. cuda::stream posee el handle subyacente cudaStream_t y lo destruye en su destructor. El cuda::stream_ref mantiene el handle sin gestionar su lifetime y es trivialmente copiable.

Los tipos _ref son esenciales para la composabilidad con código existente. Si el lifetime del handle de un stream se gestiona en otro lado, cudaStream_t convierte implícitamente a cuda::stream_ref y el handle crudo se puede recuperar con .get(). Para transferir la propiedad, cuda::stream::from_native_handle envuelve un handle crudo en el tipo con dueño y .release() devuelve la propiedad.

C++
void stream_type_example(cudaStream_t handle) {
  cuda::stream_ref non_owning{handle};
  assert(handle == non_owning.get());

  cuda::stream owning = cuda::stream::from_native_handle(handle);
  assert(handle == owning.get());
  assert(handle == owning.release());
}

El mismo patrón aplica a eventos, memory pools y otros objetos CUDA. cuda::device_ref no tiene contraparte con dueño porque no hay estado de device que poseer.

2.) Asignación de memoria

C++
auto pool = cuda::device_default_memory_pool(device);

auto A = cuda::make_buffer<int>(stream, pool, num_elements, 1);
auto B = cuda::make_buffer<int>(stream, pool, num_elements, 2);
auto C = cuda::make_buffer<int>(stream, pool, num_elements, cuda::no_init);

La siguiente sección demuestra cómo asignar e inicializar memoria de device de forma asíncrona. Aquí aparece el siguiente principio de diseño: las APIs son asíncronas por defecto. En lugar de distinguir variantes síncronas y asíncronas por nombre, CCCL Runtime usa una convención simple: si una API recibe un stream como primer argumento, opera en stream order. No planean entregar contrapartes síncronas para APIs que tienen ambas variantes en la CUDA Runtime API.