Траектория «Параллельные вычисления» · линия P · конспект 5 из 6

Данные на видеокарте из Python

О чём эта тема: Явное управление памятью устройства в Numba: to_device, copy_to_host и device-массивы, жизнь данных между запусками ядер, разделяемая память с редукцией — и честный замер, в котором копирования видны отдельной строкой.
Аннотация: P04 передавал в ядро обычные numpy-массивы, и Numba молча возила их туда-обратно на каждый вызов — узаконенная неэффективность для прототипов. Этот конспект наводит порядок: cuda.to_device поднимает массив на устройство один раз, device-массивы живут там между запусками ядер, copy_to_host забирает только результат. Правило то же, что в C06: данные должны жить на видеокарте, а не кататься вокруг каждого ядра. Вторая половина — перевод главного приёма линии C на Python: блочная редукция через cuda.shared.array и cuda.syncthreads, строка в строку с C05, плюс cuda.atomic.add для межблочной суммы — оба варианта проверены в симуляторе и сходятся с numpy. Тренажёр-счётчик: соберите конвейер из нескольких ядер и посмотрите, сколько копирований по PCIe он выполняет при неявном и явном управлении памятью.
Пререквизиты: P04 — @cuda.jit и запуск ядер; C03 и C05 — этажи памяти и древовидная редукция: здесь они возвращаются в питоновском синтаксисе.
Мотивация: Самая частая жалоба новичка: «перенёс функцию на GPU, а стало медленнее». В девяти случаях из десяти виноват не код ядра, а незаметные копирования: удобная магия неявной передачи данных оборачивается двумя пересылками по PCIe на каждый запуск. Одна привычка — поднять данные на устройство явно и держать их там до конца вычислений — часто даёт больший выигрыш, чем любая оптимизация самого ядра.

1. Явное копирование: to_device и copy_to_host

Инструментов три, и все они — методы-близнецы cudaMalloc/cudaMemcpy из C03, завёрнутые в питоновскую эргономику:

from numba import cuda
import numpy as np

a = np.arange(1000, dtype=np.float32)

d_a = cuda.to_device(a)                # выделить на устройстве + скопировать (H2D)
d_out = cuda.device_array_like(a)      # только выделить, без копирования —
                                       # для результатов, которые ядро перезапишет

vec_add[blocks, threads](d_a, d_b, d_out)   # ядру передаём device-массивы:
                                            # никаких скрытых копирований

out = d_out.copy_to_host()             # единственное копирование обратно (D2H)

Device-массив — это дескриптор памяти видеокарты: у него есть shape и dtype, его можно передавать в ядра сколько угодно раз, но напечатать или отдать в matplotlib нельзя — сначала copy_to_host. Сравните два стиля на конвейере из трёх ядер над одним массивом:

# неявно: 3 запуска × (туда + обратно) = 6 копирований по PCIe
step1[g, b](a);  step2[g, b](a);  step3[g, b](a)

# явно: 1 туда + 1 обратно = 2 копирования, ядра работают по данным на месте
d = cuda.to_device(a)
step1[g, b](d);  step2[g, b](d);  step3[g, b](d)
a = d.copy_to_host()

Замер с копированиями делается так же, как учил C06, — с той разницей, что у Numba секундомером служит обычный time.perf_counter вокруг явной синхронизации:

import time

kernel[g, b](d)                # прогрев: первый вызов включает компиляцию (P01)
cuda.synchronize()

t0 = time.perf_counter()
kernel[g, b](d)                # замеряем только ядро — данные уже на устройстве
cuda.synchronize()             # без этого замерим постановку в очередь (C06!)
t1 = time.perf_counter()

Типичная ошибка Забыть cuda.synchronize() перед остановкой секундомера. Запуск ядра асинхронен и в Numba: perf_counter покажет микросекунды «постановки в очередь» вместо реального времени — та же ловушка, что у системных часов в C06, только теперь в Python. Прогрев тоже обязателен: в первый вызов входит компиляция ядра.

2. Разделяемая память: редукция по-питоновски

