Настоящий материал представляет собой обзор шестой главы книги 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).
Прямые цитаты из книги выделены отдельными блоками с пометкой «Цитата из книги». Остальной текст — систематизированное изложение содержания главы, сопоставление с актуальными данными об архитектуре Blackwell, проверка прогнозов по открытым источникам, а также дополнительные замечания, выходящие за пределы текста главы. Цитаты приводятся в ограниченном объёме в целях обзора и критического анализа.
Шестая глава книги Fregly посвящена переходу от высокоуровневых представлений об AI-системах к конкретике исполнения на GPU. В ней разбирается модель SIMT, иерархия потоков и памяти, способы настройки запуска ядер, концепция occupancy, модель roofline и инструменты профилирования. Эта глава — фундамент для всех последующих глав, посвящённых оптимизации CUDA-ядер, распределённому обучению и inference. Ниже последовательно изложены ключевые положения главы, приведены визуализации, примеры кода и уточнения по актуальным данным об архитектуре Blackwell.
«Каждый GPU состоит из множества SM, которые примерно аналогичны ядрам CPU, но оптимизированы под параллелизм. Каждый SM на Blackwell может отслеживать до 64 варпов (групп по 32 потока)».
«Каждый GPU Blackwell включает множество SM — подобных ядрам CPU, но оптимизированных под пропускную способность. На современных GPU каждый SM отслеживает до 64 варпов (2 048 потоков) одновременно. GPU Blackwell располагают 64K 32-битных регистров на SM (256 КБ суммарно) и объединёнными 256 КБ кэша L1 и разделяемой памяти на SM. До 228 КБ (227 КБ полезных) из этой SRAM могут быть сконфигурированы как управляемая пользователем разделяемая память на SM».
Потоковый мультипроцессор (SM) — основная вычислительная единица GPU. В отличие от CPU-ядра, оптимизированного под одну последовательную задачу, SM устроен так, чтобы одновременно держать в работе сотни потоков. На Blackwell каждый SM хранит до 64 активных варпов (групп по 32 потока), что даёт 2048 резидентных потоков.
Регистровый файл — 64K 32-битных регистров (256 КБ), из которых одному потоку доступно не более 255. Унифицированный блок SRAM объёмом 256 КБ на SM делится между кэшем L1 и разделяемой памятью, причём до 228 КБ (227 КБ полезных) можно выделить под разделяемую память. Это ключевое архитектурное решение: разработчик может динамически выбирать, что важнее для конкретного ядра — кэш для глобальных данных или программно управляемая разделяемая память для обмена между потоками блока.
«Внутри SM Blackwell несколько планировщиков варпов выдают инструкции доступным конвейерам; четыре независимых планировщика варпов позволяют до четырём варпам выдавать инструкции доступным конвейерам каждый такт. Более того, каждый планировщик поддерживает dual-issue — возможность выдать две независимые инструкции (например, одну арифметическую и одну операцию с памятью) за один варп».
SM Blackwell разделён на четыре независимых планировщика варпов. Каждый такт каждый планировщик выбирает один готовый варп и может выдать две инструкции: одну арифметическую (INT32, FP32 или Tensor Core) и одну операцию с памятью (загрузка или сохранение). Теоретический пик — четыре варпа × две инструкции = восемь выдач за такт. Это возможно, только если двойная выдача (dual-issue) исходит от одного и того же варпа — комбинировать инструкции из разных варпов в одной выдаче нельзя.
Блок специальных функций (SFU) обрабатывает трансцендентные операции — синус, косинус, обратный квадратный корень и т. д. У SFU собственный конвейер, отдельный от INT32/FP32 и от загрузки/сохранения. Благодаря этому медленные трансцендентные вычисления не блокируют основные вычислительные и памятьные конвейеры, а параллелизм на уровне инструкций возрастает.
Важное уточнение для современного поколения: на Blackwell FP32 и INT32 объединены в единые CUDA-ядра, поэтому в одном такте на одном ядре выполняется либо FP32-, либо INT32-инструкция. Mixed INT32/FP32-потоки больше не выигрывают от dual-issue на одном ядре — вместо этого они используют параллелизм варпов и планировщиков. Этот нюанс меняет стратегию оптимизации ILP: вместо чередования INT и FP в одном потоке выгоднее держать в полёте две независимые FP-инструкции.
Четыре независимых планировщика, по два конвейера на каждый
| Параметр планировщика | Значение на Blackwell |
|---|---|
| Количество планировщиков | 4 |
| Варпов за такт (максимум) | 4 (по одному на планировщик) |
| Арифметических инструкций за такт | 4 |
| Операций с памятью за такт | 4 |
| Всего выдач за такт (с dual-issue) | до 8 |
| Максимум регистров на поток | 255 |
| Максимум потоков на блок | 1024 |
| Максимум варпов на SM | 64 |
SM — throughput-ориентированный процессор, разделённый на четыре независимых планировщика. Каждый такт возможна выдача до двух инструкций из одного варпа. Пик — восемь выдач за такт при полной загрузке.
Регистров на поток — до 255, потоков на блок — до 1024, варпов на SM — 64. Эти числа определяют, сколько работы поместится на один SM и как будет выглядеть occupancy.
На Blackwell FP32 и INT32 объединены в единые CUDA-ядра. Mixed-потоки следует оптимизировать через параллелизм варпов, а не через чередование типов инструкций в одном потоке.
«CUDA структурирует параллельную работу в трёхуровневую иерархию — потоки, блоки потоков (они же cooperative thread arrays [CTA]) и сетки — чтобы уравновесить программируемость и массовую пропускную способность. На нижнем уровне каждый поток выполняет код вашего ядра. Вы группируете потоки в блоки потоков размером до 1 024 потоков на современных GPU. Блоки потоков образуют сетку при запуске ядра».
«Блоки потоков подразделяются на варпы по 32 потока, которые выполняются синхронно в модели SIMT с использованием планировщика варпов».
CUDA организует параллельную работу в три уровня: поток → блок потоков (CTA) → сетка. Блок содержит до 1024 потоков; блоки образуют сетку, запускаемую одним вызовом ядра. Внутри блока потоки делят низколатентную разделяемую память и синхронизируются барьером __syncthreads(). Блоки выполняются независимо и в произвольном порядке — планировщик GPU сам распределяет их по SM.
Каждый блок делится на варпы по 32 потока, которые выполняются синхронно в модели SIMT: все потоки варпа проходят один и тот же путь управления. Если часть потоков уходит в ветвь if, а другая — в else, варп выполняет обе ветви последовательно, замаскировав неактивные линии. Это называется расхождением варпов. Разные варпы могут идти по разным ветвям без потери производительности — расхождение затрагивает только потоки внутри одного варпа.
«Традиционно потоки из разных блоков потоков не могли взаимодействовать друг с другом напрямую. Однако современные архитектуры GPU и версии CUDA поддерживают кластеры блоков потоков. Кластеры блоков потоков — это группы блоков потоков, которые могут обмениваться данными друг с другом между SM».
Кластеры блоков — относительно новая возможность, появившаяся в Hopper и развитая в Blackwell. Внутри кластера потоки из разных блоков могут обращаться к разделяемой памяти друг друга через распределённую разделяемую память (DSMEM). Это позволяет строить более крупные вычислительные конструкции, чем один блок, и синхронизировать их через барьеры уровня кластера, не прибегая к глобальной памяти.
Практический смысл: для задач вроде матричных умножений с большими тайлами можно использовать кластер из двух или четырёх блоков, чтобы эффективно увеличить размер tile и повысить арифметическую интенсивность без выхода за пределы одного SM.
Три уровня иерархии параллелизма CUDA
32 потока, движущихся в lockstep
Ветвление внутри варпа сериализует исполнение
if/else, варп выполняет обе последовательно. Разные варпы расходятся без штрафа.DSMEM: блоки кластера видят SMEM друг друга
Три уровня иерархии — поток, блок, сетка. Внутри блока потоки делят shared memory и синхронизируются через __syncthreads(). Блоки независимы. Варп — 32 потока, выполняющих одни инструкции синхронно.
Если потоки одного варпа идут по разным ветвям if/else, варп выполняет обе ветви последовательно, маскируя неактивные линии. Расхождение затрагивает только потоки внутри одного варпа.
Начиная с Hopper, блоки могут объединяться в кластеры и обращаться к shared memory друг друга через DSMEM. Это позволяет увеличить эффективный размер tile без выхода за пределы одного SM.
«Когда вы вызываете ядро из кода CPU (хоста), вы используете синтаксис <<< >>> "шеврон" для указания того, сколько потоков должно выполняться — и как они организованы — с помощью двух параметров конфигурации: blocksPerGrid для числа блоков потоков и threadsPerBlock для числа потоков внутри каждого блока».
«Критически важный аспект производительности GPU — выбор размера блока потоков, согласованного с аппаратным размером варпа в 32 потока. Как следствие, обычно выбирают размеры блоков потоков, кратные 32».
Вызов ядра в CUDA выглядит как kernel<<<blocksPerGrid, threadsPerBlock>>>(...). Выбор этих двух параметров напрямую влияет на occupancy — долю аппаратных слотов потоков, занятых активными варпами.
Размер блока должен быть кратен 32 (размеру варпа). Блок из 256 потоков — восемь полных варпов — это универсальная отправная точка. Такой размер даёт несколько преимуществ: нет недозаполненных варпов, хорошая occupancy (до восьми блоков на SM), баланс между ресурсами и параллелизмом.
Для blocksPerGrid используется формула (N + threadsPerBlock - 1) / threadsPerBlock, которая округляет вверх и гарантирует, что все элементы массива будут покрыты. Обязательна проверка границ if (idx < N): без неё «лишние» потоки последнего блока попытаются обратиться к памяти за пределами массива и вызовут cudaErrorIllegalAddress.
Host запускает kernel, GPU его исполняет
Host → launch → GPU grid
Размер блока должен быть кратен 32 (размеру варпа). Блок из 256 потоков — восемь полных варпов — это универсальная отправная точка. Причины:
__global__ void doubleKernel(float* data, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
data[idx] *= 2.0f;
}
}
int main() {
const int N = 1'000'000;
float *h_data, *d_data;
cudaMallocHost(&h_data, N * sizeof(float));
for (int i = 0; i < N; ++i) h_data[i] = 1.0f;
cudaMalloc(&d_data, N * sizeof(float));
cudaMemcpy(d_data, h_data, N * sizeof(float),
cudaMemcpyHostToDevice);
const int threadsPerBlock = 256;
const int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
doubleKernel<<>>(d_data, N);
cudaDeviceSynchronize();
cudaFree(d_data);
cudaFreeHost(h_data);
return 0;
} Проверка границ if (idx < N) критически важна: без неё «лишние» потоки последнего блока попытаются обратиться к памяти за пределами массива и вызовут cudaErrorIllegalAddress. CUDA не генерирует исключения внутри ядер — ошибка устанавливается в глобальный флаг и проявляется при следующем вызове синхронизации на хосте.
Для двумерных данных используется двумерная сетка блоков. Например, обработка изображения 1024×1024 с блоком 16×16 (256 потоков):
__global__ void double2DKernel(float* image, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
int idx = y * width + x;
image[idx] *= 2.0f;
}
}
int main() {
const int width = 1024;
const int height = 1024;
const int N = width * height;
float *h_image, *d_image;
cudaMallocHost(&h_image, N * sizeof(float));
cudaMalloc(&d_image, N * sizeof(float));
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid((width + 15) / 16, (height + 15) / 16);
double2DKernel<<>>(d_image, width, height);
cudaDeviceSynchronize();
cudaFree(d_image);
cudaFreeHost(h_image);
return 0;
} Размер блока должен быть кратен 32. 256 потоков — восьмиварповый блок — универсальная отправная точка. Формула (N + threadsPerBlock - 1) / threadsPerBlock округляет вверх и покрывает все элементы.
Обязательна проверка границ if (idx < N). Без неё лишние потоки последнего блока вызовут ошибку доступа к памяти.
Для изображений и объёмных данных используются 2D- и 3D-сетки через dim3. Это позволяет естественным образом отображать многомерные данные на потоки GPU.
«GPU предоставляет многоуровневую иерархию памяти и помогает балансировать ёмкость и скорость. Иерархия включает регистры, разделяемую память, кэши, глобальную память и специализированную TMEM на GPU Blackwell и последующих. TMEM — это выделенная внутрикристальная память объёмом ~256 КБ на SM, используемая инструкциями Tensor Core 5-го поколения на Blackwell».
Уровни памяти GPU отличаются по ёмкости, задержке и пропускной способности на несколько порядков. Понимание этой лестницы — ключ к написанию быстрых ядер.
На вершине — регистры: однотактовые, но крайне ограниченные (до 255 на поток). Чуть ниже — разделяемая память и L1-кэш: 20–30 тактов, до 228 КБ на SM. Ещё ниже — L2-кэш: 126 МБ, общий для всего GPU, ~200 тактов. На дне — глобальная память HBM3e: сотни и тысячи тактов, ~8 ТБ/с. Чем дальше от вычислительного ядра, тем больше ёмкость и тем выше задержка.
Отдельно стоит выделить Tensor Memory (TMEM) — специализированную память на 256 КБ на SM, которая используется инструкциями Tensor Core нового поколения. TMEM не адресуется обычными указателями; обмен идёт через инструкции tcgen05. Это ещё один уровень иерархии, который появился с приходом Tensor Cores и Transformer Engine.
Пирамида: от регистров к HBM
Разделяемая память и тензорная память на одном SM
tcgen05.Программно управляемый кэш на SM
Когда регистров не хватает — данные уходят в DRAM
Глобальная память — основной канал данных
| Уровень | Область | Ёмкость (Blackwell) | Задержка | Пропускная способность |
|---|---|---|---|---|
| Регистры | На поток, на SM | 64K × 32 бит (макс. 255 на поток) | 1 такт | Десятки ТБ/с на SM |
| Разделяемая память / L1 | На SM | 228 КБ (227 полезных) | 20–30 тактов | ТБ/с на SM |
| TMEM | На SM | 256 КБ | ~10 тактов | ТБ/с к Tensor Core |
| Кэш констант | На SM | ~8 КБ (для 64 КБ __constant__) | ~1 такт (broadcast) | Как у регистра |
| L2 | На весь GPU | 126 МБ | ~200 тактов | Десятки ТБ/с |
| Глобальная (HBM3e) | На устройство | 180 ГБ (B200) / ~288 ГБ (B300) | Сотни–тысячи тактов | ~8 ТБ/с |
Иерархия памяти GPU — регистры → shared/L1 → L2 → HBM. Для максимальной производительности данные должны оставаться как можно выше в этой иерархии.
Register spill — одна из главных причин падения производительности. Переполнение регистрового файла вытесняет данные в локальную память, что эквивалентно обращению к DRAM. Профилировать и устранять spill необходимо в первую очередь.
«Unified Memory (также известная как CUDA Managed Memory) даёт вам единое когерентное адресное пространство, охватывающее и CPU, и GPU, так что вам больше не нужно жонглировать отдельными буферами хоста и устройства или явно вызывать cudaMemcpy».
«Хотя работа с Unified Memory очень удобна для разработчика, она может вызывать нежелательные миграции страниц по требованию между CPU и GPU. Это порождает скрытые задержки и простои исполнения».
CUDA Managed Memory (Unified Memory) создаёт единое адресное пространство для CPU и GPU. Вы выделяете память через cudaMallocManaged() — и обе стороны видят один указатель. Страницы мигрируют по требованию: если GPU обращается к данным, находящимся в ОЗУ CPU, происходит страничная ошибка и данные переносятся через NVLink-C2C (на Grace Blackwell) или PCIe.
Удобство обманчиво. Автоматические миграции по требованию могут приводить к скрытым задержкам: если поток GPU обращается к странице, которая физически лежит в памяти CPU, ядро останавливается и ждёт переноса. На системах с PCIe это может быть на порядки медленнее, чем обращение к HBM. На Grace Blackwell с NVLink-C2C переносы быстрее, но всё равно не бесплатны.
Решение — явное управление размещением: cudaMemPrefetchAsync для упреждающей миграции страниц на GPU и cudaMemAdvise для подсказок драйверу о том, где данные будут использоваться. Это превращает миграции из синхронных страничных ошибок в асинхронные операции, которые можно перекрыть с вычислениями.
Страницы мигрируют между CPU и GPU по требованию
cudaMemPrefetchAsync — упреждающая миграция
cudaMemAdvise: preferred location, ReadMostly
cudaMemAdvise подсказывает драйверу, где выгоднее держать страницы и как их использовать. Снижает количество миграций.cudaMemPrefetchAsync(ptr, size, gpuId, stream);
cudaMemAdvise(ptr, size, cudaMemAdviseSetPreferredLocation, gpuId);
cudaMemAdvise(ptr, size, cudaMemAdviseSetReadMostly, gpuId);
cudaStreamAttachMemAsync(stream, ptr, 0, cudaMemAttachSingle);Unified Memory — удобно, но требует явного управления размещением. Автоматические миграции по страничным ошибкам могут вносить скрытые задержки.
Используйте cudaMemPrefetchAsync для упреждающей миграции и cudaMemAdvise для подсказок драйверу. Это превращает синхронные миграции в асинхронные, которые можно перекрыть с вычислениями.
«GPU поддерживают производительность за счёт одновременного выполнения множества варпов, так что когда один варп останавливается в ожидании данных, может выполняться другой. Эта способность быстро переключаться между варпами позволяет GPU скрывать задержку памяти. Как мы описывали ранее, доля ёмкости SM, фактически занятая активными варпами, называется occupancy».
«Если occupancy низкая (всего несколько активных варпов), SM может простаивать, пока один варп ждёт памяти. Это приводит к плохой загрузке SM».
Occupancy — отношение активных варпов на SM к теоретическому максимуму (64 варпа, 2048 потоков на Blackwell). Высокая occupancy позволяет планировщику моментально переключаться на готовый варп, когда текущий остановился на ожидании памяти. Это и есть скрытие задержки.
Но occupancy — не самоцель. Если ядро ограничено пропускной способностью памяти (memory-bound), доведение occupancy до 100% не поможет — узкое место в шине HBM, а не в количестве варпов. Более информативные метрики — eligible warps per cycle и warp stall reasons: сколько варпов реально готовы к исполнению и на чём останавливаются остальные.
Пример из главы показывает разницу между последовательной и параллельной реализацией сложения векторов. Последовательная версия использует один поток и загружает GPU на 1,5%. Параллельная распределяет работу и загружает GPU на 95%. Это наглядная иллюстрация того, что GPU — throughput-процессор и для его загрузки нужна массовая параллельность.
Параллельная версия делает то же за меньшее время
A + B = C — по одному потоку на элемент
| Метрика | addSequential | addParallel |
|---|---|---|
| Время ядра (мс) | 48,21 | 2,17 |
| Загрузка GPU | 1,5% | 95% |
| Occupancy | 1,3% | 38,7% |
| Эффективность варпов | 3,1% | 100% |
Для управления occupancy на уровне компилятора используются аннотации __launch_bounds__. Они подсказывают компилятору, с каким максимальным числом потоков на блок будет запускаться ядро и сколько блоков минимум должно помещаться на SM. Компилятор ограничивает регистры на поток, чтобы эти условия выполнились. Runtime-версия — CUDA Occupancy API:
__global__ __launch_bounds__(256, 16)
void myKernel(...) { /* ... */ }
int minGrid = 0, bestBlock = 0;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &bestBlock,
myKernel,
0,
0);Occupancy — средство для скрытия задержки, а не цель. Высокая occupancy помогает, только если узкое место — задержка, а не пропускная способность памяти.
eligible warps per cycle и warp stall reasons показывают, сколько варпов реально готовы к исполнению и на чём останавливаются остальные. Именно на них следует ориентироваться при тюнинге.
Для управления occupancy используйте __launch_bounds__ и CUDA Occupancy API. Помните, что увеличение occupancy за счёт регистров может привести к spill и падению производительности.
«Обращения к глобальной памяти быстрее всего, когда потоки варпа обращаются к непрерывным адресам памяти, которые аппаратура может объединить в меньшее число более крупных транзакций. Если потоки обращаются к разбросанным или невыровненным адресам, устройство не может объединить запросы в минимальное число транзакций кэш-линий, которые на современных GPU представляют собой 128-байтовые линии, состоящие из четырёх 32-байтовых секторов».
Когда 32 потока варпа одновременно читают глобальную память, GPU пытается объединить (коалесцировать) эти обращения в минимальное количество транзакций по 128 байт. Если потоки обращаются к последовательным 32-битным элементам, одна транзакция обслуживает весь варп. Если потоки разбросаны по памяти — каждый элемент требует отдельной транзакции.
Аналогичная проблема возникает в разделяемой памяти, но там она называется конфликтами банков. Разделяемая память состоит из 32 банков по 4 байта. Если несколько потоков варпа обращаются к одному банку, доступы сериализуются. Классический случай — доступ по столбцам к матрице, хранящейся по строкам. Решение — добавить padding: один лишний элемент на строку. Для матрицы 32×32, хранящейся как smem[32][33], обращения по столбцам не создают конфликтов.
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];
float val = data[threadIdx.x * 1024];
__shared__ float smem[32][33];
float val = smem[threadIdx.x][threadIdx.y];При доступе к разделяемой памяти по столбцам матрицы, хранящейся по строкам, добавьте padding — один лишний элемент на строку (int stride = width + 1). Это сдвигает банки и устраняет конфликты. Для матрицы 32×32, хранящейся как smem[32][33], обращения по столбцам не создают конфликтов.
Коалесцирование — основа пропускной способности глобальной памяти. Последовательный доступ 32 потоков к 128 смежным байтам — одна транзакция. Разбросанный доступ — 32 транзакции.
Конфликты банков убивают пропускную способность shared memory. Padding на один элемент — простейший приём их устранения.
Проектируйте структуры данных и индексацию так, чтобы 32 потока варпа обращались к смежным адресам. Это правило важнее любых микрооптимизаций.
«Модель roofline — полезная визуализация, отображающая два наложенных аппаратурой потолка производительности: одну горизонтальную линию на уровне пиковой скорости операций с плавающей точкой процессора и одну наклонную линию, задаваемую пиковой пропускной способностью памяти. Вместе они образуют огибающую "roofline", которая показывает, ограничено ли данное ядро вычислениями (compute bound) или перемещением данных (memory bound)».
Roofline — график, показывающий, что ограничивает производительность ядра: вычисления или память. По горизонтальной оси — арифметическая интенсивность (FLOP на байт трафика памяти), по вертикальной — достигнутая производительность (FLOP/с).
Две линии формируют «крышу»: горизонтальная — пик вычислений, наклонная — пиковая пропускная способность памяти. Точка пересечения — точка перегиба (ridge point).
Важное уточнение: положение точки перегиба зависит от пути исполнения. Для FP32 на CUDA-ядрах на Blackwell пик составляет ~80 TFLOP/с, а пропускная способность HBM3e — ~8 ТБ/с, что даёт точку перегиба около 10 FLOP/байт. Однако для Tensor Core-путей картина принципиально иная: FP16 dense на B200 обеспечивает ~2,25 PFLOPS, что сдвигает точку перегиба на порядок-два вправо — к сотням FLOP/байт. Для FP8 она ещё выше. Это означает, что одно и то же ядро на одном и том же чипе может быть memory-bound на CUDA-ядрах и compute-bound на Tensor Cores — без каких-либо изменений в алгоритме.
Пример из главы: ядро сложения векторов читает два float (8 байт), складывает (1 FLOP), записывает один float (4 байта). Трафик: 12 байт. Арифметическая интенсивность: 1/12 ≈ 0,083 FLOP/байт. Это в 120 раз левее точки перегиба для FP32 — ядро жёстко ограничено памятью. Чтобы сдвинуть его вправо, нужно либо уменьшить байты (перейти на FP16), либо увеличить вычисления на байт.
Наклонная — память, горизонтальная — вычисления
Трансформеры находятся в разных зонах roofline в зависимости от фазы:
Roofline определяет стратегию: левее точки перегиба — оптимизируйте память, правее — вычисления.
Точка перегиба зависит от пути исполнения. Для FP32 на CUDA-ядрах Blackwell это ~10 FLOP/байт; для Tensor Core-путей (FP16/FP8/FP4) — на порядок-два выше. Одно и то же ядро может быть memory-bound на CUDA-ядрах и compute-bound на Tensor Cores.
Prefill — compute-bound, decode — memory-bound. Это объясняет, почему для LLM-инференса нужны разные стратегии оптимизации на разных фазах.
«Поскольку приложения CUDA могут порождать тысячи потоков на одно ядро, традиционная отладка может не поймать тонкие ошибки памяти и состояния гонки. NVIDIA Compute Sanitizer — набор для проверки функциональной корректности, входящий в CUDA Toolkit, — решает эти задачи, инструментируя код во время выполнения, чтобы находить ошибки на ранних этапах разработки».
Сначала корректность, потом скорость. Compute Sanitizer проверяет память, гонки, инициализацию и синхронизацию. Nsight Systems показывает временную шкалу, Nsight Compute — счётчики одного ядра.
Nsight Systems даёт картину целиком: где GPU простаивает, где происходят синхронизации, перекрываются ли копирования с вычислениями. Nsight Compute показывает детали конкретного ядра: occupancy, memory throughput, warp stall reasons. Compute Sanitizer содержит четыре инструмента: memcheck (ошибки памяти), racecheck (гонки), initcheck (неинициализированная память), synccheck (синхронизация).
Интеграция Compute Sanitizer в CI с ключом --error-exitcode позволяет ловить регрессии корректности автоматически. Это особенно ценно при работе с кастомными CUDA-ядрами, где традиционные тесты не всегда покрывают все граничные случаи.
Асинхронный запуск возвращает управление сразу
| Инструмент | Назначение | Что проверяет |
|---|---|---|
| Nsight Systems (nsys) | Временная шкала CPU и GPU | Простои GPU, барьеры, перекрытие копирований с вычислениями, NVTX-диапазоны |
| Nsight Compute (ncu) | Счётчики одного ядра | Occupancy, memory throughput, warp stall reasons, bank conflicts, register spills |
| Compute Sanitizer memcheck | Ошибки памяти | Выход за границы, misaligned access, утечки device-памяти |
| Compute Sanitizer racecheck | Гонки в shared memory | Write-after-write, write-after-read, read-after-write |
| Compute Sanitizer initcheck | Неинициализированная память | Чтение неинициализированных глобальных данных |
| Compute Sanitizer synccheck | Синхронизация | Некорректное использование барьеров, deadlocks |
Сначала корректность, потом скорость. Compute Sanitizer проверяет память, гонки, инициализацию и синхронизацию. Nsight Systems показывает временную шкалу, Nsight Compute — счётчики одного ядра.
Интегрируйте Compute Sanitizer в CI с ключом --error-exitcode, чтобы автоматически ловить регрессии корректности. Используйте NVTX-маркеры для привязки проблем к конкретным участкам кода.
Порядок работы: сначала Nsight Systems — найти узкое место в общей картине. Затем Nsight Compute — разобраться в деталях конкретного ядра. И только потом — оптимизировать.
Книга Fregly писалась в 2025 году. Ниже представлено сопоставление положений главы 6 с открытыми данными о современном состоянии архитектуры NVIDIA и практикой профилирования.
«Четыре независимых планировщика варпов позволяют до четырём варпам выдавать инструкции доступным конвейерам каждый такт».
Подтверждено. Архитектура SM с четырьмя независимыми планировщиками остаётся основой NVIDIA GPU начиная с Volta и до Blackwell. На Blackwell количество планировщиков и логика dual-issue сохранились, но FP32 и INT32 были объединены в единые CUDA-ядра, что меняет стратегию ILP для mixed-потоков.
«Высокая occupancy — достаточно варпов, чтобы покрыть задержку, — максимизирует пропускную способность и минимизирует простои».
Подтверждено. Occupancy остаётся важной метрикой, но современные профилировщики и практика показывают, что eligible warps per cycle и warp stall reasons дают более полную картину. На Blackwell с его большим регистровым файлом и L2-кэшем 126 МБ многие ядра достигают пиковой производительности при умеренной occupancy.
«Цель — сдвинуть рабочую точку ядра вправо по roofline, чтобы повысить его арифметическую интенсивность и приблизиться к состоянию compute bound».
Подтверждено. Roofline-анализ встроен в Nsight Compute и активно используется при оптимизации LLM-ядер. Точка перегиба для Blackwell сместилась из-за роста FLOPS и пропускной способности HBM, но принцип остался тем же: сначала определите, что ограничивает ядро — память или вычисления. Важно помнить, что точка перегиба зависит от пути исполнения: FP32 на CUDA-ядрах даёт ~10 FLOP/байт, Tensor Core-пути (FP16/FP8/FP4) — на порядок-два выше.
В главе упоминается dual-issue с одной арифметической и одной памятью-инструкцией, но не уточняется, что на Blackwell FP32 и INT32 объединены в единые CUDA-ядра. Это означает, что mixed INT32/FP32-потоки больше не выигрывают от dual-issue на одном ядре. Вместо этого следует использовать ILP из двух независимых FP32- или INT32-инструкций, чтобы заполнить конвейер.
В главе упоминается TMEM как «dedicated ~256 KB per-SM on-chip memory». Это соответствует действительности: на Blackwell TMEM используется инструкциями Tensor Core нового поколения (tcgen05), и это отдельный уровень иерархии памяти, который не адресуется обычными указателями. Для программиста TMEM управляется библиотеками и компилятором, но знание о его существовании помогает понять, почему Tensor Core-ядра могут работать быстрее, чем ожидается по одной лишь пропускной способности HBM. Отдельно стоит учитывать, что точка перегиба roofline для Tensor Core-путей на порядок-два выше, чем для CUDA-ядер: FP16 dense ~2,25 PFLOPS на B200 даёт точку перегиба порядка сотен FLOP/байт, для FP8 — ещё выше. Это делает большинство matmul-ядер compute-bound на Tensor Cores даже при интенсивности, которая для CUDA-ядер была бы memory-bound.
«Используйте NVIDIA Nsight Systems/Compute или PyTorch profiler для измерения этого отношения. Стремитесь улучшить это отношение, поскольку goodput фокусируется на эффективной, полезной загрузке GPU».
Подтверждено. Nsight Systems и Nsight Compute остаются стандартом де-факто для профилирования CUDA-ядер. PyTorch profiler (на базе Kineto) интегрируется с ними, а MLPerf использует их для воспроизводимых измерений. Compute Sanitizer встроен в CI-пайплайны крупных AI-компаний.
smem[32][33], обращения по столбцам не создают конфликтов.cudaMallocAsync и cudaFreeAsync избавляют от глобальной синхронизации, которая убивает перекрытие вычислений и копирований.cudaMemPrefetchAsync и cudaMemAdvise — обязательны на системах без NVLink-C2C. Даже на Grace Blackwell предпочтительно явное управление размещением.--error-exitcode ловит регрессии корректности автоматически.Глава 6 книги Fregly — систематический разбор аппаратной основы исполнения GPU-ядер. Ключевая идея: GPU — иерархическая машина, от однотактовых регистров до тысячетактной HBM. Каждое ядро живёт на этой лестнице, и его производительность определяется тем, насколько эффективно оно использует ближние уровни и избегает дальних. Occupancy, коалесцирование, конфликты банков, roofline — это не академические концепции, а инструменты, которые при правильном применении дают ускорения в десятки раз. Дополнение по архитектуре Blackwell (объединение FP32 и INT32, TMEM, NV-HBI) и уточнение о зависимости точки перегиба roofline от пути исполнения уточняют эти принципы, но не отменяют их.
Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 6 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. Документация NVIDIA CUDA C++ Programming Guide, CUDA Toolkit, Nsight Systems, Nsight Compute, Compute Sanitizer; спецификации NVIDIA Blackwell B200/GB200; MLPerf Training/Inference v5.0.
This material is a review of Chapter 6 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 about Blackwell architecture, verification of forecasts against open sources, and additional observations that go beyond the chapter's text. Quotes are provided in limited volume for review and critical analysis purposes.
Chapter 6 of Fregly's book is dedicated to the transition from high-level views of AI systems to the concrete execution model on the GPU. It examines the SIMT model, the thread and memory hierarchy, launch configuration, occupancy, the roofline model, and profiling tools. This chapter is the foundation for all subsequent chapters on CUDA kernel optimization, distributed training, and inference. The key points of the chapter are presented below, along with visualizations, code examples, and clarifications based on current data about Blackwell architecture.
«Each GPU comprises many SMs, which are roughly analogous to CPU cores but streamlined for parallelism. Each SM can track up to 64 warps (32-thread groups) on Blackwell».
«Each Blackwell GPU includes many SMs—similar to CPU cores but optimized for throughput. On modern GPUs, each SM tracks up to 64 warps (2,048 threads) concurrently. Blackwell GPUs feature 64K 32-bit registers per SM (256 KB total) and a combined 256 KB L1 cache/shared memory per SM. Up to 228 KB (227 KB usable) of that SRAM can be configured as user-managed shared memory per SM».
The Streaming Multiprocessor (SM) is the main compute unit of the GPU. Unlike a CPU core optimized for one serial task, an SM is built to keep hundreds of threads in flight simultaneously. On Blackwell, each SM holds up to 64 active warps (groups of 32 threads), totalling 2048 resident threads.
The register file is 64K 32-bit registers (256 KB), with at most 255 available per thread. The unified 256 KB SRAM block per SM is shared between L1 cache and shared memory, with up to 228 KB (227 KB usable) allocatable to shared memory. This is a key architectural decision: the developer can dynamically choose what matters more for a specific kernel — cache for global data or software-managed shared memory for thread-block communication.
«Within a Blackwell SM, multiple warp schedulers issue instructions to the available pipelines; four independent warp schedulers allow up to four warps to issue instructions to the available pipelines on every cycle. Furthermore, each scheduler supports dual-issue capable of issuing two independent instructions (e.g., one arithmetic and one memory operation) per warp».
An SM on Blackwell is split into four independent warp schedulers. Each cycle, each scheduler picks one ready warp and can issue two instructions: one arithmetic (INT32, FP32 or Tensor Core) and one memory operation (load or store). The theoretical peak is therefore four warps × two instructions = eight issues per cycle. This is only possible if the dual-issue comes from the same warp — you cannot combine instructions from different warps in one issue.
The Special Function Unit (SFU) handles transcendental operations — sine, cosine, reciprocal square root, etc. The SFU has its own pipeline, separate from INT32/FP32 and from load/store. This keeps slow transcendental math from blocking the main compute and memory pipelines and increases instruction-level parallelism.
An important clarification for the current generation: on Blackwell FP32 and INT32 are merged into unified CUDA cores, so in a given cycle a core executes either an FP32 or an INT32 instruction. Mixed INT32/FP32 streams no longer benefit from dual-issue on a single core — instead, they use warp and scheduler parallelism. This nuance changes ILP optimization strategy: instead of interleaving INT and FP in one thread, it is better to keep two independent FP instructions in flight.
Four independent schedulers, two pipelines each
| Scheduler parameter | Value on Blackwell |
|---|---|
| Number of schedulers | 4 |
| Warps per cycle (max) | 4 (one per scheduler) |
| Arithmetic instructions per cycle | 4 |
| Memory ops per cycle | 4 |
| Total issues per cycle (with dual-issue) | up to 8 |
| Max registers per thread | 255 |
| Max threads per block | 1024 |
| Max warps per SM | 64 |
The SM is a throughput-oriented processor split into four independent schedulers. Each cycle can issue up to two instructions from one warp. Peak — eight issues per cycle at full load.
Registers per thread — up to 255, threads per block — up to 1024, warps per SM — 64. These numbers determine how much work fits on one SM and what occupancy looks like.
On Blackwell, FP32 and INT32 are merged into unified CUDA cores. Mixed streams should be optimized through warp parallelism rather than interleaving instruction types in one thread.
«CUDA structures parallel work into a three-level hierarchy—threads, thread blocks (aka cooperative thread arrays [CTAs]), and grids—to balance programmability with massive throughput. At the lowest level, each thread executes your kernel code. You group threads into thread blocks of up to 1,024 threads each on modern GPUs. Thread blocks form a grid when you launch the kernel».
«Thread blocks are subdivided into warps of 32 threads that execute in lockstep under the SIMT model using a warp scheduler».
CUDA organizes parallel work in three levels: thread → thread block (CTA) → grid. A block contains up to 1024 threads; blocks form a grid launched by a single kernel call. Inside a block, threads share low-latency shared memory and synchronize via the __syncthreads() barrier. Blocks execute independently and in arbitrary order — the GPU scheduler distributes them across SMs.
Each block is divided into warps of 32 threads that execute synchronously in the SIMT model: all threads in a warp follow the same control path. If some threads take the if branch and others the else, the warp executes both branches sequentially, masking inactive lanes. This is called warp divergence. Different warps can take different branches without penalty — divergence only affects threads within one warp.
«Traditionally, threads from different thread blocks could not work with one another directly. However, modern GPU architectures and CUDA versions support thread block clusters. Thread block clusters are groups of thread blocks that can communicate with one another across SMs».
Thread block clusters are a relatively new capability introduced with Hopper and developed further in Blackwell. Within a cluster, threads from different blocks can access each other's shared memory via distributed shared memory (DSMEM). This allows building larger compute structures than a single block and synchronizing them via cluster-level barriers without resorting to global memory.
Practical meaning: for tasks like matrix multiplication with large tiles, a cluster of two or four blocks can be used to effectively increase tile size and raise arithmetic intensity without leaving the boundaries of one SM.
Three levels of CUDA parallelism
32 threads moving in lockstep
Branching inside a warp serializes execution
if/else branches, the warp executes both sequentially. Different warps diverge with no penalty.DSMEM: blocks in a cluster see each other's SMEM
Three levels of hierarchy — thread, block, grid. Inside a block, threads share shared memory and synchronize via __syncthreads(). Blocks are independent. A warp is 32 threads executing the same instructions in sync.
If threads of one warp take different if/else branches, the warp executes both sequentially, masking inactive lanes. Divergence only affects threads within one warp.
Starting with Hopper, blocks can form clusters and access each other's shared memory via DSMEM. This allows increasing the effective tile size without leaving the boundaries of one SM.
«When you invoke a kernel from the CPU (host) code, you use the <<< >>> "chevron" syntax to specify how many threads should run—and how they're organized—using two configuration parameters: blocksPerGrid for the number of thread blocks and threadsPerBlock for the number of threads within each block».
«A critical aspect of GPU performance is choosing a thread block size that aligns with the hardware's 32-thread warp size. As such, you typically pick thread block sizes that are exact multiples of 32».
A CUDA kernel launch looks like kernel<<<blocksPerGrid, threadsPerBlock>>>(...). The choice of these two parameters directly affects occupancy — the fraction of hardware thread slots occupied by active warps.
Block size should be a multiple of 32 (warp size). A block of 256 threads — eight full warps — is a universal starting point. This size offers several advantages: no partial warps, good occupancy (up to eight blocks per SM), and a balance between resources and parallelism.
For blocksPerGrid, the formula (N + threadsPerBlock - 1) / threadsPerBlock is used, which rounds up and ensures all array elements are covered. A bounds check if (idx < N) is mandatory: without it, "extra" threads of the last block will try to access memory beyond the array and trigger cudaErrorIllegalAddress.
Host launches the kernel, GPU executes it
Host → launch → GPU grid
Block size should be a multiple of 32 (warp size). A block of 256 threads — eight full warps — is a universal starting point. Reasons:
__global__ void doubleKernel(float* data, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
data[idx] *= 2.0f;
}
}
int main() {
const int N = 1'000'000;
float *h_data, *d_data;
cudaMallocHost(&h_data, N * sizeof(float));
for (int i = 0; i < N; ++i) h_data[i] = 1.0f;
cudaMalloc(&d_data, N * sizeof(float));
cudaMemcpy(d_data, h_data, N * sizeof(float),
cudaMemcpyHostToDevice);
const int threadsPerBlock = 256;
const int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
doubleKernel<<>>(d_data, N);
cudaDeviceSynchronize();
cudaFree(d_data);
cudaFreeHost(h_data);
return 0;
} The bounds check if (idx < N) is critical: without it, "extra" threads of the last block will try to access memory beyond the array and trigger cudaErrorIllegalAddress. CUDA does not throw exceptions inside kernels — the error is set in a global flag and surfaces at the next host-side synchronization call.
For two-dimensional data, a two-dimensional grid of blocks is used. For example, processing a 1024×1024 image with a 16×16 block (256 threads):
__global__ void double2DKernel(float* image, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
int idx = y * width + x;
image[idx] *= 2.0f;
}
}
int main() {
const int width = 1024;
const int height = 1024;
const int N = width * height;
float *h_image, *d_image;
cudaMallocHost(&h_image, N * sizeof(float));
cudaMalloc(&d_image, N * sizeof(float));
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid((width + 15) / 16, (height + 15) / 16);
double2DKernel<<>>(d_image, width, height);
cudaDeviceSynchronize();
cudaFree(d_image);
cudaFreeHost(h_image);
return 0;
} Block size must be a multiple of 32. 256 threads — an eight-warp block — is a universal starting point. The formula (N + threadsPerBlock - 1) / threadsPerBlock rounds up and covers all elements.
A bounds check if (idx < N) is mandatory. Without it, extra threads of the last block will trigger a memory access error.
For images and volumetric data, 2D and 3D grids are used via dim3. This naturally maps multidimensional data onto GPU threads.
«The GPU provides a multilevel memory hierarchy and helps balance capacity and speed. The hierarchy includes registers, shared memory, caches, global memory, and a specialized TMEM on Blackwell GPUs and beyond. TMEM is a dedicated ~256 KB per-SM on-chip memory used by Blackwell's 5th-generation Tensor Core instructions».
GPU memory levels differ in capacity, latency and bandwidth by orders of magnitude. Understanding this ladder is the key to writing fast kernels.
At the top are registers: single-cycle, but extremely limited (up to 255 per thread). Slightly below — shared memory and L1 cache: 20–30 cycles, up to 228 KB per SM. Further down — L2 cache: 126 MB, shared across the entire GPU, ~200 cycles. At the bottom — global HBM3e memory: hundreds and thousands of cycles, ~8 TB/s. The further from the compute core, the greater the capacity and the higher the latency.
Separately, Tensor Memory (TMEM) should be noted — a dedicated 256 KB per-SM memory used by next-generation Tensor Core instructions. TMEM is not addressed by ordinary pointers; exchange goes through tcgen05 instructions. This is another level of the hierarchy that appeared with Tensor Cores and Transformer Engine.
A pyramid: from registers to HBM
Shared memory and tensor memory on one SM
tcgen05 instructions.A software-managed cache on the SM
When registers run out — data goes to DRAM
Global memory — the main data channel
| Level | Scope | Capacity (Blackwell) | Latency | Bandwidth |
|---|---|---|---|---|
| Registers | Per thread, on SM | 64K × 32-bit (max 255 per thread) | 1 cycle | Tens of TB/s per SM |
| Shared memory / L1 | Per SM | 228 KB (227 usable) | 20–30 cycles | TB/s per SM |
| TMEM | Per SM | 256 KB | ~10 cycles | TB/s to Tensor Cores |
| Constant cache | Per SM | ~8 KB (for 64 KB __constant__) | ~1 cycle (broadcast) | Register-like |
| L2 | GPU-wide | 126 MB | ~200 cycles | Tens of TB/s |
| Global (HBM3e) | Device-wide | 180 GB (B200) / ~288 GB (B300) | Hundreds–thousands of cycles | ~8 TB/s |
The GPU memory hierarchy is registers → shared/L1 → L2 → HBM. For maximum performance, data should stay as high in this hierarchy as possible.
Register spill is one of the main causes of performance degradation. Overflowing the register file forces data into local memory, equivalent to a DRAM access. Spills must be profiled and eliminated first.
«Unified Memory (also known as CUDA Managed Memory) gives you a single, coherent address space that spans both CPU and GPU, so you no longer have to juggle separate host and device buffers or issue explicit cudaMemcpy calls».
«While accessing Unified Memory is super developer-friendly, it can cause unwanted on-demand page migrations between the CPU and the GPU. This will introduce hidden latency and execution stalls».
CUDA Managed Memory (Unified Memory) creates a single address space for CPU and GPU. You allocate via cudaMallocManaged() — and both sides see one pointer. Pages migrate on demand: if the GPU accesses data sitting in CPU RAM, a page fault occurs and data is transferred over NVLink-C2C (on Grace Blackwell) or PCIe.
The convenience is deceptive. Automatic on-demand migrations can introduce hidden latency: if a GPU thread accesses a page physically residing in CPU memory, the kernel stalls and waits for the transfer. On PCIe systems this can be orders of magnitude slower than accessing HBM. On Grace Blackwell with NVLink-C2C, transfers are faster, but still not free.
The solution is explicit placement management: cudaMemPrefetchAsync for proactive page migration to the GPU and cudaMemAdvise for driver hints about where data will be used. This turns migrations from synchronous page faults into asynchronous operations that can be overlapped with computation.
Pages migrate between CPU and GPU on demand
cudaMemPrefetchAsync — proactive migration
cudaMemAdvise: preferred location, ReadMostly
cudaMemAdvise tells the driver where to keep pages and how they will be used. Reduces the number of migrations.cudaMemPrefetchAsync(ptr, size, gpuId, stream);
cudaMemAdvise(ptr, size, cudaMemAdviseSetPreferredLocation, gpuId);
cudaMemAdvise(ptr, size, cudaMemAdviseSetReadMostly, gpuId);
cudaStreamAttachMemAsync(stream, ptr, 0, cudaMemAttachSingle);Unified Memory is convenient but requires explicit placement management. Automatic page-fault migrations can introduce hidden latency.
Use cudaMemPrefetchAsync for proactive migration and cudaMemAdvise for driver hints. This turns synchronous migrations into asynchronous ones that can be overlapped with computation.
«GPUs sustain performance by running many warps concurrently so that when one warp stalls waiting for data, another warp can run. This ability to rapidly switch between warps allows a GPU to hide memory latency. As we described earlier, the fraction of an SM's capacity actually occupied by active warps is called occupancy».
«If occupancy is low (just a few active warps), an SM may sit idle while one warp is waiting on memory. This leads to poor SM utilization».
Occupancy is the ratio of active warps per SM to the theoretical maximum (64 warps, 2048 threads on Blackwell). High occupancy lets the scheduler instantly switch to a ready warp when the current one stalls on memory. This is latency hiding.
But occupancy is not a goal in itself. If the kernel is memory-bound, pushing occupancy to 100% won't help — the bottleneck is the HBM bus, not the number of warps. More informative metrics are eligible warps per cycle and warp stall reasons: how many warps are actually ready to execute and what the rest are stalling on.
The chapter's example shows the difference between sequential and parallel vector addition. The sequential version uses one thread and loads the GPU at 1.5%. The parallel version distributes the work and loads the GPU at 95%. This illustrates that the GPU is a throughput processor and requires massive parallelism to be loaded.
The parallel version does the same work in less time
A + B = C — one thread per element
| Metric | addSequential | addParallel |
|---|---|---|
| Kernel time (ms) | 48.21 | 2.17 |
| GPU utilization | 1.5% | 95% |
| Occupancy | 1.3% | 38.7% |
| Warp efficiency | 3.1% | 100% |
To control occupancy at the compiler level, __launch_bounds__ annotations are used. They tell the compiler the maximum number of threads per block the kernel will be launched with and the minimum number of blocks that should fit per SM. The compiler limits registers per thread so these conditions are met. The runtime version is the CUDA Occupancy API:
__global__ __launch_bounds__(256, 16)
void myKernel(...) { /* ... */ }
int minGrid = 0, bestBlock = 0;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &bestBlock,
myKernel,
0,
0);Occupancy is a means to hide latency, not a goal. High occupancy helps only if the bottleneck is latency, not memory bandwidth.
Eligible warps per cycle and warp stall reasons show how many warps are actually ready to execute and what the rest are stalling on. These should guide tuning.
Use __launch_bounds__ and the CUDA Occupancy API to control occupancy. Remember that increasing occupancy at the cost of registers can lead to spills and performance degradation.
«Global memory accesses are fastest when threads in a warp access contiguous memory addresses that the hardware can combine into fewer, larger transactions. If threads access scattered or misaligned addresses, the device cannot coalesce requests into the minimal number of cache line transactions, which on modern GPUs are 128-byte lines composed of four 32-byte sectors».
When 32 threads of a warp read global memory simultaneously, the GPU tries to coalesce those accesses into as few 128-byte transactions as possible. If threads access consecutive 32-bit elements, one transaction serves the whole warp. If threads are scattered across memory, each element requires its own transaction.
A similar problem arises in shared memory, where it is called bank conflicts. Shared memory consists of 32 banks of 4 bytes each. If several threads of a warp access the same bank, the accesses serialize. A classic case is column access to a row-major matrix. The solution is to add padding: one extra element per row. For a 32×32 matrix stored as smem[32][33], column accesses cause no conflicts.
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];
float val = data[threadIdx.x * 1024];
__shared__ float smem[32][33];
float val = smem[threadIdx.x][threadIdx.y];When accessing shared memory by columns of a row-major matrix, add padding — one extra element per row (int stride = width + 1). This shifts the banks and eliminates conflicts. For a 32×32 matrix stored as smem[32][33], column accesses cause no conflicts.
Coalescing is the basis of global memory bandwidth. Sequential access of 32 threads to 128 adjacent bytes — one transaction. Scattered access — 32 transactions.
Bank conflicts kill shared memory bandwidth. One element of padding is the simplest fix.
Design data structures and indexing so that 32 threads of a warp access adjacent addresses. This rule is more important than any micro-optimization.
«A roofline model is a useful visualization that charts two hardware-imposed performance ceilings: one horizontal line at the processor's peak floating-point rate and one diagonal line set by the peak memory bandwidth. Together, these form a "roofline" envelope that reveals whether a given kernel is limited by computation (compute bound) or data movement (memory bound)».
Roofline is a plot showing what limits kernel performance: compute or memory. The horizontal axis is arithmetic intensity (FLOP per byte of memory traffic), the vertical axis is achieved performance (FLOP/s).
Two lines form the "roof": the horizontal one is peak compute, the sloped one is peak memory bandwidth. Their intersection is the ridge point.
An important clarification: the position of the ridge point depends on the execution path. For FP32 on CUDA cores on Blackwell, peak is ~80 TFLOP/s, and HBM3e bandwidth is ~8 TB/s, giving a ridge point around 10 FLOP/byte. However, for Tensor Core paths the picture is fundamentally different: FP16 dense on B200 delivers ~2.25 PFLOPS, shifting the ridge point one to two orders of magnitude to the right — to hundreds of FLOP/byte. For FP8 it is even higher. This means the same kernel on the same chip can be memory-bound on CUDA cores and compute-bound on Tensor Cores — without any algorithmic change.
The chapter's example: a vector addition kernel reads two floats (8 bytes), adds them (1 FLOP), writes one float (4 bytes). Traffic: 12 bytes. Arithmetic intensity: 1/12 ≈ 0.083 FLOP/byte. This is 120× to the left of the FP32 ridge point — the kernel is strictly memory-bound. To move it to the right, one must either reduce bytes (e.g., switch to FP16) or increase compute per byte.
The slope is memory, the ceiling is compute
Transformers live in different roofline zones depending on the phase:
Roofline determines strategy: left of the ridge point — optimize memory, right — optimize compute.
The ridge point depends on the execution path. For FP32 on Blackwell CUDA cores it is ~10 FLOP/byte; for Tensor Core paths (FP16/FP8/FP4) it is one to two orders of magnitude higher. The same kernel can be memory-bound on CUDA cores and compute-bound on Tensor Cores.
Prefill is compute-bound, decode is memory-bound. This explains why LLM inference requires different optimization strategies at different phases.
«Since CUDA applications can spawn thousands of threads per kernel, traditional debugging may fail to catch subtle memory bugs and race conditions. NVIDIA Compute Sanitizer, a functional-correctness suite included with the CUDA Toolkit, addresses these challenges by instrumenting code at runtime to find errors early in development».
Correctness first, then speed. Compute Sanitizer checks memory, races, initialization, and synchronization. Nsight Systems shows the timeline, Nsight Compute shows single-kernel counters.
Nsight Systems gives the big picture: where the GPU is idle, where synchronization happens, whether copies overlap with compute. Nsight Compute shows details of a specific kernel: occupancy, memory throughput, warp stall reasons. Compute Sanitizer contains four tools: memcheck (memory errors), racecheck (races), initcheck (uninitialized memory), synccheck (synchronization).
Integrating Compute Sanitizer into CI with the --error-exitcode flag allows catching correctness regressions automatically. This is especially valuable when working with custom CUDA kernels, where traditional tests do not always cover all edge cases.
An async launch returns control immediately
| Tool | Purpose | What it checks |
|---|---|---|
| Nsight Systems (nsys) | CPU/GPU timeline | GPU idle time, barriers, overlap of copies and compute, NVTX ranges |
| Nsight Compute (ncu) | Single-kernel counters | Occupancy, memory throughput, warp stall reasons, bank conflicts, register spills |
| Compute Sanitizer memcheck | Memory errors | Out-of-bounds, misaligned access, device memory leaks |
| Compute Sanitizer racecheck | Shared memory races | Write-after-write, write-after-read, read-after-write |
| Compute Sanitizer initcheck | Uninitialized memory | Reads of uninitialized global data |
| Compute Sanitizer synccheck | Synchronization | Incorrect barrier use, deadlocks |
Correctness first, then speed. Compute Sanitizer checks memory, races, initialization, and synchronization. Nsight Systems shows the timeline, Nsight Compute shows single-kernel counters.
Integrate Compute Sanitizer into CI with the --error-exitcode flag to catch correctness regressions automatically. Use NVTX markers to tie problems to specific code regions.
Workflow order: first Nsight Systems — find the bottleneck in the overall picture. Then Nsight Compute — understand the details of a specific kernel. Only after that — optimize.
Fregly's book was written in 2025. Below is a comparison of Chapter 6's provisions with open data about the current state of NVIDIA architecture and profiling practice.
«Four independent warp schedulers allow up to four warps to issue instructions to the available pipelines on every cycle».
Confirmed. The SM architecture with four independent warp schedulers has remained the foundation of NVIDIA GPUs from Volta through Blackwell. On Blackwell the number of schedulers and dual-issue logic are preserved, but FP32 and INT32 were merged into unified CUDA cores, which changes ILP strategy for mixed streams.
«High occupancy—enough warps to cover latency—maximizes throughput and minimizes idle stalls».
Confirmed. Occupancy remains an important metric, but modern profilers and practice show that eligible warps per cycle and warp stall reasons give a more complete picture. On Blackwell with its large register file and 126 MB L2 cache, many kernels achieve peak performance at moderate occupancy.
«The goal is to push the kernel's operational point to the right on the roofline to increase its arithmetic intensity and move closer to becoming compute bound».
Confirmed. Roofline analysis is built into Nsight Compute and is actively used in LLM kernel optimization. For Blackwell the ridge point has shifted due to increased FLOPS and HBM bandwidth, but the principle remains the same: first determine what limits the kernel — memory or compute. It is important to remember that the ridge point depends on the execution path: FP32 on CUDA cores gives ~10 FLOP/byte, Tensor Core paths (FP16/FP8/FP4) give one to two orders of magnitude higher.
The chapter mentions dual-issue with one arithmetic and one memory instruction, but does not specify that on Blackwell FP32 and INT32 are merged into unified CUDA cores. This means mixed INT32/FP32 streams no longer benefit from dual-issue on a single core. Instead, ILP from two independent FP32 or INT32 instructions should be used to fill the pipeline.
The chapter mentions TMEM as «dedicated ~256 KB per-SM on-chip memory». This is correct: on Blackwell TMEM is used by next-generation Tensor Core instructions (tcgen05), and it is a separate level of the memory hierarchy not addressed by ordinary pointers. For the programmer TMEM is managed by libraries and the compiler, but knowing about its existence helps understand why Tensor Core kernels can run faster than expected from HBM bandwidth alone. Separately, it should be noted that the roofline ridge point for Tensor Core paths is one to two orders of magnitude higher than for CUDA cores: FP16 dense ~2.25 PFLOPS on B200 gives a ridge point of hundreds of FLOP/byte, for FP8 even higher. This makes most matmul kernels compute-bound on Tensor Cores even at intensities that would be memory-bound for CUDA cores.
«Use NVIDIA Nsight Systems/Compute or PyTorch profiler to measure this ratio. Strive to improve this ratio, as goodput focuses on effective, useful GPU utilization».
Confirmed. Nsight Systems and Nsight Compute remain the de facto standard for CUDA kernel profiling. PyTorch profiler (based on Kineto) integrates with them, and MLPerf uses them for reproducible measurements. Compute Sanitizer is built into CI pipelines at large AI companies.
smem[32][33], column accesses cause no conflicts.cudaMallocAsync and cudaFreeAsync eliminate global synchronization that kills overlap of compute and copies.cudaMemPrefetchAsync and cudaMemAdvise are mandatory on systems without NVLink-C2C. Even on Grace Blackwell explicit placement management is preferable.--error-exitcode catches correctness regressions automatically.Chapter 6 of Fregly's book is a systematic breakdown of the hardware basis of GPU kernel execution. The key idea: the GPU is a hierarchical machine, from single-cycle registers to thousand-cycle HBM. Every kernel lives on this ladder, and its performance is determined by how efficiently it uses the near levels and avoids the far ones. Occupancy, coalescing, bank conflicts, roofline — these are not academic concepts but tools that, when applied correctly, yield speedups of tens of times. The addition on Blackwell architecture (FP32/INT32 merge, TMEM, NV-HBI) and the clarification on the execution-path dependence of the roofline ridge point refine these principles but do not overturn them.
Sources: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Chapter 6 — quotes provided in limited volume for review and critical analysis purposes. NVIDIA CUDA C++ Programming Guide, CUDA Toolkit, Nsight Systems, Nsight Compute, Compute Sanitizer documentation; NVIDIA Blackwell B200/GB200 specifications; MLPerf Training/Inference v5.0.
