cuda-core 1.0 — пишем CUDA-ядра на Python без C++ (ну почти)

11 мая 2026 года NVIDIA выпустила cuda-core v1.0.0 — первый стабильный релиз библиотеки, которая даёт Python-разработчикам прямой доступ к CUDA Runtime без тяжелых C++ обвязок.

Мы взяли 3 видеокарты (4090, 3090, A100 80Gb) и протестировали работу библиотеки на каждой.

cuda-core — это Pythonic-обёртка над CUDA Runtime. Она закрывает ту нишу, которую раньше занимали pycuda или ручные вызовы через ctypes (компиляция ядер прямо из Python, управление памятью на GPU, запуск ядер без C++ расширений). Версия 1.0.0 фиксирует публичный API — теперь можно применять библиотеку в продакшн-зависимостях.

В экосистеме CUDA Python два слоя:

Для кого полезно:

• Исследователи, прототипирующие нестандартные операции для обучения моделей

• Инференс-инженеры, оптимизирующие горячие пути под конкретное железо

• Авторы ML-библиотек, которым нужен GPU-доступ без C++ build system

Что нового в v1.0.0

Кэш компиляции ядер

Компиляция CUDA-ядра через NVRTC занимает 10–50 мс. При каждом старте приложения это складывается в секунды. v1.0.0 добавляет два класса:

• InMemoryProgramCache — in-process LRU-кэш на время жизни процесса

• FileStreamProgramCache — персистентный кэш на диске между запусками

from cuda.core import Device, Program, ProgramOptions, ObjectCodefrom cuda.core.utils import InMemoryProgramCached = Device(0)d.set_current()cuda_src = r"""extern "C" __global__ void saxpy(    float alpha, const float* __restrict__ x,    const float* __restrict__ y, float* out, int n) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < n) out[idx] = alpha * x[idx] + y[idx];}"""arch = f'sm_{d.compute_capability.major}{d.compute_capability.minor}'cache = InMemoryProgramCache()cache_key = f'saxpy_{arch}'# Первый запуск — компилируем и кладём в кэш (~15 мс)if cache.get(cache_key) is None:    prog = Program(cuda_src, 'c++', options=ProgramOptions(arch=arch))    cache[cache_key] = prog.compile('cubin')# Все последующие — мгновенно (~0.05 мс, 300× быстрее)kernel = ObjectCode.from_cubin(cache.get(cache_key)).get_kernel('saxpy')

Green Contexts: делим GPU между задачами

Green Contexts — механизм разделения SM-ресурсов GPU между задачами без физической изоляции. RTX 4090 имеет 128 SM; можно нарезать их на партиции и запускать задачи параллельно с гарантированными ресурсами для каждой.

Типичный сценарий: два независимых инференс-запроса должны работать на одном GPU без взаимного вытеснения за SM. Без Green Contexts задачи конкурируют за ресурс неконтролируемо. С ними — каждая получает свой выделенный кусок.

