CUDA без C++: ядра на Python

1. То же ядро, другой язык

Возьмём сложение векторов из C04 и поставим два листинга рядом. C++:

__global__ void vecAdd(const float* a, const float* b, float* c, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}
// запуск: vecAdd<<<grid, block>>>(da, db, dc, n);

Python с Numba:

from numba import cuda
import numpy as np

@cuda.jit                          # «__global__ по-питоновски»
def vec_add(a, b, c):
    i = cuda.grid(1)               # = blockIdx.x*blockDim.x + threadIdx.x
    if i < c.size:                 # та же защита от хвостовых тредов
        c[i] = a[i] + b[i]

n = 1000
a = np.arange(n, dtype=np.float32)
b = 2 * a
c = np.zeros(n, dtype=np.float32)

threads = 256
blocks = (n + threads - 1) // threads          # округление вверх — как в C04
vec_add[blocks, threads](a, b, c)              # конфигурация — в квадратных скобках

print(np.allclose(c, 3 * a))                   # True — проверено

Отличий по существу три. Конфигурация запуска пишется в квадратных скобках перед аргументами — kernel[blocks, threads] вместо тройных угловых. Формула глобального индекса спрятана в cuda.grid(1) — но никуда не делась: развёрнутые blockIdx/blockDim/threadIdx тоже доступны, и для двумерного грида px, py = cuda.grid(2) возвращает пару координат сразу. И третье: мы передали в ядро обычные numpy-массивы — Numba молча скопировала их на устройство и обратно. Удобно для старта, дорого для скорости: явное управление копированием данных — тема P05.

2. Словарь соответствий C++ ↔ Numba

Всё, что встречалось в линии C, переводится механически:

Не переводится (и это важно знать заранее): внутри ядра нельзя создавать numpy-массивы, вызывать не-device-функции, использовать списки и словари — ядро подчиняется тем же ограничениям nopython-режима, что и @njit (P01), плюс своим: код одного треда, скаляры и индексация. Печать print внутри ядра работает — как printf в C01, с той же оговоркой про синхронизацию.

3. Симулятор: CUDA без видеокарты

У Numba есть режим, которого лишена CUDA C++: симулятор. Если до импорта numba установить переменную окружения, все ядра будут выполняться обычным Python-кодом на CPU — тред за тредом, блок за блоком:

# вариант 1: в терминале
#   Windows (PowerShell):  $env:NUMBA_ENABLE_CUDASIM = "1"
#   Linux/macOS:           export NUMBA_ENABLE_CUDASIM=1

# вариант 2: в начале скрипта, ДО импорта numba
import os
os.environ["NUMBA_ENABLE_CUDASIM"] = "1"
from numba import cuda

print(cuda.is_available())      # True — «устройство» есть, хоть и нарисованное

Симулятор чудовищно медленный — он для корректности, не для скорости. Зато внутри «ядра» работают print для любого треда, точки останова отладчика и вся интроспекция Python: можно остановиться в треде (5, 3) блока (2, 1) и рассмотреть его переменные — роскошь, недоступная настоящему железу. Рабочий цикл получается такой:

Все Python-примеры этой и следующих страниц выполнены именно так — в симуляторе, на машине без NVIDIA (результаты сверены с эталонами; скрипты — в папке сode). Настоящую скорость даст Colab из C01: тот же блокнот, только вместо %%writefile и nvcc — просто ячейки с @cuda.jit; Toolkit там уже стоит, T4 к вашим услугам.

4. Тренажёр: переводчик C++ → Numba

Контрольные вопросы

• Что общего и что различно у @cuda.jit-ядра и __global__-функции?
  Общее — всё существенное: код одного треда, запуск гридом из блоков, те же threadIdx/blockIdx, невозможность вернуть значение, результат через переданные массивы. Различия синтаксические: декоратор вместо спецификатора, kernel[grid, block] вместо тройных скобок, cuda.grid(1) как готовая формула индекса.
• Что делает cuda.grid(1) и cuda.grid(2)?
  cuda.grid(1) возвращает глобальный индекс треда — blockIdx.x·blockDim.x + threadIdx.x. cuda.grid(2) — пару (x, y) для двумерного грида: обе координаты по той же формуле по своим осям. Это сокращения; развёрнутые встроенные переменные тоже доступны.
• Что происходит, когда в @cuda.jit-ядро передают обычный numpy-массив?
  Numba неявно копирует его в память устройства перед запуском и результат обратно после. Удобно для прототипа, но копирование происходит на каждый вызов ядра — для цепочек ядер и замеров нужно явное управление через cuda.to_device / copy_to_host (P05).
• Для чего симулятор CUDASIM подходит, а для чего категорически нет?
  Подходит для проверки корректности и отладки: ядра выполняются интерпретатором на CPU, работают print и отладчик, видеокарта не нужна вовсе. Не подходит для замеров: выполнение в тысячи раз медленнее и устроено принципиально иначе — тайминги симулятора не значат ничего.
• Почему ограничения @njit (P01) действуют и внутри @cuda.jit, и какие добавляются?
  Ядро тоже компилируется в машинный код (PTX) через тот же вывод типов — динамические структуры Python в него не помещаются. Добавляется специфика устройства: нельзя создавать массивы внутри ядра, вызывать можно только device-функции, стиль — скаляры и индексация, код одного треда.

Источники

1. Numba for CUDA GPUs. Overview; Writing CUDA Kernels // Numba Documentation : [сайт]. — URL: https://numba.readthedocs.io/en/latest/cuda/kernels.html (дата обращения: 09.07.2026).
2. Debugging CUDA Python with the the CUDA Simulator // Numba Documentation : [сайт]. — URL: https://numba.readthedocs.io/en/latest/cuda/simulator.html (дата обращения: 09.07.2026).
3. Supported Python features in CUDA Python // Numba Documentation : [сайт]. — URL: https://numba.readthedocs.io/en/latest/cuda/cudapysupported.html (дата обращения: 09.07.2026).