IВведение: от оптимизации ядра к оркестрации конвейера

Глава 12 книги Fregly завершает блок, посвящённый низкоуровневой оптимизации CUDA. Если предыдущие главы были сосредоточены на внутриядерных техниках — арифметической интенсивности, Tensor Cores, специализации варпов и персистентных ядрах, — то глава 12 поднимается на уровень выше: как организовать параллельную работу между ядрами, между GPU и хостом, и между самими GPU. Ключевая идея главы: производительность GPU определяется не только плотностью загрузки SM, но и одновременной работой всех аппаратных подсистем — вычислительных SM, DMA-движков и интерконнектов.

IIДинамическое планирование с помощью атомарных очередей задач

Анатомия простоя SM

Неравномерное распределение работы между потоками может привести к тому, что одни SM будут простаивать, пока другие всё ещё заняты. Это растрачивает вычислительные ресурсы и снижает общую пропускную способность.

Дисбаланс часто возникает, когда разные потоки или блоки обрабатывают переменный объём работы из-за зависящих от данных циклов или условных нагрузок. Некоторые блоки завершаются быстро, оставляя свои SM простаивающими, в то время как другие SM продолжают выполнять более длительные блоки.

На современных GPU с сотнями SM периоды простоя могут оставить множество SM без работы. Рассмотрим конкретный пример: GPU NVIDIA H100 содержит 132 SM, каждый из которых может одновременно выполнять до 64 варпов (2048 потоков). Если нагрузка такова, что половина блоков завершается за 1 мс, а вторая половина — за 10 мс, то в течение 9 мс половина SM (66 из 132) полностью простаивает. Это эквивалентно потере ~33 трлн потенциальных FLOPS в секунду — сравнимо с производительностью отдельного GPU предыдущего поколения.

К моменту завершения самой длительной работы часть GPU уже простаивала. Это снижает достигнутую заполняемость (achieved occupancy) — среднюю долю активных варпов на SM относительно аппаратного максимума. Помните, что вы можете использовать Nsight Systems для профилирования и наглядного отображения этих пробелов на временной шкале GPU.

Чтобы сопоставить пробелы на временной шкале с конкретными участками кода, вставьте маркеры диапазона NVTX вокруг значимых участков работы GPU:

#include <nvtx3/nvToolsExt.h>

nvtxRangePush("computeKernel_phase1");
computeKernel<<<grid, block>>>(d_in, d_out, N);
cudaStreamSynchronize(stream);
nvtxRangePop();

Атомарные счётчики: аппаратные основы

Атомарные счётчики — это основа для атомарных очередей, которые позволяют динамически распределять работу.

На современных GPU (архитектуры Ampere, Hopper, Blackwell) глобальные атомарные операции обслуживаются и сериализуются в кэше L2 на устройстве. Это снижает задержку по сравнению с обращениями к DRAM, когда целевая строка находится в кэше. Типичная задержка atomicAdd в L2 составляет 20–40 нс, в то время как атомарная операция с промахом в DRAM — 200–400 нс. Атомарные счётчики всё ещё несут задержку и сериализуются при конкуренции. Но неконкурирующие операции atomicAdd выполняются чрезвычайно быстро, оставаясь на кристалле.

В разделе Memory Workload Analysis в Nsight Compute вы найдёте счётчик atomic_transactions и atomic_transactions_per_request. Счётчик атомарных транзакций представляет общее количество атомарных транзакций в кэше L2, включая любые повторы, вызванные конкуренцией. Метрика атомарных транзакций на запрос, или коэффициент конкуренции, представляет среднее количество L2-транзакций, генерируемых каждой инструкцией atomicAdd.

Когда каждый atomicAdd выполняет ровно одну L2-транзакцию, ваш atomic_transactions_per_request держится около 1,0, что означает минимальные затраты. Если это отношение поднимается выше 1,0, это сигнализирует о том, что потоки простаивают и повторяют атомарные обновления вместо полезной работы. Каждый повтор указывает на конкуренцию.

Дополнительная оптимизация: __reduce_add_sync для суммирования внутри варпа

Варп-редукция __reduce_add_sync (sm_60+ для целочисленных типов, CUDA 9; redux.sync — sm_80+) выполняет редукцию суммы по варпу за одну инструкцию без необходимости в atomicAdd:

unsigned mask = 0xffffffff;
int partial = my_value;
int warp_sum = __reduce_add_sync(mask, partial);
int lane = threadIdx.x & 31;
if (lane == 0) {
    atomicAdd(&global_counter, warp_sum);
}

Этот подход снижает количество глобальных атомарных операций в 32 раза по сравнению с попотоковой схемой — без необходимости в __shfl_sync и ручном выборе лидера.

Пакетирование: амортизация атомарных операций

