Современные графические процессоры NVIDIA — от Hopper до Blackwell — способны выполнять десятки терафлопс в формате FP32 и ещё больше с использованием Tensor Core. Но между теоретическим пиком и реальной производительностью стоит один барьер: подсистема памяти. Данные должны поступать в вычислительные блоки достаточно быстро, чтобы процессор не простаивал. Именно здесь большинство ядер теряют производительность — не из-за нехватки арифметики, а из-за неэффективного обращения к памяти.

IПамять как узкое место

GPU располагает многоуровневой иерархией памяти, и цена доступа к каждому уровню различается на порядки. Регистры — самый быстрый уровень: задержка составляет один цикл, но их объём ограничен (до 65 536 32-битовых регистров на SM у Blackwell). Разделяемая память (shared memory) обеспечивает задержку около 30 циклов и объём до 228 КБ на SM. Кэш L2, общий для всех SM, добавляет примерно 200 циклов задержки. И, наконец, глобальная память — HBM3e на Blackwell — обеспечивает задержку около 800 циклов, но при этом обладает наибольшим объёмом: до 192 ГБ на устройство и пропускной способностью до 8 ТБ/с (16 ТБ/с в двухчиповых конфигурациях Grace Blackwell).

Ключевое наблюдение: задержка доступа к HBM примерно в 800 раз больше, чем к регистру. Если каждый warp будет ждать 800 циклов на каждую загрузку, производительность упадёт кратно. Поэтому вся оптимизация сводится к одному принципу: меньше обращаться к HBM, больше — к ближним уровням иерархии, и не давать GPU простаивать в ожидании данных.

Каждый приём, рассмотренный ниже, иллюстрируется кодом на CUDA C++ и PyTorch, а также метриками профилировщика Nsight Compute, чтобы было видно не только «что изменить», но и «почему это помогает».

IIОбъединение загрузок: почему 32 потока должны идти рядом

Обращения к глобальной памяти работают быстрее всего, когда 32 потока warp'а читают непрерывный участок памяти. Аппаратура GPU способна объединить 32 отдельных запроса в минимальное число 128-байтовых транзакций — при идеальном выравнивании это одна транзакция на warp. Каждая 128-байтовая кэш-линия состоит из четырёх 32-байтовых секторов, и оптимальный доступ укладывается ровно в четыре сектора на запрос.

Если же потоки разбегаются по разрозненным адресам — шаговая индексация, нерегулярный доступ, невыровненный старт, — аппаратура не может объединить запросы. Вместо одной широкой транзакции warp выдаёт множество мелких, каждая из которых тянет из HBM 32-байтовый сектор, но полезные байты в нём — лишь малая часть. В худшем случае число секторов на запрос может приближаться к 32.

В Nsight Compute раздел Memory Workload Analysis покажет пониженную эффективность загрузки глобальной памяти (Global Memory Load Efficiency), повышенное число чтений секторов DRAM и среднее количество секторов на запрос выше 4,0. Процент пропускной способности DRAM останется значительно ниже пикового, подтверждая, что warp тратит циклы в ожидании памяти вместо нагрузки на ALU.

Чтобы выйти из режима ограниченности памятью, можно реорганизовать данные так, чтобы каждый из 32 потоков warp'а загружал непрерывные элементы. Либо индексируйте массив как input[idx], где idx = blockIdx.x * blockDim.x + threadIdx.x, либо перейдите к компоновке «структура массивов» (SoA) вместо «массива структур» (AoS), чтобы поток i всегда обращался к элементу i.

Пример «до» (C++): необъединённый доступ с шагом

#include <cuda_runtime.h>
#include <iostream>

__global__ void uncoalescedCopy(const float* __restrict__ in,
                                float* __restrict__ out,
                                int N, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx * stride];
    }
}

Пример «до» (PyTorch)

import torch

def uncoalesced_copy(input_tensor, stride):
    flat = input_tensor.contiguous().view(-1)
    idx = torch.arange(0, flat.numel(), stride,
                       device=flat.device, dtype=torch.long)
    return torch.index_select(flat, 0, idx)

В необъединённой версии запрос памяти каждого warp'а разбивается в среднем на 8 отдельных 32-байтовых секторов. Шаблон доступа охватывает две кэш-линии на warp, раздувая количество транзакций и истощая конвейер памяти. Nsight Compute сообщает лишь около 25% от устойчивого пикового значения пропускной способности DRAM.

Пример «после» (C++): объединённый доступ