import numpy as npimport ctypesimport threadingfrom cuda.core import (    Device, SMResourceOptions, ContextOptions,    Program, ProgramOptions, LaunchConfig, launch, ObjectCode,    DeviceMemoryResource, DeviceMemoryResourceOptions)from cuda.core.utils import InMemoryProgramCachefrom cuda.bindings import runtime as cudartN = 5_000_000d = Device(0)d.set_current()total_sm = d.resources.sm.sm_countprint(f"GPU: {d.name}, Total SMs: {total_sm}")  # Total SMs: 128# Делим 128 SM на два раздела по 64sm_opts = SMResourceOptions(count=[total_sm // 2, total_sm // 2])groups, _ = d.resources.sm.split(sm_opts)ctx_a = d.create_context(ContextOptions([groups[0]]))ctx_b = d.create_context(ContextOptions([groups[1]]))print(f"Context A: {ctx_a.resources.sm.sm_count} SMs, is_green={ctx_a.is_green}")print(f"Context B: {ctx_b.resources.sm.sm_count} SMs, is_green={ctx_b.is_green}")# Context A: 64 SMs, is_green=True# Context B: 64 SMs, is_green=Truecuda_src = r"""extern "C" __global__ void saxpy(    float alpha, const float* __restrict__ x,    const float* __restrict__ y, float* out, int n) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < n) out[idx] = alpha * x[idx] + y[idx];}"""arch = f"sm_{d.compute_capability.major}{d.compute_capability.minor}"cache = InMemoryProgramCache()cache_key = f"saxpy_{arch}"if cache.get(cache_key) is None:    prog = Program(cuda_src, "c++", options=ProgramOptions(arch=arch))    cache[cache_key] = prog.compile("cubin")def run_task(ctx, task_name, alpha, fill_x, fill_y):    d.set_current()  # нужен primary context в каждом потоке    # stream, созданный через green context, привязан к его SM-партиции    stream = ctx.create_stream()    mr = DeviceMemoryResource(0, DeviceMemoryResourceOptions())    kernel = ObjectCode.from_cubin(cache.get(cache_key)).get_kernel("saxpy")    size = N * 4  # float32    buf_x = mr.allocate(size, stream=stream)    buf_y = mr.allocate(size, stream=stream)    buf_o = mr.allocate(size, stream=stream)    x_h = np.full(N, fill_x, dtype=np.float32)    y_h = np.full(N, fill_y, dtype=np.float32)    cudart.cudaMemcpy(buf_x.handle, x_h.ctypes.data, size, cudart.cudaMemcpyKind.cudaMemcpyHostToDevice)    cudart.cudaMemcpy(buf_y.handle, y_h.ctypes.data, size, cudart.cudaMemcpyKind.cudaMemcpyHostToDevice)    cfg = LaunchConfig(grid=(N + 255) // 256, block=256)    launch(stream, cfg, kernel,           np.float32(alpha), buf_x.handle, buf_y.handle, buf_o.handle, ctypes.c_int(N))    stream.sync()    out_h = np.empty(N, dtype=np.float32)    cudart.cudaMemcpy(out_h.ctypes.data, buf_o.handle, size, cudart.cudaMemcpyKind.cudaMemcpyDeviceToHost)    print(f"  {task_name}: result={out_h[0]:.1f}")    buf_x.close(); buf_y.close(); buf_o.close()    stream.close()# Запускаем два ядра параллельно — каждое на своих 64 SMt_a = threading.Thread(target=run_task, args=(ctx_a, "Task-A (64 SM)", 3.0, 2.0, 1.0))t_b = threading.Thread(target=run_task, args=(ctx_b, "Task-B (64 SM)", 5.0, 4.0, 2.0))t_a.start(); t_b.start()t_a.join(); t_b.join()# Task-A (64 SM): result=7.0# Task-B (64 SM): result=22.0ctx_a.close(); ctx_b.close()

Где доступно: Ada Lovelace (RTX 40xx, sm_89) и новее, CUDA 12.4+.

Расширенный NVML

Модуль system получил GPU-мониторинг в реальном времени:

from cuda.core import system, Deviceprint(f"Devices: {system.get_num_devices()}")d = Device(0)d.set_current()sd = d.to_system_device()util = sd.utilizationmem  = sd.memory_infoprint(f"{sd.name}: GPU {util.gpu}%, MEM {util.memory}%")print(f"Memory: {mem.used // 10242} MB / {mem.total // 10242} MB")

Также добавлены: MIG-режим, NVLink (версия, состояние), список запущенных compute-процессов с потреблением памяти.

Ускорение StridedMemoryView для PyTorch

StridedMemoryView получил fast path через AOT Inductor — 7–20× более быстрое построение view для PyTorch-тензоров без копирования данных.

Полный рабочий пример: SAXPY на GPU

Полный цикл: инициализация → компиляция → память → запуск → проверка. Протестировано на RTX 4090 / CUDA 13.0.