Оптимизация здесь заключается в амортизации атомарных операций путём захвата работы пакетами. То есть вместо того, чтобы каждый поток или варп выполнял atomicAdd, вы группируете несколько задач на одну атомарную операцию. Вот пример до и после с размером пакета 32:

// До пакетирования
int idx = atomicAdd(&queue_head, 1);
if (idx < N) process(data[idx]);

// После пакетирования
const int batchSize = 32;
int start = atomicAdd(&queue_head, batchSize);
for (int i = start; i < start + batchSize && i < N; ++i) {
    process(data[i]);
}

Теперь одно атомарное обновление выдаёт варпу целый срез работы — 32 элемента — прежде чем снова обратиться к счётчику. Вы по-прежнему платите одну L2-транзакцию за пакет, но выполняете в 32 раза больше полезной работы между ними.

На практике только один поток на варп выполняет этот atomicAdd, чтобы получить начальный индекс следующего пакета. Затем он транслирует его остальным потокам варпа (например, с помощью __shfl_sync). Весь варп обрабатывает эти 32 элемента параллельно. Это даёт одну атомарную операцию на варп вместо одной на поток, резко снижая конкуренцию.

В Nsight Compute вы увидите, как atomic_transactions резко падает, а количество транзакций на запрос возвращается к 1,0. Это доказывает, что вы обменяли дорогостоящую конкуренцию на устойчивые вычисления.

Динамические очереди на уровне варпов

Давайте теперь используем глобальный атомарный счётчик для координации динамической очереди задач. Цель — использовать атомарный счётчик и atomicAdd для балансировки произвольных нагрузок между всеми SM так, чтобы ни один поток или варп не простаивал.

В следующем примере кода (computeKernel) каждый поток вычисляет разное количество итераций на основе idx % 256. Потоки с малым значением idx % 256 выполняют очень мало работы, в то время как потоки с большим значением — много работы. В результате такого дисбаланса потоки завершаются в разное время, и некоторые SM простаивают в ожидании завершения самых длинных потоков. Вот код, использующий статическую неравномерную нагрузку:

// uneven_static.cu
#include <cuda_runtime.h>
#include <cmath>

__global__ void computeKernel(const float* input, float* output, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        int work = idx % 256;
        float result = 0.0f;
        for (int i = 0; i < work; ++i) {
            result += sinf(input[idx]) * cosf(input[idx]);
        }
        output[idx] = result;
    }
}

int main() {
    const int N = 1 << 20;
    float* h_in = nullptr;
    float* h_out = nullptr;
    cudaMallocHost(&h_in, N * sizeof(float));
    cudaMallocHost(&h_out, N * sizeof(float));

    for (int i = 0; i < N; ++i) h_in[i] = float(i) / N;

    float *d_in, *d_out;
    cudaMalloc(&d_in, N * sizeof(float));
    cudaMalloc(&d_out, N * sizeof(float));
    cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);

    dim3 block(256), grid((N + 255) / 256);
    computeKernel<<<grid, block>>>(d_in, d_out, N);
    cudaDeviceSynchronize();

    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_in);
    cudaFree(d_out);
    cudaFreeHost(h_in);
    cudaFreeHost(h_out);
    return 0;
}

В оптимизированной версии с динамической диспетчеризацией задач единый глобальный счётчик (в памяти устройства) используется как очередь задач на уровне варпов. Мы превращаем счётчик в постоянную очередь задач, используя пакетные атомарные операции. Варпы, которые завершаются раньше, немедленно захватывают следующий пакет вместо простоя:

// uneven_dynamic.cu
#include <cuda_runtime.h>

__device__ unsigned int globalIndex = 0;

__global__ void computeKernelDynamicBatch(const float* input,
                                          float* output,
                                          int N) {
    int lane = threadIdx.x & (warpSize - 1);

    while (true) {
        unsigned mask = __activemask();
        int leader = __ffs(mask) - 1;

        unsigned int base = 0;
        if (lane == leader) {
            base = atomicAdd(&globalIndex, warpSize);
        }

        base = __shfl_sync(mask, base, leader);
        unsigned int idx = base + lane;

        if (idx >= (unsigned)N) break;

        float s = sinf(input[idx]);
        float c = cosf(input[idx]);
        int work = idx % 256;
        float result = 0.0f;

        #pragma unroll 1
        for (int i = 0; i < work; ++i) {
            result += s * c;
        }

        output[idx] = result;
    }
}

Иерархические и блочно-локальные очереди

Одиночная глобальная очередь хорошо работает для умеренных дисбалансов, но при экстремальной конкуренции (сотни активных варпов одновременно обращаются к одному счётчику) может потребоваться иерархический подход.

Идея: каждый блок потоков сначала обрабатывает локальный пул задач, выделенный ему при запуске. Когда локальный пул исчерпан, блок атомарно запрашивает следующий чанк из глобальной очереди. Это двухуровневая схема:

__global__ void hierarchicalDynamicKernel(const float* input,
                                           float* output,
                                           int N,
                                           int chunkPerBlock) {
    extern __shared__ int blockOffset[];

    int tid = threadIdx.x;
    int lane = tid & 31;

    if (tid == 0) {
        blockOffset[0] = atomicAdd(&globalIndex, chunkPerBlock);
    }
    __syncthreads();

    int localStart = blockOffset[0];
    int processed = 0;

    while (true) {
        int localIdx = localStart + processed + tid;
        if (localIdx >= (unsigned)N) {
            if (tid == 0) {
                blockOffset[0] = atomicAdd(&globalIndex, chunkPerBlock);
            }
            __syncthreads();
            localStart = blockOffset[0];
            processed = 0;
            localIdx = localStart + tid;
            if (localIdx >= (unsigned)N) break;
        }
        if (localIdx < (unsigned)N) {
            float s = sinf(input[localIdx]);
            float c = cosf(input[localIdx]);
            int work = localIdx % 256;
            float result = 0.0f;
            for (int i = 0; i < work; ++i) {
                result += s * c;
            }
            output[localIdx] = result;
        }
        processed += blockDim.x;
    }
}

Здесь chunkPerBlock — размер чанка, выделяемого каждому блоку за один atomicAdd. При chunkPerBlock = 256 (по одному элементу на поток в блоке из 256) и 4096 блоков на сетку, глобальный atomicAdd выполняется 4096 раз вместо 1 млн раз — снижение конкуренции в 244 раза.

Оценка производительности и метрики профилирования

В простом микробенчмарке с N = 2²⁰ и work = idx % 256 статически назначенное ядро заняло около 200 мс, тогда как версия с динамической очередью выполнилась примерно за 100 мс — ускорение в 2 раза. Для воспроизводимости: тест проводился на NVIDIA H100 SXM (80 ГБ HBM3), CUDA 12.4, частота ядра 1830 МГц, nvcc -O3 -arch=sm_90.

В Nsight Compute проверьте: Warp Stall Reasons — если преобладает Stall Long Scoreboard (stall_long_sb, ожидание памяти), пакетирование может усугубить задержку; Register Pressure — дополнительный цикл while(true) увеличивает использование регистров; Achieved Occupancy — если динамическая версия показывает заполняемость ниже 60%, возможно, размер пакета слишком мал.

IIICUDA Graphs

Мотивация: накладные расходы запуска

Когда ваш конвейер состоит из множества ядер, копирований, записей событий потока и обратных вызовов, их последовательный запуск на хосте на каждой итерации влечёт накладные расходы CPU. Типичная задержка запуска одного ядра с хоста составляет 5–10 мкс на современном GPU. Если ваш конвейер содержит 50 ядер (типично для inference-движков), это 250–500 мкс чистых накладных расходов на итерацию — при том, что вычисления могут занимать всего 1–2 мс. Накладные расходы составляют 12–25% общего времени.

CUDA Graphs позволяют захватить весь рабочий процесс один раз и воспроизводить его повторно практически без накладных расходов CPU — типичная задержка воспроизведения графа из 50 узлов составляет 2–5 мкс, что в разы (до 5–10× при коротких ядрах, в зависимости от нагрузки) меньше, чем при последовательных индивидуальных запусках.

Зачем использовать CUDA Graphs: сокращение накладных расходов на запуск; лучшая планировка на GPU; зависимости известны заранее; корректное перекрытие асинхронных копирований и ядер.

Захват графа: потоковый и явный API

CUDA Graphs поддерживают два способа создания графа: потоковый захват (stream capture) и явное построение (explicit graph API).

Потоковый захват — самый простой способ:

cudaStream_t stream;
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
cudaGraph_t graph;
cudaGraphExec_t instance;

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

kernelA<<<grid, block, 0, stream>>>(d_X);
kernelB<<<grid, block, 0, stream>>>(d_Y);
kernelC<<<grid, block, 0, stream>>>(d_Z);

cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&instance, graph, nullptr, nullptr, 0);

for (int iter = 0; iter < 100; ++iter) {
    cudaGraphLaunch(instance, stream);
}

cudaStreamSynchronize(stream);
cudaGraphExecDestroy(instance);
cudaGraphDestroy(graph);

Явное построение даёт более тонкий контроль — вы строите граф узел за узлом, явно указывая зависимости:

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);

cudaGraphNode_t nodeA, nodeB, nodeC;
cudaKernelNodeParams paramsA = {}, paramsB = {}, paramsC = {};

paramsA.func = (void*)kernelA;
paramsA.gridDim = grid;
paramsA.blockDim = block;

cudaGraphAddKernelNode(&nodeA, graph, nullptr, 0, &paramsA);

cudaGraphNode_t depsB[] = {nodeA};
cudaGraphAddKernelNode(&nodeB, graph, depsB, 1, &paramsB);