Быстрая память блока из C03/C05 доступна и здесь: cuda.shared.array объявляет массив на блок, cuda.syncthreads — тот самый барьер. Древовидная редукция C05 переводится дословно (проверено в симуляторе: сумма частичных сумм совпала с numpy):

from numba import cuda, float32

@cuda.jit
def block_sum(a, partial):
    sm = cuda.shared.array(256, dtype=float32)   # «общий стол» блока; размер —
    tid = cuda.threadIdx.x                        # константа времени компиляции
    i = cuda.grid(1)

    sm[tid] = a[i] if i < a.size else 0.0       # хвостовые треды кладут ноль
    cuda.syncthreads()

    stride = cuda.blockDim.x // 2
    while stride > 0:                             # дерево: 256 → 128 → … → 1
        if tid < stride:
            sm[tid] += sm[tid + stride]
        cuda.syncthreads()                        # барьер вне if — правило C05 в силе
        stride //= 2

    if tid == 0:
        partial[cuda.blockIdx.x] = sm[0]          # частичная сумма блока

Частичные суммы можно досложить на хосте (partial.sum() — их всего десятки), а можно прямо в ядре атомарным сложением — перевод atomicAdd из C05:

@cuda.jit
def total_atomic(a, out):
    i = cuda.grid(1)
    if i < a.size:
        cuda.atomic.add(out, 0, a[i])   # (массив, индекс, слагаемое)

Оба варианта проверены: 499 500 — сумма чисел 0…999 — сходится и у дерева, и у атомарной версии, и у numpy. Все предостережения линии C сохраняют силу вместе с приёмами: атомарные операции сериализуются (по одной на блок — дёшево, по одной на тред — очередь), сложение float неассоциативно (сравнивать с допуском), а syncthreads в расходящейся ветке вешает блок. Python сокращает записи — физика железа остаётся той же.

3. Тренажёр: счётчик копирований PCIe

Тренажёр · сколько копирований делает ваш код

Конвейер обрабатывает массив в 100 МБ цепочкой ядер. Выберите число ядер и стиль управления памятью — схема покажет каждую операцию, а счётчик оценит время на шине PCIe (при типичных 12 ГБ/с в одну сторону).

ядер в конвейере: 3 явное управление (to_device / copy_to_host)

операции:

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

На каждый запуск ядра Numba неявно копирует массив на устройство и обратно — два копирования по PCIe, которые в цепочке ядер множатся. Допустимо для прототипов и разовых запусков; в рабочем коде данные поднимают на устройство один раз через to_device.
to_device выделяет память устройства и копирует туда содержимое массива (H2D). device_array_like только выделяет память той же формы и типа, без копирования — правильный выбор для выходных массивов, которые ядро всё равно перезапишет: копировать мусор незачем.
Прогревочный вызов (первый включает компиляцию — P01), затем perf_counter вокруг запуска с cuda.synchronize() перед остановкой секундомера: запуск асинхронен, и без синхронизации замеряется постановка в очередь, а не работа (та же ловушка, что в C06). Данные при этом должны уже лежать на устройстве, иначе в замер попадут копирования.
Разделяемая память распределяется на этапе компиляции ядра: от её объёма на блок зависит, сколько блоков поместится на SM (C03). Поэтому размер — литерал или глобальная константа, известная компилятору, а не значение из аргумента.
Неявные копирования numpy-массивов вокруг каждого запуска (лечится to_device); замер без прогрева — в него вошла компиляция; задача слишком мала или бедна вычислениями на байт — по правилам рентабельности C06 ей место на CPU.

Источники

Memory management // Numba Documentation : [сайт]. — URL: https://numba.readthedocs.io/en/latest/cuda/memory.html (дата обращения: 09.07.2026).
GPU Reduction // Numba Documentation : [сайт]. — URL: https://numba.readthedocs.io/en/latest/cuda/reduction.html (дата обращения: 09.07.2026).
CUDA C++ Best Practices Guide. Data Transfer Between Host and Device // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/ (дата обращения: 09.07.2026).