Skip to content

English · Español

Lab 00 — Hola, CUDA: comprobación del toolchain

Objetivo: verificar que el toolchain de la GPU en la nube funciona de extremo a extremo. Escribir un kernel CUDA trivial (suma vectorial), compilar, lanzar, validar contra NumPy. La idea es aflorar problemas de tooling antes de que los labs más duros dependan de ellos.

Tiempo estimado: 1–2 horas (la mayor parte es entorno).

Prerrequisito: Fase 23 completa. Instancia de GPU en la nube disponible (alquilada según el ritual de phase-23/lab/00-provision-cloud-gpu.md). cupy o cuda-python instalado.

Lo que produces

Un directorio experiments/24-hello-cuda/ que contiene:

  • vec_add.cu — el código fuente del kernel.
  • run.py — lanzamiento + comprobación de corrección.
  • manifest.json — ver plantilla más abajo.
  • README.md — 1 párrafo: qué ruta de build (cupy.RawKernel o cuda-python), versión del driver, escollos encontrados.

El kernel

__global__ void vec_add(const float* a, const float* b, float* c, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) c[tid] = a[tid] + b[tid];
}

TODOs

Bloque A — elegir la ruta de build

  • Decide: cupy.RawKernel (la más simple; JIT de Python a PTX) o cuda-python con nvrtc (más flexible; más cercana a producción). Por defecto: cupy.RawKernel.
  • Instala la ruta elegida vía uv pip install. Fija la versión en pyproject.toml.
  • Verifica que nvidia-smi muestra tu GPU. Verifica que nvcc --version o la ruta JIT funciona.

Bloque B — lanzar el kernel

  • run.py: reserva dos arrays host de tamaño \(N = 2^{20}\) con np.random.default_rng(42).
  • Copia al device (cupy.asarray o equivalente).
  • Lanza con <<<grid, 256>>> donde grid = (N + 255) // 256.
  • Copia el resultado de vuelta al host. Compara con a + b de NumPy. Aserta igualdad exacta (la suma en float32 es lo bastante asociativa para una sola operación).

Bloque C — cronometrar (a la ligera)

  • Cronometra el kernel con eventos CUDA (NO con time.perf_counter de Python — el lanzamiento es asíncrono; estarías midiendo la submission, no la ejecución).
  • Imprime el ancho de banda alcanzado: \(3 \cdot N \cdot 4\) bytes / tiempo. Compara con el pico de HBM (nvidia-smi -q | grep "Memory Bus" o la spec publicada de la GPU en la nube).
  • Esperado: la suma vectorial está limitada por ancho de banda; debería alcanzar un 60–80% del pico de HBM. Si está <20%, algo va mal (probablemente cudaMemcpy está dentro de la ventana de cronometraje).

Bloque D — manifest

{
  "experiment": "24-hello-cuda",
  "date": "YYYY-MM-DD",
  "seed": 42,
  "gpu": {"model": null, "compute_capability": null, "driver": null},
  "build_path": "cupy.RawKernel",
  "versions": {"python": "3.11.x", "cupy": null, "cuda_runtime": null},
  "results": {
    "N": 1048576,
    "kernel_time_us": null,
    "achieved_bandwidth_gbs": null,
    "peak_bandwidth_gbs": null,
    "fraction_of_peak": null,
    "correctness": "passed | failed"
  }
}

Restricciones

  • No cronometres cudaMemcpy dentro del tiempo del kernel. Usa eventos CUDA que rodeen únicamente el lanzamiento del kernel.
  • No avances si la corrección falla. La suma vectorial es el kernel más simple posible; si está mal, el toolchain o la sintaxis del lanzamiento están rotos.
  • No tunees. Esto es una comprobación de toolchain, no un benchmark de rendimiento. ≥20% del pico es de sobra.

Condiciones de parada

Hecho cuando:

  1. run.py se ejecuta de extremo a extremo en la GPU en la nube.
  2. La aserción de corrección pasa.
  3. El ancho de banda alcanzado es ≥20% del pico (chequeo de cordura).
  4. manifest.json commiteado.
  5. La facturación de la instancia en la nube registrada en tu log de costes de la fase 23.

Escollos

  • Overhead de la primera llamada a cupy.RawKernel. La compilación JIT ocurre en la primera llamada (~1–2 s). Cronometra la segunda llamada. O usa cp.cuda.compile_with_cache(...).
  • Desajuste de versión de cupy con el CUDA runtime instalado. cupy-cuda12x no es la misma wheel que cupy-cuda11x. Hazla coincidir con el driver.
  • Off-by-one en grid. N // 256 se pierde la cola cuando N % 256 != 0. Usa (N + 255) // 256 y la guarda if (tid < N).
  • Carrera en el puntero del device. Devolver el resultado antes de cudaDeviceSynchronize o de sincronizar el stream da datos rancios. cupy sincroniza en .get(); verifícalo.

Cuándo consultar solutions/

Tras cumplir todas las condiciones de parada. La referencia recorre el montaje de cupy.RawKernel y un cronometraje conocido-bueno de la suma vectorial.

Siguiente lab: lab/01-naive-kernel.md.