cudaGraphNode_t depsC[] = {nodeB};
cudaGraphAddKernelNode(&nodeC, graph, depsC, 1, &paramsC);

Пулы памяти и стабильность указателей

Важным соображением является управление памятью с CUDA Graphs. Операции с памятью внутри CUDA Graph подчиняются тем же правилам, что и в потоках CUDA. Если вы выделяете память GPU внутри захвата, это выделение становится частью выполнения графа.

Обычно вы хотите избегать выделения памяти GPU внутри графа и предварительно выделять память вне графа. Многие фреймворки, такие как PyTorch, используют статические пулы памяти с CUDA Graphs:

pool = torch.cuda.graph_pool_handle()
g1 = torch.cuda.CUDAGraph()
g2 = torch.cuda.CUDAGraph()

with torch.cuda.graph(g1, pool=pool):
    # операции для графа 1
    ...

with torch.cuda.graph(g2, pool=pool):
    # операции для графа 2
    ...

Динамическое обновление графа

После того как вы записали CUDA Graph, вам не нужно выбрасывать его только из-за изменения параметров запуска. API обновления графа включает cudaGraphExecUpdate и низкоуровневый cudaGraphExecKernelNodeSetParams.

cudaKernelNodeParams newParams = {};
newParams.func = (void*)kernelB;
newParams.gridDim = dim3((batchSize + 255) / 256);
newParams.blockDim = dim3(256);
newParams.kernelParams = kernelArgs;

cudaGraphExecKernelNodeSetParams(instance, nodeB, &newParams);
cudaGraphLaunch(instance, stream);

PyTorch, инференс-движки и CUDA Graphs

AI-фреймворки, такие как PyTorch, используют CUDA Graphs под капотом для статических участков моделей глубокого обучения:

import torch

X = torch.randn(1<<20, device='cuda')

def opA(x): return x * 1.1
def opB(x): return x + 2.0
def opC(x): return x.sqrt()

static_x = torch.empty_like(X)
static_y = torch.empty_like(X)
static_z = torch.empty_like(X)
static_w = torch.empty_like(X)

_ = opC(opB(opA(X)))
torch.cuda.synchronize()

static_x.copy_(X)

g = torch.cuda.CUDAGraph()
stream = torch.cuda.Stream()
torch.cuda.synchronize()

with torch.cuda.graph(g, stream=stream):
    torch.mul(static_x, 1.1, out=static_y)
    torch.add(static_y, 2.0, out=static_z)
    torch.sqrt(static_z, out=static_w)

for i in range(100):
    g.replay()

CUDA Graphs в тензорных движках: Triton и CuDNN

Triton генерирует CUDA-ядра, которые могут быть захвачены в графы точно так же, как ядра на C++:

import triton
import triton.language as tl

@triton.jit
def triton_kernel(x_ptr, y_ptr, N):
    pid = tl.program_id(0)
    offsets = pid * 256 + tl.arange(0, 256)
    mask = offsets < N
    x = tl.load(x_ptr + offsets, mask=mask)
    tl.store(y_ptr + offsets, x * 1.1, mask=mask)

g = torch.cuda.CUDAGraph()
with torch.cuda.graph(g):
    triton_kernel[(grid,)](x_buf, y_buf, N)
g.replay()

IVЗапуск CUDA Graph с инициированием с устройства

Мотивация: устранение хоста с критического пути

С помощью запуска CUDA Graph с инициированием с устройства работающее ядро GPU может запустить предварительно записанный граф непосредственно на устройстве, полностью минуя хост.

Чтобы включить запуск с устройства, сначала захватите граф как обычно на хосте. Затем создайте его экземпляр с флагом cudaGraphInstantiateFlagDeviceLaunch и загрузите с помощью cudaGraphUpload:

cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr,
                     cudaGraphInstantiateFlagDeviceLaunch);
cudaGraphUpload(graphExec, stream);

Запуски графов с устройства убирают планирование хоста с критического пути. На практике они показывают примерно вдвое меньшую задержку запуска по сравнению с эквивалентными запусками со стороны хоста. Задержка запуска с устройства не зависит от количества узлов или параллельных ветвей в графе.

Режимы запуска: fire-and-forget, tail, sibling

Fire-and-forget (cudaStreamGraphFireAndForget): дочерний граф начинает выполняться немедленно и одновременно с запускающим родительским ядром. Родитель не ждёт завершения дочернего. Лимит: до 120 общих fire-and-forget графов одновременно.

Tail (cudaStreamGraphTailLaunch): выполнение графа откладывается до завершения запускающего ядра. Граф ставится в очередь как продолжение. Лимит: до 255 ожидающих tail-запусков; для self-tail-launch — только один одновременно.

Tail-запуски особенно мощны при реализации планировщиков, резидентных на GPU. Постоянное ядро-«планировщик» запускает граф через tail, затем перезапускает себя:

__global__ void schedulerKernel(/* ... */) {
    cudaGraphExec_t self = cudaGetCurrentGraphExec();

    unsigned int taskIdx = atomicAdd(&queueHead, 1);
    if (taskIdx >= N_tasks) return;

    cudaGraphLaunch(decodingGraphExec, cudaStreamGraphTailLaunch);
    cudaGraphLaunch(self, cudaStreamGraphTailLaunch);
}

Sibling (cudaStreamGraphFireAndForgetAsSibling): запущенный граф выполняется как равноправный с родительским графом, а не как дочерний.

Резидентный планировщик LLM-инференса

Рассмотрим практический паттерн, объединяющий атомарные очереди и CUDA Graphs с device-side tail-запуском для планирования инференса LLM внутри постоянного ядра.

Лёгкое постоянное ядро-планировщик использует atomicAdd(&queueHead, 1) для захвата следующего элемента работы. Затем оно запускает через tail предварительно захваченный CUDA Graph декодирования и немедленно возвращается к следующему элементу в очереди.

Связывая tail-запуски в цепочку, каждый токен обрабатывается от начала до конца на устройстве. Поскольку CPU никогда не возвращается на критический путь, SM остаются загруженными, задержка на токен снижается, и можно адаптироваться к разным длинам последовательностей и размерам батчей на лету.

Отладка device-side графов

Отладка графов, запущенных с устройства, может быть сложной. Рекомендации:

• Nsight Systems показывает дочерние графы на временной шкале GPU как отдельные потоки работы.
• NVTX-маркеры в родительском ядре до и после вызовов cudaGraphLaunch помогают отметить, где происходят запуски.
• Проверка ошибок device-side: используйте cudaGetLastError() внутри устройства (доступно с CUDA 12.0).
• Логирование: printf внутри устройства работает, но может существенно замедлить выполнение.

VУсловные узлы графа

Типы условных узлов: IF, IF/ELSE, WHILE, SWITCH

В традиционном CUDA Graph каждый узел и его зависимости фиксированы на момент захвата, что вынуждает всю логику принятия решений возвращать на хост. Условные узлы графа ломают эту жёсткость, откладывая принятие решений на сам GPU — на основе небольшого «дескриптора условия», связанного с узлом.

Пример создания условного узла IF:

__global__ void setHandle(cudaGraphConditionalHandle handle,
                          int *data, size_t N) {
    constexpr int threshold = 123456;
    unsigned int flag =
        (reduce_sum(data, N) > threshold) ? 1u : 0u;
    cudaGraphSetConditional(handle, flag);
}

__global__ void bodyKernel() {
    printf("Условное тело выполнено на GPU!\n");
}

Вложенные условия и многоуровневые решения

Условные узлы могут быть вложенными. Например, тело узла WHILE может содержать узел IF. Это позволяет реализовать многоуровневую логику принятия решений без обращений к CPU:

┌──────────────────────────────────────┐
│  WHILE (пока есть невыполненные задачи)│
│  ┌──────────────────────────────────┐ │
│  │  IF (тип задачи == "compression")│ │
│  │  ┌────────────┐ ┌──────────────┐ │ │
│  │  │ LZ-сжатие  │ │ Huff-сжатие  │ │ │
│  │  └────────────┘ └──────────────┘ │ │
│  └──────────────────────────────────┘ │
│  ┌──────────────────────────────────┐ │
│  │  ELSE IF (тип == "decompression") │ │
│  │  ┌────────────────────────────┐  │ │
│  │  │  Декомпрессия              │  │ │
│  │  └────────────────────────────┘  │ │
│  └──────────────────────────────────┘ │
└──────────────────────────────────────┘

Практический пример: адаптивный выбор ядра

Рассмотрим конвейер сжатия на GPU, где ядро должно выбирать между разными алгоритмами в зависимости от содержимого данных:

__global__ void classifyData(cudaGraphConditionalHandle handle,
                               int* data, size_t N) {
    if (threadIdx.x == 0) {
        float entropy = compute_entropy(data, N);

        unsigned int algo;
        if (entropy < 2.0f)      algo = 0; // LZ-сжатие
        else if (entropy < 5.0f) algo = 1; // Huffman
        else                    algo = 2; // Без сжатия

        cudaGraphSetConditional(handle, algo);
    }
}

VIДинамический параллелизм

Когда нужен DP, а когда достаточно графа

Динамический параллелизм (DP) даёт вашим ядрам GPU способность порождать новую работу для себя, вместо того чтобы ждать CPU. В то время как CUDA Graphs требуют, чтобы вы знали весь конвейер заранее, DP позволяет работающему «родительскому» ядру исследовать свои выходные данные и решать на лету, сколько «дочерних» ядер запустить.

Управление ресурсами: стек и pending launch count

Каждый дочерний запуск использует ресурсы планирования GPU и требует дополнительного пространства стека. Чтобы избежать ошибок «переполнения стека», увеличьте размер стека:

cudaDeviceSetLimit(cudaLimitStackSize, 1024);
cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 8192);

Сравнение host-launched и device-launched

Версия с DP и device-side запуском:

// dp_device_launched.cu
// Компилировать с -rdc=true
__global__ void childKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) data[idx] = data[idx] * data[idx];
}

__global__ void parentKernel(float* data, int n) {
    if (blockIdx.x == 0 && threadIdx.x == 0) {
        const int threadsPerBlock = 256;
        const int firstHalfCount = n / 2;
        const int secondHalfCount = n - firstHalfCount;
        const int blocksFirst = (firstHalfCount + threadsPerBlock - 1)
            / threadsPerBlock;
        const int blocksSecond = (secondHalfCount + threadsPerBlock - 1)
            / threadsPerBlock;

        childKernel<<<blocksFirst, threadsPerBlock>>>(data, firstHalfCount);
        childKernel<<<blocksSecond, threadsPerBlock>>>(
            data + firstHalfCount, secondHalfCount);
    }
}

DP + NVSHMEM: нерегулярные задачи в кластере

Динамический параллелизм особенно мощен в сочетании с NVSHMEM для нерегулярных задач в кластере. Например, в алгоритме обхода графа каждый GPU обрабатывает свою часть графа. Когда обнаруживаются рёбра, ведущие к вершинам на других GPU, ядро может:

1. Через NVSHMEM поместить данные в удалённый GPU.
2. Через DP запустить дочернее ядро для обработки локальной части.
3. Дочернее ядро на удалённом GPU активируется по сигналу NVSHMEM.

VIIОркестрация на нескольких GPU и узлах кластера

Peer-to-peer копирования и GPUDirect

Когда вы масштабируетесь с одного GPU до многих, основная цель остаётся прежней: поддерживать занятость каждого устройства, скрывая перемещение данных за полезной работой.

cudaMemcpyPeerAsync(dest_ptr, dest_gpu, src_ptr, src_gpu, size, comm_stream);

NCCL: коллективы и интеграция с CUDA Graphs

Когда вам нужна коллективная коммуникация, такая как all-reduce градиентов в PyTorch Distributed Data Parallel (DDP), вы запускаете неблокирующие процедуры NCCL в отдельном потоке:

ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum,
              comm, comm_stream);

NCCL-операции могут быть захвачены в CUDA Graphs:

cudaStreamBeginCapture(comm_stream, cudaStreamCaptureModeGlobal);
ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum,
              comm, comm_stream);
cudaStreamEndCapture(comm_stream, &graph);

NVSHMEM: односторонняя коммуникация PGAS

Для рабочих нагрузок, требующих сверхплотной, управляемой событиями координации, NVSHMEM рассматривает каждый GPU как обрабатывающий элемент (PE) в разделяемом глобальном адресном пространстве (PGAS). GPU может напрямую писать в память другого GPU из кода устройства, минуя CPU.

Паттерн «отправить и сигнализировать»:

#include <nvshmem.h>
#include <nvshmemx.h>

__device__ int *remote_flag;
__device__ float *remote_data;

__global__ void sender_kernel(float *local_data, int dest_pe) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float value = local_data[idx];

    nvshmem_float_p(remote_data + 1, value, dest_pe);
    nvshmem_quiet();
    nvshmem_int_p(remote_flag + 0, 1, dest_pe);
}

__global__ void receiver_kernel(float *recv_buffer) {
    nvshmem_int_wait_until(remote_flag + 0, NVSHMEM_CMP_EQ, 1);
    float val = remote_data[1];
    recv_buffer[0] = val * 2.0f;
}

Work stealing на множестве GPU

NVSHMEM-атомарные операции позволяют каждому PE динамически захватывать работу из глобального счётчика — реализуя work stealing полностью на устройстве:

__global__ void work_steal_kernel(/*...*/, int *queue_head, Task *tasks) {
    while (true) {
        int idx = nvshmem_int_atomic_inc(queue_head);
        if (idx >= N_tasks) break;
        // Обрабатываем tasks[idx]...
    }
}

Кооперативные запуски и барьеры

Для сценариев, требующих минимального джиттера, запустите кооперативное ядро, охватывающее все GPU, с использованием nvshmemx_collective_launch():

__global__ void synchronized_step_kernel(/*...*/) {
    nvshmem_barrier_all();
    // Все GPU продолжают синхронно
}

Паттерн масштабирования на N GPU

Независимо от того, используете ли вы peer-копирования, кольца NCCL или односторонние операции NVSHMEM, паттерн масштабирования всегда один и тот же: система должна однократно диспетчеризовать ядра, конвейеризовать и перекрывать передачи данных с вычислениями и масштабироваться линейно — особенно с хостом CPU вне критического пути.

Формула эффективности масштабирования:

Speedup(N) = 1 / (f_compute/N + f_comm)