__global__ void coalescedCopy(const float* __restrict__ in,
                              float* __restrict__ out,
                              int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = in[idx];
    }
}

Объединённая версия на PyTorch — это просто out = inp.clone(), поскольку clone() на непрерывном тензоре использует векторизованное копирование памяти под капотом. Можно также применить torch.compile с mode="max-autotune", чтобы TorchInductor выбрал объединённые и векторизованные расписания.

IIIВекторизация: меньше инструкций — больше данных

Объединение — это аппаратная оптимизация времени выполнения: GPU сам «склеивает» 32 запроса по 4 байта в широкую транзакцию. Но под капотом остаются 32 отдельные инструкции загрузки, по одной на поток. Векторизация устраняет этот избыток: каждый поток читает не один float, а сразу четыре — типом float4 (16 байт).

Когда все 32 потока warp'а выдают загрузку float4, они вместе извлекают 32 × 16 = 512 байт непрерывных данных, которые контроллер памяти разбивает ровно на четыре 128-байтовые транзакции. Каждая транзакция состоит из четырёх 32-байтовых секторов, каждый полностью используемый — идеальные 4,0 сектора на транзакцию. По сравнению со скалярным случаем число инструкций загрузки на поток падает в 4 раза.

Для использования векторизованных загрузок указатели данных должны быть выровнены по ширине вектора. Функции выделения памяти CUDA, такие как cudaMalloc, возвращают указатели, выровненные как минимум по 256 байт. Однако добавление смещения элемента может нарушить выравнивание:

auto ptr4 = reinterpret_cast<const float4*>(ptr);

Пример «после» (C++): векторизованное копирование с float4

__global__ void copyVector16B(const float4* __restrict__ in,
                              float4* __restrict__ out,
                              int N4) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N4) {
        out[idx] = in[idx];
    }
}

32-байтовые векторные загрузки на Blackwell (CUDA 13)

На архитектуре Blackwell доступен ещё более широкий вектор — 32 байта на поток (8 float). CUDA не предоставляет встроенного типа float8, поэтому нужно определить собственную структуру с 32-байтовым выравниванием:

struct alignas(32) float8 { float v[8]; };

__global__ void copyVector32B(const float8* __restrict__ in,
                              float8* __restrict__ out, int N8) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N8) {
        out[idx] = in[idx];
    }
}

При 32-байтовых поточных загрузках warp перемещает 1024 байта (32 потока × 32 байта), что отображается на 8 × 128-байтовых линий. На Hopper та же операция потребует двух 16-байтовых загрузок на поток, поскольку Hopper не поддерживает 32-байтовые векторные инструкции.

IVТайлинг: загрузи один раз — используй много раз

Объединение и векторизация ускоряют саму передачу данных, но не решают проблему повторных загрузок. Наивное умножение матриц N × N загружает каждый элемент матрицы A из HBM N раз — по разу для каждой строки B, с которой он умножается. Это приводит к N–1 избыточным загрузкам на элемент.

Тайлинг решает проблему: блок потоков загружает подматрицу (тайл) 32 × 32 из A и B в разделяемую память ровно один раз, а затем все потоки блока используют эти данные для 32 × 32 = 1024 умножений-накоплений. Каждый элемент A и B теперь загружается из HBM один раз на тайл вместо 32 раз.

#define TILE_SIZE 32

__global__ void tiledMatMul(const float* A, const float* B,
                            float* C, int N) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE];
    __shared__ float sB[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < N; t += TILE_SIZE) {
        if (row < N && (t + threadIdx.x) < N)
            sA[threadIdx.y][threadIdx.x] = A[row * N + t + threadIdx.x];
        else sA[threadIdx.y][threadIdx.x] = 0.0f;

        if ((t + threadIdx.y) < N && col < N)
            sB[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
        else sB[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads();

        for (int k = 0; k < TILE_SIZE; ++k)
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];

        __syncthreads();
    }

    if (row < N && col < N) C[row * N + col] = sum;
}

Падение пропускной способности DRAM с 90% до 25% — это не ухудшение, а улучшение. Наивное ядро жадно пожирает полосу, но большая часть байтов бесполезна. Тайловое ядро потребляет меньше полосы, потому что каждый байт используется многократно. Арифметическая интенсивность возрастает с 1,5 FLOP на байт до 8 FLOP на байт, и ядро переходит из режима «ограничено памятью» в режим «ограничено вычислениями».

VКонфликты банков: невидимый тормоз разделяемой памяти

