Настоящий материал представляет собой обзор главы 11 книги Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026 (Copyright © 2026 Flux Capacitor, LLC).
Прямые цитаты из книги выделены отдельными блоками с пометкой «Цитата из книги». Остальной текст — систематизированное изложение содержания главы, сопоставление с актуальными данными, проверка прогнозов, а также технические детали, выходящие за пределы исходного текста и основанные на официальной документации NVIDIA и открытых источниках. Цитаты приводятся в ограниченном объёме в целях обзора и критического анализа.
Глава 11 книги Fregly завершает блок, посвящённый низкоуровневой оптимизации CUDA. Если главы 9 и 10 были сосредоточены на внутриядерных техниках — арифметической интенсивности, Tensor Cores, специализации варпов и персистентных ядрах, — то глава 11 поднимается на уровень выше: как организовать параллельную работу между ядрами, между GPU и хостом, и между самими GPU. Ключевая идея главы: производительность GPU определяется не только плотностью загрузки SM, но и одновременной работой всех аппаратных подсистем — вычислительных SM, DMA-движков и интерконнектов.
«До сих пор мы фокусировались на внутриядерных инструментах — двойной буферизации cuda::pipeline, специализации варпов, персистентных ядрах и кластерах блоков потоков — чтобы удерживать SM занятыми в пределах одного ядра. В этой главе мы сохраняем эти ядра и показываем, как организовать конвейер между ядрами и батчами с помощью потоков CUDA, событий и потоково-упорядоченного аллокатора памяти».
Глава делит материал на две взаимодополняющие части: механику потоков и аллокатора (первые разделы) и практические сценарии — двухуровневые конвейеры, кластеры, PDL и много-GPU оркестрацию (последующие разделы). В отличие от глав 9 и 10, где рассматривались отдельные ядра, здесь акцент смещается на взаимодействие всех аппаратных подсистем GPU одновременно.
«Поток CUDA — это упорядоченная очередь операций: запусков ядер, копирований памяти, аллокаций и освобождений. Операции внутри одного потока выполняются строго в порядке постановки (FIFO). Операции в разных потоках могут выполняться параллельно, если позволяют аппаратные ресурсы».
В главе подчёркивается критическая роль флага cudaStreamNonBlocking. Без него поток подчиняется правилам синхронизации с устаревшим потоком по умолчанию (stream 0), что уничтожает весь параллелизм. Отдельно рассматриваются приоритеты потоков: cudaStreamCreateWithPriority позволяет задать приоритет, который планировщик GPU учитывает при выборе, какой блок потоков запустить на освободившемся SM. Приоритет невытесняющий — работающее ядро не прерывается.
Один CPU-поток ставит ядра в два потока CUDA, ядра выполняются параллельно
ker_1 и ker_3 идут в stream1, ker_A и ker_2 — в stream2. Хост-код cpu_code_1() и cpu_code_2() выполняется параллельно с GPU-работой. Без флага cudaStreamNonBlocking этот параллелизм был бы уничтожен неявной синхронизацией с stream 0.
cudaStream_t stream1, stream2;
cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);
myKernel<<<grid, block, 0, stream1>>>(...);
cudaMemcpyAsync(dest, src, size, cudaMemcpyHostToDevice, stream2);Поток CUDA — упорядоченная FIFO-очередь. Флаг cudaStreamNonBlocking критичен для предотвращения неявной синхронизации с stream 0.
cudaStreamCreateWithPriority для latency-sensitive операций, в частности для NCCL. PyTorch использует высокоприоритетные потоки для NCCL-коммуникаций.
Явное управление потоками — первый шаг к эффективному использованию всех подсистем GPU.
«GPU имеет как минимум два независимых DMA-движка: один для Host→Device, другой для Device→Host. Вычислительные SM работают независимо от обоих. Это даёт потенциально трёхстороннее перекрытие: вычислительное ядро для батча N работает на SM, H2D-копирование батча N+1 идёт на DMA H2D, D2H-копирование батча N-1 идёт на DMA D2H».
В главе формулируются четыре условия для корректного перекрытия: pinned-память хоста, неблокирующие потоки, асинхронные аллокации и достаточные ресурсы SM. Отдельно оговаривается порог насыщения: когда вычисления полностью загружают пропускную способность памяти, а конвейер копирования пытается использовать ту же шину, обе операции замедляются примерно до 50% каждая.
Вычислительное ядро, H2D-копирование и D2H-копирование работают параллельно
Четыре условия перекрытия: pinned-память, неблокирующие потоки, асинхронные аллокации, достаточные ресурсы SM.
Когда вычисления насыщают пропускную способность памяти, конвейер копирования замедляется: обе операции делят шину.
2–4 потоков достаточно для большинства конвейеров. Большее количество увеличивает накладные расходы и риск превышения лимита резидентных сеток.
«cudaMalloc — блокирующая операция уровня устройства. Перед возвратом она выполняет глобальную синхронизацию: все ожидающие ядра и копирования во всех потоках должны завершиться… В конвейерной нагрузке вызов cudaMalloc остановит все потоки. Перекрытие, на которое вы рассчитывали, будет уничтожено».
Глава подробно разбирает внутреннее устройство пула памяти: каждый поток имеет свою очередь операций, аллокация в потоке A не блокирует поток B, освобождение возвращает блок в пул. Атрибут cudaMemPoolAttrReleaseThreshold управляет тем, сколько памяти пул удерживает перед возвратом ОС. Отдельно рассматривается интеграция с PyTorch: стандартный кэширующий аллокатор уже потоково-осведомлённый; включение бэкенда cudaMallocAsync дополнительно снижает фрагментацию.
Три потока аллоцируют и освобождают память в общем пуле без глобальной синхронизации
cudaMallocAsync(&d_data1, dataSizeBytes, stream1);
computeKernel<<<gridDim, blockDim, 0, stream1>>>(d_data1, d_result1);
cudaFreeAsync(d_data1, stream1);cudaMalloc — блокирующая операция уровня устройства. cudaMallocAsync ставит аллокацию в очередь потока без глобальной синхронизации.
cudaMemPoolAttrReleaseThreshold управляет удержанием памяти. PyTorch включается через PYTORCH_ALLOC_CONF=backend:cudaMallocAsync.
Для конвейерных нагрузок используется cudaMallocAsync. Это устраняет глобальные синхронизации.
«События CUDA — это легковесные маркеры, которые поток записывает в определённой точке своей очереди команд. Другие потоки могут ждать этих маркеров, не блокируя CPU или всё устройство… cudaLaunchHostFunc позволяет зарегистрировать функцию хоста, которая будет вызвана после завершения работы в потоке GPU».
Глава подчёркивает флаг cudaEventDisableTiming, снижающий накладные расходы, если событие используется только для синхронизации. Отдельно рассматривается критическое ограничение cudaLaunchHostFunc: из колбэка нельзя вызывать никакие CUDA API — это может привести к взаимоблокировке.
Только consumer-поток ждёт событие; CPU свободен
cudaEventDisableTiming снижает накладные расходы, если событие используется только для синхронизации.
cudaEvent_t doneEvent;
cudaEventCreateWithFlags(&doneEvent, cudaEventDisableTiming);
producerKernel<<<grid, block, 0, stream0>>>(...);
cudaEventRecord(doneEvent, stream0);
cudaStreamWaitEvent(stream1, doneEvent, 0);
consumerKernel<<<grid, block, 0, stream1>>>(...);События CUDA — легковесные маркеры для точечной синхронизации. cudaEventDisableTiming снижает накладные расходы.
Из cudaLaunchHostFunc нельзя вызывать CUDA API.
Для координации зависимостей используется cudaStreamWaitEvent — блокируется только конкретный поток.
«PDL (Programmatic Dependent Launch) — это механизм, позволяющий одному ядру сигнализировать другому о завершении, не дожидаясь полного сворачивания. Ядро B может начать выполнение ещё до того, как ядро A полностью завершится, если ядро A уже подготовило все необходимые данные».
PDL описывается как три компонента: cudaTriggerProgrammaticLaunchCompletion() в первичном ядре, cudaGridDependencySynchronize() в зависимом ядре и cudaLaunchAttributeProgrammaticStreamSerialization на стороне хоста. Эта комбинация позволяет Kernel B перекрывать эпилог Kernel A. Современные исследования показывают, что PDL даёт ускорение до 10–15% на фазах генерации токенов и до 33% при комбинировании с CUDA Graphs на последовательных слоях LLM.
B перекрывает эпилог A без обхода через CPU
cudaGridDependencySynchronize() для ожидания сигнала. Размещение триггера — это trial and error: слишком ранний триггер вызывает конкуренцию за ресурсы; слишком поздний не даёт выигрыша.
__global__ void kernel_A(float* d_ptr) {
cudaTriggerProgrammaticLaunchCompletion();
}
__global__ void kernel_B(float* d_ptr) {
cudaGridDependencySynchronize();
}
cudaLaunchAttribute attrs[1];
attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attrs[0].val.programmaticStreamSerializationAllowed = 1;PDL — межъядерное перекрытие без обхода через CPU. Три компонента: триггер, синхронизация, атрибут запуска.
PDL даёт 10–15% ускорения на генерации токенов, до 33% при комбинации с CUDA Graphs на последовательных слоях.
PDL сокращает задержку между последовательными ядрами, но требует ручной настройки размещения триггера.
«Кластеры блоков потоков позволяют группировать блоки потоков, которые могут обмениваться данными через Distributed Shared Memory (DSMEM). Один ведущий блок загружает тайл один раз, а остальные блоки в кластере читают данные из его разделяемой памяти через cluster.map_shared_rank… Кооперативные и кластерные запуски сериализуются с другими кооперативными ядрами в том же потоке».
Кластеры дают переиспользование тайлов между блоками через DSMEM; TMA-мультикаст позволяет одному DMA-запросу раздать тайл нескольким CTA. Без мультикаста каждый блок выдаёт собственный TMA-запрос — четыре отдельных чтения глобальной памяти. С мультикастом — одно чтение и вещание через DSMEM. Потоки CUDA позволяют перекрывать кластерное ядро в одном потоке с аллокациями, копированиями и некооперативными ядрами в других потоках.
Ведущий CTA загружает один тайл и вещает его остальным CTA через DSMEM
map_shared_rank. Перед использованием обязателен cluster.sync(). Кластеры ко-планируются в GPC: если размер кластера не делит число SM на GPC, часть SM простаивает.
cudaStreamBeginCapture(captureStream, cudaStreamCaptureModeGlobal);
cudaMallocAsync(&dptr, size, captureStream);
kernelA<<<grid, block, 0, captureStream>>>(dptr);
cudaStreamEndCapture(captureStream, &graph);
cudaGraphInstantiate(&execGraph, graph, nullptr, nullptr, 0);
cudaGraphLaunch(execGraph, captureStream);Кластеры дают переиспользование тайлов через DSMEM. CUDA Graphs фиксируют аллокации как узлы графа.
Комбинация CUDA Graphs + cudaMallocAsync устраняет повторяющиеся накладные расходы на аллокацию при каждой итерации.
Кооперативные и кластерные запуски сериализуются с другими кооперативными ядрами в том же потоке.
«PyTorch использует высокоприоритетные потоки для NCCL-коммуникаций».
В PyTorch DistributedDataParallel NCCL-коммуникации запускаются на выделенных высокоприоритетных потоках. Параметр is_high_priority_stream=True в ProcessGroupNCCL.Options позволяет драйверу CUDA приоритезировать NCCL-ядра, когда есть ожидающие вычислительные ядра. Это перекрытие может скрыть 70–90% времени коммуникаций на NVLink-соединённых GPU. Граф зависимостей через события связывает цепочку A→B→D, тогда как поток C работает полностью параллельно.
Поток A → поток B → поток D синхронизированы событиями; поток C независим
nccl_options = dist.ProcessGroupNCCL.Options(is_high_priority_stream=True);
dist.init_process_group("nccl", pg_options=nccl_options);
backwardLayer<<<grid, block, 0, computeStream>>>(...);
cudaEventRecord(gradReady, computeStream);
cudaStreamWaitEvent(ncclStream, gradReady, 0);
ncclAllReduce(grad, grad, size, ncclFloat, ncclSum, comm, ncclStream);NCCL на высокоприоритетном потоке с cudaStreamWaitEvent позволяет перекрывать коммуникации с вычислениями.
Перекрытие может скрыть 70–90% времени коммуникаций на NVLink-соединённых GPU.
is_high_priority_stream=True — стандартная практика в production-кластерах.
«Нагрузки обучения и инференса LLM — это сценарии, где межъядерный параллелизм через потоки CUDA даёт наибольший выигрыш. Каждый слой трансформера требует дополнительной памяти для промежуточных вычислений… Поскольку разные мини-батчи могут различаться по длине, размеры scratch-буферов варьируются».
В главе приводятся четыре сценария применения потоково-упорядоченного аллокатора в LLM: scratch-буферы переменного размера, растущий KV-кэш при авторегрессивной декодировке, послойная конвейеризация и перекрытие вычислений с NCCL-коммуникациями. С cudaMallocAsync расширение KV-кэша выполняется в том же потоке, который запускает ядро attention, без остановки других потоков.
Пока не все потоки достигли __syncthreads() — никто не продолжает
cuda::pipeline с точечными producer/consumer handoff'ами.
В LLM потоки CUDA применяются для scratch-буферов, растущего KV-кэша, послойной конвейеризации и перекрытия NCCL.
Для per-batch scratch-буферов используется cudaMallocAsync: каждый мини-батч получает ровно столько памяти, сколько нужно.
Перекрытие computation и communication достигается через cudaEventRecord + cudaStreamWaitEvent. Это позволяет all-reduce идти параллельно с вычислением следующего слоя.
«В конвейерной нагрузке вызов cudaMalloc остановит все потоки».
Подтверждено. К 2026 году cudaMallocAsync стал стандартом де-факто для production-нагрузок. PyTorch, TensorFlow и JAX поддерживают его. Исследования IEEE показывают, что асинхронный аллокатор достигает микросекундной задержки против 100–500 мкс у cudaMalloc. PyTorch PR #188110 добавил механизм trim-and-retry для обработки OOM в stream-ordered pool.
«PyTorch использует высокоприоритетные потоки для NCCL-коммуникаций».
Подтверждено. Параметр is_high_priority_stream доступен в ProcessGroupNCCL.Options с флагом по умолчанию False, но в production-кластерах он включается. Это позволяет NCCL-ядрам получать приоритет при конкуренции с вычислительными ядрами.
«Ядро B может начать выполнение ещё до того, как ядро A полностью завершится».
Подтверждено. Исследования на RTX PRO 6000 показывают 10% ускорения на фазах генерации токенов. На DGX Spark — 4–5%. Triton PR #6394 демонстрирует до 33% ускорения на последовательных слоях LLM при комбинации PDL + CUDA Graphs.
«Один ведущий блок загружает тайл один раз, а остальные блоки читают данные из его разделяемой памяти».
Подтверждено. TMA-мультикаст позволяет одному DMA-запросу раздать тайл всем CTA кластера. Без мультикаста каждый блок выдаёт собственный TMA-запрос — четыре отдельных чтения глобальной памяти. С мультикастом — одно чтение и вещание через DSMEM.
Взаимодействие с CUDA Graphs. При захвате графа cudaMallocAsync записывается как узел аллокации, а cudaFreeAsync — как узел освобождения. При воспроизведении графовые аллокации переиспользуют те же адреса. Это означает, что cudaMallocAsync и CUDA Graphs усиливают друг друга: аллокации становятся частью графа, устраняя накладные расходы при каждой итерации.
PDL требует ручной настройки. Размещение триггера launch completion — это trial and error. Слишком ранний триггер вызывает конкуренцию за ресурсы; слишком поздний не даёт выигрыша. В некоторых ядрах (например, mul_mat_vec_q) ранний триггер даёт отрицательный эффект.
Кластеры ко-планируются в GPC. Если размер кластера не делит число SM на GPC, часть SM простаивает. Blackwell вводит preferred/fallback размеры кластеров для смягчения этой квантованности.
Многопроцессная аллокация. В production-нагрузках на одном GPU часто работают несколько процессов (MPS или MIG). В этом случае каждый процесс имеет свой пул памяти, и переиспользование блоков между процессами невозможно. В мультитенантных inference-серверах применяется MIG для изоляции пулов или MPS с настройкой активных потоков.
Nsight Systems и NVTX. Временная шкала Nsight Systems может содержать тысячи событий, и идентификация проблем требует навигации по NVTX-диапазонам. Без NVTX временная шкала остаётся нечитаемой для многопоточных конвейеров.
NIXL для disaggregated inference. NIXL — библиотека для point-to-point передач в disaggregated inference — использует потоки CUDA для асинхронных передач и интегрируется с потоково-упорядоченным аллокатором. Для inference с передачей KV-кэша между prefill и decode-узлами NIXL является предпочтительным инструментом, а не NCCL send/recv.
cudaStreamNonBlocking критичен для предотвращения неявной синхронизации.cudaMallocAsync вместо cudaMalloc. Устраняет глобальные синхронизации в конвейере.cudaEventDisableTiming снижает накладные расходы.cudaMallocAsync. Аллокации фиксируются в графе, устраняя повторяющиеся накладные расходы.cudaMalloc в конвейере.Глава 11 книги Fregly — систематический разбор межъядерного параллелизма в CUDA. Ключевая идея: производительность GPU определяется не только плотностью загрузки SM, но и тем, насколько эффективно используются все подсистемы одновременно. Потоки CUDA, потоково-упорядоченный аллокатор, события, PDL, кластеры и графы — инструменты, позволяющие всем аппаратным блокам работать параллельно.
Проверка прогнозов показывает: книга оказалась точной. Потоково-упорядоченный аллокатор стал стандартом; NCCL на высокоприоритетных потоках используется в production; PDL даёт до 33% ускорения на последовательных слоях; TMA-мультикаст подтверждён исследованиями. Главный вывод остаётся неизменным: явное управление потоками, потоково-упорядоченная аллокация и точечная синхронизация через события — три принципа, переживающих поколения.
Межъядерный параллелизм с потоками CUDA — не дополнительная оптимизация, а фундамент производительности современных GPU. Без потоков GPU большую часть времени простаивает: пока копируются данные, SM молчит; пока работает ядро, DMA-движки бездействуют; пока один GPU ждёт данные от другого, его вычислительные ресурсы не используются.
Потоки CUDA, потоково-упорядоченный аллокатор, события, PDL, кластеры блоков потоков и графы CUDA — инструменты, которые позволяют всем аппаратным блокам GPU работать одновременно. В нагрузках LLM, где конвейеризация батчей, слоёв и коммуникаций определяет сквозную задержку, владение этими инструментами — разница между 30% и 90% утилизации GPU.
По мере развития оборудования параллелизм становится всё более многоуровневым: внутри варпа (Tensor Cores), внутри блока (разделяемая память), внутри кластера (DSMEM), между ядрами (PDL, потоки) и между GPU (NCCL, NIXL). Понимание того, как эти уровни наслаиваются, — ключ к извлечению максимальной производительности из высокопроизводительного AI-оборудования.
Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 11 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. Официальная документация NVIDIA CUDA 13.1 (Stream-Ordered Memory Allocator, CUDA Graphs); IEEE IPDPSW 2026 (Quantitative Characterization of CUDA Memory Allocators); GitHub PR #22522 (llama.cpp PDL), PR #6394 (Triton PDL), PR #188110 (PyTorch cudaMallocAsync trim-and-retry); DAK arXiv 2604.26074 (TMA Multicast).
This material is a review of Chapter 11 of Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026 (Copyright © 2026 Flux Capacitor, LLC).
Direct quotes from the book are marked with separate blocks labeled «Quote from the book». The rest of the text is a systematized exposition of the chapter's content, a comparison with current data, forecast verification, and technical details beyond the original text based on official NVIDIA documentation and open sources. Quotes are provided in limited volume for review and critical analysis purposes.
Chapter 11 of Fregly's book completes the block dedicated to low-level CUDA optimization. While Chapters 9 and 10 focused on intra-kernel techniques — arithmetic intensity, Tensor Cores, warp specialization, and persistent kernels — Chapter 11 rises to the next level: how to organize parallel work between kernels, between GPU and host, and between GPUs themselves. The key idea: GPU performance is determined not only by how tightly individual SMs are loaded, but also by how efficiently all hardware subsystems — compute SMs, DMA engines and interconnects — are used simultaneously.
«So far, we have focused on the intra-kernel tools — cuda::pipeline double-buffering, warp specialization, persistent kernels, and thread-block clusters — to keep the SMs busy for a single kernel. In this chapter we keep those kernels and show how to pipeline across kernels and batches with CUDA streams, events, and the stream-ordered memory allocator».
The chapter divides the material into two complementary parts: the mechanics of streams and the allocator (first sections) and practical scenarios — two-level pipelines, clusters, PDL and multi-GPU orchestration (later sections). Unlike Chapters 9 and 10, which examined individual kernels, here the emphasis shifts to the interaction of all GPU hardware subsystems simultaneously.
«A CUDA stream is an ordered queue of operations: kernel launches, memory copies, allocations, and deallocations. Operations within a stream execute strictly in issue order (FIFO). Operations in different streams may execute in parallel if hardware resources allow».
The chapter emphasizes the critical role of the cudaStreamNonBlocking flag. Without it, the stream is subject to stream 0 synchronization rules, which destroys all parallelism. Stream priorities via cudaStreamCreateWithPriority are considered separately: the scheduler uses priority when choosing which thread block to launch on a freed SM. Priority is non-preemptive — a running kernel is not interrupted.
One CPU thread enqueues kernels into two CUDA streams; kernels execute in parallel
ker_1 and ker_3 go into stream1, ker_A and ker_2 — into stream2. Host code cpu_code_1() and cpu_code_2() runs in parallel with GPU work. Without the cudaStreamNonBlocking flag, this parallelism would be destroyed by implicit synchronization with stream 0.
cudaStream_t stream1, stream2;
cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);
myKernel<<<grid, block, 0, stream1>>>(...);
cudaMemcpyAsync(dest, src, size, cudaMemcpyHostToDevice, stream2);A CUDA stream is an ordered FIFO queue. cudaStreamNonBlocking is critical to prevent implicit synchronization with stream 0.
cudaStreamCreateWithPriority for latency-sensitive operations, in particular NCCL. PyTorch uses high-priority streams for NCCL communications.
Explicit stream management is the first step to efficient utilization of all GPU subsystems.
«The GPU has at least two independent DMA engines: one for Host→Device, one for Device→Host. Compute SMs operate independently of both, giving potentially three-way overlap: the compute kernel for batch N runs on SMs, H2D copy of batch N+1 goes on DMA H2D, D2H copy of batch N-1 goes on DMA D2H».
The chapter formulates four conditions for correct overlap: pinned host memory, non-blocking streams, async allocations, and sufficient SM resources. The chapter separately stipulates the saturation threshold: when computations fully load memory bandwidth and the copy pipeline tries to use the same bus, both operations slow down to about 50% each.
Compute kernel, H2D copy, and D2H copy run in parallel
Four conditions for overlap: pinned memory, non-blocking streams, async allocations, sufficient SM resources.
When computations saturate memory bandwidth, the copy pipeline slows down: both operations share the bus.
2–4 streams suffice for most pipelines. More increases overhead and risks exceeding the resident grid limit.
«cudaMalloc is a device-wide blocking operation. Before returning, it performs global synchronization: all pending kernels and copies in all streams must complete… In a pipelined workload, a cudaMalloc call will stop all streams. The overlap you counted on is destroyed».
The chapter explains the internal structure of the memory pool: each stream has its own operation queue, an allocation in stream A does not block stream B, a free returns the block to the pool. The cudaMemPoolAttrReleaseThreshold attribute controls memory retention. PyTorch integration is considered separately: the default caching allocator is already stream-aware; enabling the cudaMallocAsync backend further reduces fragmentation.
Three streams allocate and free memory in a shared pool without global synchronization
cudaMallocAsync(&d_data1, dataSizeBytes, stream1);
computeKernel<<<gridDim, blockDim, 0, stream1>>>(d_data1, d_result1);
cudaFreeAsync(d_data1, stream1);cudaMalloc is a device-wide blocking operation. cudaMallocAsync enqueues allocation on the stream without global synchronization.
cudaMemPoolAttrReleaseThreshold controls memory retention. PyTorch enables it via PYTORCH_ALLOC_CONF=backend:cudaMallocAsync.
Pipelined workloads use cudaMallocAsync. This eliminates global synchronizations.
«CUDA events are lightweight markers that a stream records at a point in its command queue. Other streams can wait on these markers without blocking the CPU or the whole device… cudaLaunchHostFunc lets you register a host function that will be called after work in a GPU stream completes».
The chapter emphasizes the cudaEventDisableTiming flag, which reduces overhead when the event is used only for synchronization. A critical limitation of cudaLaunchHostFunc is considered separately: you cannot call any CUDA API from the callback — this may deadlock.
Only the consumer stream waits for the event; the CPU is free
cudaEventDisableTiming flag reduces overhead when the event is used only for synchronization.
cudaEvent_t doneEvent;
cudaEventCreateWithFlags(&doneEvent, cudaEventDisableTiming);
producerKernel<<<grid, block, 0, stream0>>>(...);
cudaEventRecord(doneEvent, stream0);
cudaStreamWaitEvent(stream1, doneEvent, 0);
consumerKernel<<<grid, block, 0, stream1>>>(...);CUDA events are lightweight markers for fine-grained synchronization. cudaEventDisableTiming reduces overhead.
CUDA APIs cannot be called from cudaLaunchHostFunc.
For coordinating dependencies, cudaStreamWaitEvent is used — it blocks only the specific stream.
«PDL (Programmatic Dependent Launch) lets one kernel signal another about completion without waiting for a full drain. Kernel B may begin before Kernel A fully finishes, provided A has already produced the necessary data».
PDL is described as three components: cudaTriggerProgrammaticLaunchCompletion() in the primary kernel, cudaGridDependencySynchronize() in the dependent kernel, and cudaLaunchAttributeProgrammaticStreamSerialization on the host side. Modern research shows PDL delivers up to 10–15% speedup on token generation and up to 33% when combined with CUDA Graphs on successive LLM layers.
B overlaps A's epilogue without a CPU round-trip
cudaGridDependencySynchronize() to wait for the signal. Trigger placement is trial and error: too early causes resource contention; too late gives no benefit.
__global__ void kernel_A(float* d_ptr) {
cudaTriggerProgrammaticLaunchCompletion();
}
__global__ void kernel_B(float* d_ptr) {
cudaGridDependencySynchronize();
}
cudaLaunchAttribute attrs[1];
attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attrs[0].val.programmaticStreamSerializationAllowed = 1;PDL is inter-kernel overlap without a CPU round-trip. Three components: trigger, synchronization, launch attribute.
PDL delivers 10–15% speedup on token generation, up to 33% when combined with CUDA Graphs on successive layers.
PDL reduces latency between successive kernels but requires manual placement tuning of the trigger.
«Thread block clusters let groups of thread blocks share data through Distributed Shared Memory (DSMEM). A leader block loads a tile once, and other blocks read from its shared memory via cluster.map_shared_rank… Cooperative and cluster launches serialize with other cooperative kernels in the same stream».
Clusters provide tile reuse across blocks via DSMEM; TMA multicast lets one DMA request broadcast a tile to multiple CTAs. Without multicast, each block issues its own TMA request — four separate global memory reads. With multicast — one read and broadcast via DSMEM. CUDA streams allow overlapping a cluster kernel in one stream with allocations, copies, and non-cooperative kernels in other streams.
The leader CTA loads a tile once and broadcasts it to other CTAs via DSMEM
map_shared_rank. cluster.sync() is mandatory before use. Clusters are co-scheduled within a GPC: if the cluster size doesn't divide the number of SMs per GPC, some SMs idle.
cudaStreamBeginCapture(captureStream, cudaStreamCaptureModeGlobal);
cudaMallocAsync(&dptr, size, captureStream);
kernelA<<<grid, block, 0, captureStream>>>(dptr);
cudaStreamEndCapture(captureStream, &graph);
cudaGraphInstantiate(&execGraph, graph, nullptr, nullptr, 0);
cudaGraphLaunch(execGraph, captureStream);Clusters provide tile reuse via DSMEM. CUDA Graphs fix allocations as graph nodes.
The combination CUDA Graphs + cudaMallocAsync eliminates repeated memory allocation overhead on each iteration.
Cooperative and cluster launches serialize with other cooperative kernels in the same stream.
«PyTorch uses high-priority streams for NCCL communications».
In PyTorch DistributedDataParallel, NCCL communications run on dedicated high-priority streams. The is_high_priority_stream=True option in ProcessGroupNCCL.Options lets the CUDA driver prioritize NCCL kernels when there are waiting compute kernels. This overlap can hide 70–90% of communication time on NVLink-connected GPUs. The dependency graph via events links the chain A→B→D, while stream C runs fully in parallel.
Stream A → stream B → stream D synchronized by events; stream C is independent
nccl_options = dist.ProcessGroupNCCL.Options(is_high_priority_stream=True);
dist.init_process_group("nccl", pg_options=nccl_options);
backwardLayer<<<grid, block, 0, computeStream>>>(...);
cudaEventRecord(gradReady, computeStream);
cudaStreamWaitEvent(ncclStream, gradReady, 0);
ncclAllReduce(grad, grad, size, ncclFloat, ncclSum, comm, ncclStream);NCCL on a high-priority stream with cudaStreamWaitEvent allows overlapping communications with computations.
Overlap can hide 70–90% of communication time on NVLink-connected GPUs.
is_high_priority_stream=True is standard practice in production clusters.
«LLM training and inference workloads are scenarios where inter-kernel parallelism via CUDA streams delivers the greatest benefit. Each transformer layer requires additional memory for intermediate computations… Since different mini-batches may vary in length, scratch buffer sizes vary».
The chapter cites four scenarios for using the stream-ordered allocator in LLMs: variable-size scratch buffers, growing KV cache during autoregressive decoding, layer pipelining, and overlapping computation with NCCL communications. With cudaMallocAsync, KV cache expansion happens in the same stream that launches the attention kernel, without stopping other streams.
Until all threads reach __syncthreads(), nobody continues
cuda::pipeline with fine-grained producer/consumer handoffs is used.
In LLMs, CUDA streams are used for scratch buffers, growing KV cache, layer pipelining, and NCCL overlap.
For per-batch scratch buffers, cudaMallocAsync is used: each mini-batch gets exactly the memory it needs.
Overlapping computation and communication is achieved via cudaEventRecord + cudaStreamWaitEvent. This lets all-reduce proceed in parallel with computing the next layer.
«In a pipelined workload, a cudaMalloc call will stop all streams».
Confirmed. By 2026, cudaMallocAsync became the de facto standard for production workloads. PyTorch, TensorFlow, and JAX support it. IEEE research shows the async allocator achieves microsecond-scale latency versus 100–500 µs for cudaMalloc. PyTorch PR #188110 added a trim-and-retry mechanism for handling OOM in the stream-ordered pool.
«PyTorch uses high-priority streams for NCCL communications».
Confirmed. The is_high_priority_stream option is available in ProcessGroupNCCL.Options with a default of False, but it is enabled in production clusters. This allows NCCL kernels to get priority when competing with compute kernels.
«Kernel B may begin before Kernel A fully finishes».
Confirmed. Studies on RTX PRO 6000 show 10% speedup on token generation. On DGX Spark — 4–5%. Triton PR #6394 demonstrates up to 33% speedup on successive LLM layers when combining PDL + CUDA Graphs.
«A leader block loads a tile once, and other blocks read from its shared memory».
Confirmed. TMA multicast lets one DMA request broadcast a tile to all CTAs in a cluster. Without multicast, each block issues its own TMA request — four separate global memory reads. With multicast — one read and broadcast via DSMEM.
Interaction with CUDA Graphs. During stream capture, cudaMallocAsync is recorded as an allocation node, cudaFreeAsync as a free node. On replay, graph allocations reuse the same addresses. This means cudaMallocAsync and CUDA Graphs reinforce each other: allocations become part of the graph, eliminating overhead on each iteration.
PDL requires manual tuning. Placement of the launch completion trigger is trial and error. Too early a trigger causes resource contention; too late gives no benefit. In some kernels (e.g., mul_mat_vec_q), an early trigger has a negative effect.
Clusters co-schedule within a GPC. If the cluster size doesn't divide the number of SMs per GPC, some SMs idle. Blackwell introduces preferred/fallback cluster sizes to soften this quantization.
Multi-process allocation. In production workloads, multiple processes often run on one GPU (MPS or MIG). In this case, each process has its own memory pool, and block reuse between processes is impossible. In multi-tenant inference servers, MIG is used for pool isolation or MPS with active thread tuning.
Nsight Systems and NVTX. The Nsight Systems timeline may contain thousands of events, and identifying problems requires navigating NVTX ranges. Without NVTX, the timeline remains unreadable for multithreaded pipelines.
NIXL for disaggregated inference. NIXL — a library for point-to-point transfers in disaggregated inference — uses CUDA streams for asynchronous transfers and integrates with the stream-ordered allocator. For inference with KV cache transfer between prefill and decode nodes, NIXL is the preferred tool rather than NCCL send/recv.
cudaStreamNonBlocking is critical to prevent implicit synchronization.cudaMallocAsync instead of cudaMalloc. Eliminates global synchronizations in the pipeline.cudaEventDisableTiming reduces overhead.cudaMallocAsync. Allocations are fixed in the graph, eliminating repeated overhead.cudaMalloc in the pipeline.Chapter 11 of Fregly's book is a systematic review of inter-kernel parallelism in CUDA. The key idea: GPU performance is determined not only by how tightly individual SMs are loaded, but also by how efficiently all subsystems are used simultaneously. CUDA streams, the stream-ordered allocator, events, PDL, clusters, and graphs are tools that let all hardware blocks work in parallel.
Forecast verification shows: the book proved accurate. The stream-ordered allocator became standard; NCCL on high-priority streams is used in production; PDL delivers up to 33% speedup on successive layers; TMA multicast is confirmed by research. The main takeaway remains unchanged: explicit stream management, stream-ordered allocation, and fine-grained event synchronization are three principles that survive generations.
Inter-kernel parallelism with CUDA streams is not an add-on optimization but the foundation of modern GPU performance. Without streams, the GPU idles most of the time: while data is copied, SMs are silent; while a kernel runs, DMA engines idle; while one GPU waits for data from another, its compute resources go unused.
CUDA streams, the stream-ordered allocator, events, PDL, thread block clusters, and CUDA graphs are tools that let all hardware blocks work simultaneously. In LLM workloads where pipelining batches, layers, and communications determines end-to-end latency, mastery of these tools is the difference between 30% and 90% GPU utilization.
As hardware evolves, parallelism becomes more layered: within a warp (Tensor Cores), within a block (shared memory), within a cluster (DSMEM), between kernels (PDL, streams), and between GPUs (NCCL, NIXL). Understanding how these levels stack is the key to extracting maximum performance from high-performance AI hardware.
Sources: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Chapter 11 — quotes provided in limited volume for review and critical analysis purposes. Official NVIDIA CUDA 13.1 documentation (Stream-Ordered Memory Allocator, CUDA Graphs); IEEE IPDPSW 2026 (Quantitative Characterization of CUDA Memory Allocators); GitHub PR #22522 (llama.cpp PDL), PR #6394 (Triton PDL), PR #188110 (PyTorch cudaMallocAsync trim-and-retry); DAK arXiv 2604.26074 (TMA Multicast).