где f_compute — доля вычислений, f_comm — доля коммуникации, N — количество GPU. При f_comm = 0,2 и N = 8 ускорение составит 1/(0,1 + 0,2) = 3,3×, а не 8×.

VIIIПринятие решений на основе Roofline-модели

Основы Roofline

Roofline-модель — это визуальная модель производительности, которая связывает достигнутую производительность (FLOPS) с операционной интенсивностью (FLOPS/байт) и аппаратными пределами.

Performance = min(Peak FLOPS, OI × Peak Memory BW)

Операционная интенсивность (OI) — это количество операций с плавающей точкой, приходящихся на каждый байт, прочитанный из памяти:

OI = FLOPs / Bytes accessed

Например, для матричного умножения M × N × K: FLOPs = 2MNK, Bytes (FP16) = 2(MK + KN + MN). При M = N = K = 4096: OI ≈ 1365 FLOPS/байт.

Балансировочная точка для H100 (FP16): 989 TFLOPS / 3,35 TB/s ≈ 295 FLOPS/байт. Матричное умножение (OI ~1365) находится далеко справа — ограничение вычислениями. Поэлементное сложение (OI ~0,083) — далеко слева — ограничение памятью.

Сценарий A: ограничение памятью

Когда ваше ядро ограничено памятью (OI < ridge point), производительность ограничена пропускной способностью памяти. Приоритеты:

1. Слияние ядер — объединение нескольких поэлементных операций в одно ядро увеличивает OI.
2. Перекрытие копирований с вычислениями — используйте потоки CUDA для сокрытия задержки передачи данных.
3. CUDA Graphs — пакетирование операций снижает накладные расходы CPU.
4. Размещение данных — cudaMemcpyAsync с pinned memory.

Сценарий B: ограничение вычислениями

Когда ваше ядро ограничено вычислениями (OI > ridge point), производительность ограничена пиковой вычислительной мощностью:

1. Снижение накладных расходов запуска — CUDA Graphs с device-side запуском.
2. Динамический параллелизм — для адаптивных алгоритмов.
3. Динамическое планирование — балансировка нагрузки через атомарные очереди.
4. Tensor Core утилизация — wmma или mma для FP16/BF16/FP8.

Сценарий C: ограничение задержкой запуска

Когда время выполнения ядра сопоставимо с накладными расходами запуска (< 50 мкс):

1. CUDA Graphs — захват всего конвейера в один граф (2–5 мкс вместо 250+ мкс).
2. Device-side запуск графов — ещё ниже задержка (~2 мкс).
3. Условные узлы — перенос ветвления на GPU.
4. Постоянные ядра — ядро, которое не завершается, а ждёт новую работу.

Сценарий D: ограничение коммуникацией

В многопроцессорных системах производительность может быть ограничена пропускной способностью межсоединений:

1. Перекрытие коммуникации и вычислений — отдельные потоки для NCCL и вычислений.
2. NVSHMEM — односторонние операции для тонкой синхронизации.
3. Конвейеризация — разбиение батча на микробатчи.
4. NCCL в CUDA Graphs — устранение накладных расходов CPU.

Дерево решений

                    ┌─ Узкое место ─┐
                    │               │
            ┌── Память ──┐    ┌── Вычисления ──┐
            │            │    │                  │
     Слияние ядер   Перекрытие   CUDA Graphs    Динамическое
     + Graphs       копирований  + DP           планирование
                    │            │
              ┌── Запуск ──┐    ┌── Коммуникация ──┐
              │            │    │                    │
         Device-side    Условные  NVSHMEM +      NCCL в Graphs
         запуск графов   узлы     перекрытие     + конвейеризация

IXДополнительные темы

Cooperative Groups и кооперативные ядра

Cooperative Groups — это расширение CUDA, которое предоставляет более гибкие примитивы синхронизации, чем барьеры __syncthreads(). Кооперативные ядра запускают ядро, которое может синхронизироваться на уровне всей сетки через cg::this_grid().sync():

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void cooperativeKernel(float* data, int N, int iterations) {
    cg::grid_group grid = cg::this_grid();

    for (int iter = 0; iter < iterations; ++iter) {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        if (idx < N) data[idx] = process(data[idx]);
        grid.sync();
    }
}

void* args[] = {&d_data, &N, &iterations};
cudaLaunchCooperativeKernel((void*)cooperativeKernel,
                             dim3(grid), dim3(block), args, 0, stream);

Ограничение: кооперативные ядра требуют, чтобы все блоки могли быть запущены одновременно. Это ограничивает максимальное количество блоков произведением max_blocks_per_SM × SM_count. На H100 это 32 × 132 = 4224 блоков максимум.

Triton и высокоуровневые альтернативы CUDA C++

Triton — это язык программирования GPU-ядер, разработанный OpenAI, который компилируется в эффективный PTX/SASS. Он предоставляет высокоуровневые абстракции для работы с блоками данных, автоматического управления памятью и векторизации.