import numpy as npimport ctypesfrom cuda.core import (    Device, DeviceMemoryResource, DeviceMemoryResourceOptions,    Program, ProgramOptions, LaunchConfig, launch, ObjectCode)from cuda.core.utils import InMemoryProgramCachefrom cuda.bindings import runtime as cudartN = 10_000_000FLOAT_BYTES = 4# 1. Устройство и стримd = Device(0)d.set_current()stream = d.create_stream()print(f"GPU: {d.name}, SM: {d.properties.multiprocessor_count}")# 2. Memory poolmr = DeviceMemoryResource(0, DeviceMemoryResourceOptions())# 3. Kernel: SAXPY = alpha * X + Ycuda_src = r"""extern "C" __global__ void saxpy(    float alpha, const float* __restrict__ x,    const float* __restrict__ y, float* out, int n) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;    if (idx < n) out[idx] = alpha * x[idx] + y[idx];}"""# 4. Компиляция с кэшемarch = f'sm_{d.compute_capability.major}{d.compute_capability.minor}'cache = InMemoryProgramCache()cache_key = f'saxpy_{arch}'if cache.get(cache_key) is None:    prog = Program(cuda_src, 'c++', options=ProgramOptions(arch=arch))    cache[cache_key] = prog.compile('cubin')kernel = ObjectCode.from_cubin(cache.get(cache_key)).get_kernel('saxpy')# 5. GPU-памятьsize_bytes = N * FLOAT_BYTESbuf_x   = mr.allocate(size_bytes, stream=stream)buf_y   = mr.allocate(size_bytes, stream=stream)buf_out = mr.allocate(size_bytes, stream=stream)# 6. H→Dx_host = np.ones(N, dtype=np.float32) * 2.0y_host = np.ones(N, dtype=np.float32) * 1.0alpha  = np.float32(3.0)cudart.cudaMemcpy(buf_x.handle, x_host.ctypes.data, size_bytes, cudart.cudaMemcpyKind.cudaMemcpyHostToDevice)cudart.cudaMemcpy(buf_y.handle, y_host.ctypes.data, size_bytes, cudart.cudaMemcpyKind.cudaMemcpyHostToDevice)# 7. Запускcfg = LaunchConfig(grid=(N + 255) // 256, block=256)launch(stream, cfg, kernel, alpha, buf_x.handle, buf_y.handle, buf_out.handle, ctypes.c_int(N))# 8. D→H + проверкаstream.sync()out_host = np.empty(N, dtype=np.float32)cudart.cudaMemcpy(out_host.ctypes.data, buf_out.handle, size_bytes, cudart.cudaMemcpyKind.cudaMemcpyDeviceToHost)assert np.allclose(out_host, 3.0 * x_host + y_host)print(f"SAXPY OK: {out_host[0]:.1f}")  # 7.0buf_x.close(); buf_y.close(); buf_out.close()stream.close()

Время ядра на 10M элементов — ~0.13 мс, эффективная полоса пропускания ~920 GB/s (91% от пика RTX 4090).

Полный код примеров и бенчмарков: github.com/IntelionCloud/research-examples/cuda-core/

Benchmark: пропускная способность памяти

Реальные измерения SAXPY на cuda-core v1.0.0 (100M float32 элементов, 5 прогонов, warmup):

Все три GPU использовали 100% своих SM: при 100M элементах и block=256 получается 390 625 блоков — с избытком для загрузки 108 SM (A100), 128 SM (RTX 4090) и 82 SM (RTX 3090). Разные проценты эффективности отражают архитектуру памяти, а не SM-утилизацию: GDDR6X (RTX 4090, 3090) хорошо насыщается простым последовательным стримингом, HBM2e (A100) оптимизирован под массивно-параллельный доступ и матричные операции — SAXPY его не раскрывает (SAXPY = Single-precision A·X Plus Y — операция из стандарта Basic Linear Algebra Subprograms, 1970-е). На transformer attention и крупных matmul картина была бы другой.

cuda-core vs pycuda

Ограничения

• CUDA 12+ обязателен

• Green Contexts — Ada Lovelace (sm_89) и новее, CUDA 12.4+

• Process checkpointing — только Linux

Итог

cuda-core v1.0.0 — официальный и теперь стабильный Python-интерфейс к CUDA. Кэш компиляции убирает latency cold start, Green Contexts дают детерминированное разделение GPU между задачами, NVML — мониторинг из той же библиотеки. Для тех, кто писал кастомные ядра через pycuda или ctypes — время мигрировать.

# CUDA 12 (драйвер 525+)pip install "cuda-core[cu12]"# CUDA 13 (драйвер 570+)pip install "cuda-core[cu13]"