Разделяемая память на GPU разбита на 32 банка с шириной 4 байта (адреса отображаются по модулю 128 байт). Если несколько потоков warp'а обращаются к одному банку, возникает конфликт — обращения сериализуются. 32-сторонний конфликт означает, что 32 параллельных чтения превращаются в 32 последовательных, уничтожая преимущество разделяемой памяти в скорости.

Классический сценарий конфликта — транспозиция матрицы с тайлом 32 × 32. Запись в разделяемую память идёт по строкам (tile[ty][tx]) и объединена. Но чтение — по столбцам (tile[tx][ty]): при фиксированном threadIdx.x и варьируемом threadIdx.y все 32 потока попадают в один банк, потому что строки отстоят на 128 байт. Это полный 32-сторонний конфликт.

Аналогичная проблема возникает при шаге, равном числу банков:

__shared__ float arr[32 * warpCount];
float x = arr[threadIdx.x * 32];  // 32-way conflict: все в банке 0

Пример «до» (C++): наивная транспозиция

#define TILE_DIM 32

__global__ void transposeNaive(const float *idata,
                               float *odata, int width) {
    __shared__ float tile[TILE_DIM][TILE_DIM];
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    tile[threadIdx.y][threadIdx.x] = idata[y * width + x];
    __syncthreads();
    odata[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

Пример «после» (C++): транспозиция с дополнением

#define TILE_DIM 32
#define PAD 1

__global__ void transposePadded(const float *idata,
                                float *odata, int width) {
    __shared__ float tile[TILE_DIM][TILE_DIM + PAD];
    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;
    tile[threadIdx.y][threadIdx.x] = idata[y * width + x];
    __syncthreads();
    odata[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

Альтернатива дополнению — swizzling, преобразование индексов, которое «перемешивает» потоки по банкам без увеличения расхода памяти:

int bank_idx = threadIdx.x ^ (threadIdx.y & 0x1F);

Библиотека CUTLASS от NVIDIA активно использует swizzling в своих итераторах тайлов. Дополнение проще, swizzling экономнее по памяти — выбор зависит от контекста.

VIWarp shuffle: обмен через регистры, мимо памяти

Если потоки одного warp'а должны обменяться данными, разделяемая память — не единственный путь. Интринсики __shfl_sync позволяют передавать значения напрямую между регистрами потоков, минуя разделяемую память. Конфликты банков при этом невозможны — банки памяти не задействованы.

unsigned mask = __activemask();
float val = threadVal;

for (int offset = 16; offset > 0; offset >>= 1) {
    float other = __shfl_down_sync(mask, val, offset);
    val += other;
}

Каждый __shfl_down_sync напрямую читает регистр другой полосы, уменьшая число активных потоков вдвое на каждом шаге, пока полоса 0 не накопит полную сумму. За 5 итераций (16 → 8 → 4 → 2 → 1) все 32 потока получают общую сумму. Поскольку вся коммуникация остаётся в регистрах, нет трафика разделяемой памяти, нет конфликтов банков и нет лишних синхронизаций.

Ограничение: shuffle работает только внутри одного warp'а (32 потока). Для обмена между warp'ами по-прежнему нужна разделяемая или глобальная память с надлежащей синхронизацией.

VIIКэш данных только для чтения: бесплатный ускоритель

На современных GPU загрузки из глобальной памяти автоматически кэшируются в L2 и часто в L1. Но есть специализированный путь данных только для чтения, который имеет более низкую задержку для широковещательных обращений и не вытесняет другие кэшированные данные. Это не то же самое, что 64-килобайтный кэш константной памяти — кэш только для чтения значительно больше и предназначен для больших массивов неизменяемых данных, таких как векторы внедрения (embedding) в NLP-моделях.

__global__ void lookup(const float* __restrict__ table,
                       float* out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        int t = idx % T;
        out[idx] = table[t];
    }
}

Без __restrict__ компилятор не может доказать, что table не алиасит с out, и использует обычный путь глобальной памяти. С __restrict__ загрузки идут через кэш только для чтения, который поддерживает широковещание: если все 32 потока читают один и тот же адрес, значение доставляется за один цикл вместо 32 отдельных обращений.

VIIITMA и асинхронные конвейеры

Самый передовой приём — асинхронная передача данных с помощью Tensor Memory Accelerator (TMA). TMA — это аппаратный DMA-движок, который перемещает данные между глобальной памятью и разделяемой памятью без участия вычислительных блоков SM. Пока TMA загружает следующий тайл, SM вычисляет на предыдущем. TMA способна на массовые копии 1D–5D с произвольными шагами, включая объединение и многомерные передачи.

Чтобы запустить передачу TMA, используется cuda::memcpy_async() вместе с cuda::pipeline:

cuda::memcpy_async(sharedBuf,
                   globalPtr + offset,
                   cuda::aligned_size_t<16>(bytes),
                   pipe);
pipe.producer_commit();

Полный пример двойной буферизации с TMA

__shared__ float tile0[TILE_SIZE];
__shared__ float tile1[TILE_SIZE];
float* tiles[2] = { tile0, tile1 };

__shared__ cuda::pipeline_shared_state<
    cuda::thread_scope_block, 2> state;
auto pipe = cuda::make_pipeline(
    cuda::this_thread_block(), &state);

pipe.producer_acquire();
cuda::memcpy_async(tiles[0], global_ptr,
    cuda::aligned_size_t<32>{bytes}, pipe);
pipe.producer_commit();

for (int t = 1; t < nTiles; ++t) {
    pipe.consumer_wait();
    processTile(tiles[(t - 1) & 1]);
    pipe.consumer_release();

    pipe.producer_acquire();
    cuda::memcpy_async(tiles[t & 1],
        global_ptr + t * TILE_SIZE,
        cuda::aligned_size_t<32>{bytes}, pipe);
    pipe.producer_commit();
}

pipe.consumer_wait();
processTile(tiles[(nTiles - 1) & 1]);
pipe.consumer_release();

IXЧек-лист оптимизации памяти GPU

Все приёмы можно свести к нескольким принципам:

• Объединяйте загрузки. Потоки warp'а должны читать соседние адреса. Переходите от AoS к SoA, убирайте шаговые индексы. Цель — 4,0 сектора на запрос, эффективность загрузки около 100%.
• Векторизуйте. Используйте float4 (16 байт на поток) на Hopper или 32-байтовые векторы на Blackwell (CUDA 13). Проверяйте выравнивание указателей — cudaMalloc возвращает выравнивание не менее 256 байт, но смещения могут его нарушить.
• Тайлируйте. Загружайте тайлы 32 × 32 в разделяемую память один раз, используйте многократно. Подбирайте размер тайла под бюджет регистров и shared memory.
• Устраняйте конфликты банков. Дополняйте массивы (33 вместо 32 столбцов) или используйте swizzling индексов. Стоимость дополнения — около 3% памяти, выигрыш — до 3× по времени.
• Используйте shuffle. Для межпоточного обмена внутри warp'а — __shfl_sync, __shfl_down_sync, __shfl_xor_sync. Конфликты банков невозможны, синхронизация не нужна.
• Помечайте read-only данные как const __restrict__ — это включает кэш только для чтения и сокращает DRAM-трафик в 3 раза.
• Асинхронные конвейеры. cuda::memcpy_async с TMA и двойной буферизацией скрывает задержку HBM за вычислениями.
• Профилируйте. Nsight Compute — для поточных метрик (эффективность загрузки, конфликты банков, активность SM, причины простоев warp'ов). Nsight Systems — для временных шкал перекрытия и простоев на уровне системы.

XПроверка прогнозов и актуальные данные

Глава 7 книги Fregly была написана в 2025 году. С тех пор появились новые данные, которые позволяют проверить прогнозы и уточнить цифры.

XIЗаключение

Производительность GPU-ядра определяется не количеством арифметических инструкций, а тем, насколько хорошо оно использует иерархию памяти. Ядро, которое при каждом умножении обращается к HBM, будет тратить 800 циклов на ожидание — независимо от того, насколько быстр его вычислительный блок. Ядро, которое загружает тайл один раз, переиспользует его через shared memory, обменивается значениями через регистры и асинхронно предзагружает следующий тайл через TMA, превращает эти 800 циклов ожидания в 800 циклов полезной работы.

Переход от «ограничено памятью» к «ограничено вычислениями» — это и есть цель оптимизации. Когда узкое место смещается с памяти на арифметику, производительность ядра начинает приближаться к теоретическому пику GPU. Каждый из описанных приёмов — объединение, векторизация, тайлинг, устранение конфликтов банков, shuffle, кэш только для чтения, TMA-конвейеры — это один шаг на этом пути. И каждый шаг подтверждается конкретными метриками профилировщика.

Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 7 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. Документация NVIDIA CUDA C++ Programming Guide; NVIDIA Nsight Compute documentation; NVIDIA Blackwell architecture whitepaper; CUTLASS documentation; Triton documentation.