Пример fused-ядра на Triton для scaled dot-product attention:

import triton
import triton.language as tl

@triton.jit
def flash_attention(Q, K, V, O, sm_scale,
                    N_ctx, BLOCK_M: tl.constexpr,
                    BLOCK_N: tl.constexpr, BLOCK_D: tl.constexpr):
    start_m = tl.program_id(0)
    off_hz = tl.program_id(1)
    off_qk = off_hz * N_ctx

    q = tl.load(Q + off_qk + start_m * BLOCK_M + ...)
    tl.store(O + off_qk + start_m * BLOCK_M + ..., result)

Stream-Ordered Memory Allocator

CUDA 11.2 представил stream-ordered memory allocator (cudaMallocAsync / cudaFreeAsync), который упрощает управление памятью в динамических рабочих нагрузках и интегрируется с CUDA Graphs:

cudaStream_t stream;
cudaMallocAsync(&ptr, size, stream);
kernel<<<grid, block, 0, stream>>>(ptr);
cudaFreeAsync(ptr, stream);

Профилирование device-side запусков

Nsight Systems: дочерние графы отображаются как отдельные строки на временной шкале GPU. Используйте NVTX-маркеры. Включите --gpu-metrics-device=all.

Nsight Compute: дочерние ядра, запущенные через DP, профилируются индивидуально. Используйте --kernel-id.

Практический чек-лист для device-side запусков:

1. Убедитесь, что граф загружен через cudaGraphUpload.
2. Проверьте, что cudaGraphInstantiate вызван с флагом cudaGraphInstantiateFlagDeviceLaunch.
3. Убедитесь, что количество одновременных запусков не превышает лимиты.
4. Проверьте использование стека устройства.
5. Убедитесь, что хост не синхронизируется внутри цикла device-side запусков.

XКлючевые выводы

• Динамическое планирование с атомарными очередями в кэше L2. Атомарные операции в кэше L2 на современных GPU исключительно быстры. Используйте пакетные инкременты для балансировки нерегулярных нагрузок. Это может повысить пропускную способность до 2× в случаях экстремального дисбаланса, но обычно между 10% и 30%.
• CUDA Graphs для пакетирования конвейеров. Захват повторяющихся последовательностей ядер, копирований и событий в один граф. Это снижает накладные расходы запуска в разы (до 5–10× при коротких ядрах) и устраняет пробелы планирования CPU.
• Оркестрация на стороне устройства. Запускайте работу с самого GPU через tail-запуск CUDA Graph или динамический параллелизм. Условные узлы (IF, WHILE, SWITCH) переносят логику ветвления на GPU.
• Многопроцессорное перекрытие. Всегда перекрывайте коммуникацию с вычислениями. Используйте отдельные потоки для peer-to-peer передач, коллективов NCCL, CUDA-aware MPI или односторонних операций NVSHMEM.
• NVSHMEM для тонкой координации. Для нерегулярных, зависящих от данных рабочих нагрузок NVSHMEM предоставляет односторонние операции put/get и тонкую синхронизацию через сигналы.
• Roofline как метод выбора. Определите, ограничена ли ваша нагрузка памятью, вычислениями, задержкой запуска или коммуникацией. Выбирайте инструмент соответственно.

XIЗаключение

В этой главе мы вышли за пределы оптимизаций отдельных ядер и исследовали сквозные техники оркестрации. Мы рассмотрели, как запускать работу полностью на устройстве с помощью динамического параллелизма, захватывать сложные рабочие процессы в CUDA Graphs, переносить логику ветвления через условные узлы и координировать множество GPU с помощью NCCL и NVSHMEM.

Каждая техника преследует одну и ту же цель: поддерживать подачу работы в каждый движок, скрывать задержку и устранять разрывы хост-устройство, чтобы ваше оборудование работало на полную мощность.

Современные платформы NVIDIA всё больше размывают границу между CPU и GPU. Суперчипы Grace Blackwell и Vera Rubin соединяют CPU с несколькими GPU с помощью когерентного NVLink с огромной пропускной способностью. В таких системах традиционное разделение «хост управляет, устройство вычисляет» уступает место парадигме, в которой GPU сами управляют своим конвейером — от планирования и диспетчеризации до синхронизации и коммуникации.

Это не просто технологическая эволюция — это смена парадигмы программирования. Инженеры, которые освоят методы device-side оркестрации, получат возможность извлекать производительность, недоступную тем, кто остаётся в рамках хост-управляемой модели. Инструменты, которые мы рассмотрели — атомарные очереди, CUDA Graphs с device-side запуском, условные узлы, динамический параллелизм и NVSHMEM — это строительные блоки для конвейеров следующего поколения, в которых GPU работает как самостоятельный сопроцессор, а не как подчинённый периферийный блок.