Настройка GPU Shared Last Level Cache для достижения максимальной производительности

GPU Shared Last Level Cache (LLC, или L2-кеш) определяет, насколько эффективно стриминг-мультипроцессоры обмениваются данными с памятью. Неправильная конфигурация этого ресурса обнуляет прирост от мощного GPU. Задержки растут, пропускная способность падает, а параллельные воркнагрузки начинают мешать друг другу. В этой статье разобрано устройство LLC, методы настройки, инструменты профилирования и типичные ошибки конфигурации, с которыми сталкиваются разработчики ML, системные администраторы и инженеры по графике.

Что такое Shared LLC и как он устроен внутри современных GPU

Все SM в GPU обращаются к памяти через единый L2-кеш. В современных архитектурах он разбит на несколько физических секций, распределённых по кристаллу. NVIDIA A100 располагает 40 МБ L2, H100 получил 50 МБ. У AMD MI300X роль LLC выполняет огромный L3 (Infinity Cache) объёмом 256 МБ. Каждое обращение к памяти, которое не нашлось в локальных кешах конкретного SM, уходит именно в этот общий LLC. Только если и там промах, запрос уходит в HBM или GDDR.

Это делает LLC критическим узлом всей подсистемы памяти. Его пропускная способность на порядок выше, чем у внешней памяти. H100 L2 обеспечивает около 12 ТБ/с внутренней пропускной способности, тогда как HBM3 снаружи даёт около 3.35 ТБ/с. Промах в LLC в несколько раз дороже попадания по задержке и по нагрузке на шину.

При конкурентной нагрузке без явного управления LLC возникает классическая проблема cache thrashing. Несколько задач или процессов попеременно вытесняют кеш-линии друг друга. SM 0 загрузил данные, SM 1 вытеснил их своими, SM 0 вынужден снова читать из HBM. В итоге ни один процесс не получает пользы от кеша, а шина памяти работает на полную нагрузку впустую.

Физически LLC разделён на банки (cache slices), каждый из которых обслуживает определённое подмножество адресов через хеш-функцию. Горячие данные, попадающие в один банк, перегружают именно его, тогда как соседние банки простаивают. Понимание этой особенности важно при выборе размера тайлов и layout тензоров.

Как NVIDIA добавила управляемое резервирование LLC начиная с архитектуры Ampere

До Ampere управлять L2 программно было практически невозможно. Кеш работал автоматически по LRU-политике, и разработчик не мог зарезервировать его часть под конкретные данные. Единственным инструментом оставались подсказки через текстурные юниты и флаги кеширования в PTX-инструкциях, которые давали ограниченный эффект.

С архитектурой Ampere и CUDA 11.2 появился механизм cudaAccessPolicyWindow. Он позволяет указать диапазон виртуальных адресов и задать для него политику кеширования. Данные в этом диапазоне будут помечаться как persisting и вытесняться из LLC в последнюю очередь. Всё остальное получает метку streaming и вытесняется при первой возможности.

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

size_t reserveSize = min(
    (size_t)(prop.l2CacheSize * 0.75),
    prop.persistingL2CacheMaxSize
);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, reserveSize);

cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr  = dataPtr;
attr.accessPolicyWindow.num_bytes = dataSize;
attr.accessPolicyWindow.hitRatio  = 1.0f;
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;

cudaStreamSetAttribute(
    stream,
    cudaStreamAttributeAccessPolicyWindow,
    &attr
);

Параметр hitRatio управляет долей строк внутри окна, которые фактически получают метку persisting. Значение 1.0 означает, что все они остаются в кеше как можно дольше. Если размер окна превышает размер зарезервированной зоны LLC, hitRatio позволяет случайным образом занять только её часть, избегая полного вытеснения других данных.

Важно понимать, что резервирование является рекомендацией, а не жёсткой гарантией. GPU не блокирует зарезервированную зону от других кернелов. Он лишь приоритизирует eviction таким образом, чтобы persisting-данные вытеснялись последними. Это принципиально отличается от MIG, где изоляция физическая.

Как MIG и vGPU делят LLC между изолированными разделами без конкуренции за ресурс

В режиме MIG (Multi-Instance GPU, доступен с A100 и H100) LLC физически разбивается между GPU-инстансами. Каждый инстанс получает свою гарантированную долю кеша и памяти без пересечения с соседними. Это не виртуализация поверх общего ресурса, а именно физический раздел кристалла.

nvidia-smi -i 0 -mig 1

nvidia-smi mig -cgi 9,9 -C

nvidia-smi mig -lgi

Каждый созданный GPU Instance (GI) получает изолированную долю LLC. Инстанс типа 3g.40gb на A100 получает ровно половину L2 и половину HBM, и никакой другой инстанс не может вытеснить его данные из кеша. Это критически важно для облачных сред, где несколько арендаторов используют один физический GPU.

При использовании vGPU через гипервизор картина другая. Профили типа A100-40C дают полную изоляцию, аналогичную MIG. Профили типа A100-1Q делят L2 динамически без жёстких гарантий. В многопользовательских сценариях это может приводить к непредсказуемым всплескам латентности именно из-за конкуренции за LLC. Выбор профиля vGPU напрямую влияет на предсказуемость производительности.

Для контейнерных сред с несколькими CUDA-процессами на одном GPU без MIG рекомендуется явно использовать cudaAccessPolicyWindow и ограничивать persisting-зону так, чтобы суммарный резерв всех процессов не превышал 75% от общего L2. Иначе конкуренция за persisting-зону создаёт ситуацию хуже, чем полное отсутствие резервирования.

Настройка LLC для ML-инференса когда веса модели должны оставаться в кеше между запросами

При инференсе ключевые кандидаты для кеширования это веса модели. Для небольших моделей они целиком помещаются в LLC, и каждый новый запрос находит их там без обращения к HBM. При больших моделях задача сложнее: нужно выбрать, какой слой или какую часть весов зафиксировать в кеше.

Стратегия для TensorRT и PyTorch заключается в прогреве кеша перед началом инференса и в явном резервировании зоны под тензоры весов. Прогрев нужен потому, что сразу после загрузки модели LLC заполнен данными от предыдущих операций. Несколько холостых прогонов вытесняют лишнее и заполняют LLC актуальными весами.

import torch

model = MyModel().cuda().eval()
dummy = torch.randn(batch_size, seq_len, device='cuda')

with torch.no_grad():
    for _ in range(20):
        _ = model(dummy)

torch.cuda.synchronize()

После прогрева LLC содержит горячие веса. Теперь нужно убедиться, что входные данные запросов не вытесняют их. Для этого входные тензоры помечаются как streaming, а тензоры весов как persisting через CUDA API или через аннотации в кастомных CUDA-кернелах.

Для TensorRT ситуация удобнее: начиная с версии 8.6 движок автоматически использует cudaAccessPolicyWindow для весов при наличии достаточного L2. Для явного управления используется IBuilderConfig::setL2LimitForTiling.

При батчевом обучении картина другая. Градиенты и активации постоянно меняются между итерациями, и зафиксировать их в LLC не имеет смысла. Здесь важнее coalesced-доступ к памяти, правильный размер тайла матричного умножения и минимизация bank conflicts в shared memory SM. LLC при обучении работает как буфер пропускной способности, а не как хранилище рабочего набора.

Как профилировать LLC с помощью Nsight Compute и CUPTI чтобы увидеть реальный эффект настроек

Без профилирования невозможно понять, работает ли конфигурация LLC так, как ожидается. Nsight Compute и CUPTI предоставляют метрики непосредственно по L2.

ncu --metrics \
  l2_global_load_bytes,\
  l2_global_store_bytes,\
  l2_hit_rate,\
  l2_tex_read_hit_rate,\
  l2_read_throughput,\
  lts__t_sectors_op_read_hit_rate \
  --target-processes all \
  python inference.py

Ключевой показатель это l2_hit_rate. Значение ниже 60% при повторяющихся паттернах доступа сигнализирует о thrashing или неправильном размере окна persistence. Для инференса с фиксированными весами hit rate должен быть выше 80%, при хорошей настройке достигает 90% и выше.

l2_global_load_bytes показывает объём данных, запрошенных у L2. Если этот показатель резко выше, чем реально используемый объём данных, память читается нелинейно или с padding-ом, и кеш-линии используются неэффективно.

lts__t_sectors_op_read_hit_rate это более детальная метрика, доступная с Ampere. Она показывает hit rate на уровне секторов (32-байтные подстроки кеш-линий) и позволяет обнаружить частичное использование кеш-линий, которое не видно в агрегированном l2_hit_rate.

nsys profile \
  --trace=cuda,nvtx \
  --gpu-metrics-set=ga10x \
  -o report \
  python inference.py

nsys stats report.nsys-rep --report cuda_gpu_mem_time_sum

Nsight Systems добавляет временную шкалу, на которой видно, в какие моменты GPU ждёт данных из памяти, а в какие работает с данными из LLC. Длинные серые провалы на шкале SM Active при одновременно высоком Memory Throughput указывают на то, что вычисления простаивают в ожидании данных, и LLC не справляется с задачей.

Для AMD GPU аналогичный инструмент это rocprof с метриками L2CacheHit, L2CacheHitCount и MemUnitBusy. Запуск выглядит так:

rocprof --stats \
  -i metrics.txt \
  --timestamp on \
  python inference.py

В metrics.txt указываются нужные счётчики. AMD ROCm 5.x добавил поддержку rocprofv2 с более удобным форматом вывода и поддержкой MI300X.

Типичные ошибки при настройке LLC и конкретные способы их диагностики и исправления

Первая и самая распространённая ошибка это слишком большое persisting-окно. Если зарезервировать под одно окно 100% LLC, остальным стримам и кернелам кеш не достанется. При нескольких параллельных CUDA-стримах это создаёт очередь за ресурсом хуже, чем без резервирования вообще. Рабочее правило: persisting-зона не должна превышать 75% L2, если запущены несколько стримов.

Вторая ошибка это игнорирование выравнивания данных. LLC работает с кеш-линиями по 128 байт. Если тензор выровнен по 64 байта, каждое обращение захватывает половину строки и фактически тратит половину пропускной способности впустую.

t = torch.randn(1024, 1024, device='cuda')
print(t.data_ptr() % 128)
# должно быть 0

Если результат не 0, тензор нужно перевыделить с явным выравниванием через torch.empty с параметром memory_format или через кастомный аллокатор.

Третья ошибка это отсутствие сброса persisting-политики между запусками. После завершения задачи политика остаётся активной для следующего кернела в том же стриме.

cudaStreamAttrValue resetAttr = {};
resetAttr.accessPolicyWindow.num_bytes = 0;
cudaStreamSetAttribute(
    stream,
    cudaStreamAttributeAccessPolicyWindow,
    &resetAttr
);

Четвёртая ошибка это смешивание MIG и CUDA persistence API без учёта реального размера доступного L2. Внутри MIG-инстанса API резервирования работает в рамках выделенной ему доли. Если код написан с расчётом на полный LLC и запускается в инстансе с 10 МБ, окно persistence молча обрезается.

size_t available;
cudaDeviceGetLimit(&available, cudaLimitPersistingL2CacheSize);
printf("Available L2: %zu bytes\n", available);

Эту проверку нужно делать при инициализации и адаптировать размер окна динамически.

Настройка LLC на AMD GPU через ROCm и отличия от подхода NVIDIA

На AMD GPU (MI300X, RX 7900) управление кешем организовано иначе. ROCm не предоставляет прямого аналога cudaAccessPolicyWindow, но даёт доступ к флагам кеширования на уровне инструкций через GCN/CDNA ISA.

В HIP-коде управление осуществляется через атрибуты памяти при аллокации и через флаги __builtin_amdgcn_s_dcache_inv() для инвалидации кеша:

hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, 0);

void* ptr;
hipExtMallocWithFlags(&ptr, size, hipDeviceMallocUncached);

__builtin_amdgcn_s_dcache_inv();

MI300X имеет принципиально другую архитектуру LLC. Роль общего кеша здесь выполняет L3 (Infinity Cache) объёмом 256 МБ. Локальный L2 составляет всего 32 МБ и физически разбит между несколькими XCD (Accelerator Complex Die). Каждый XCD получает только 4 МБ собственного L2. Данные, к которым обращаются SM с разных XCD, передаются через Infinity Fabric к общему L3, что добавляет латентность. Это означает, что для MI300X правильное распределение данных по XCD важнее, чем в монолитных архитектурах NVIDIA.

ROCm 6.x добавил hipMemAdvise с атрибутом hipMemAdviseSetPreferredLocation для явного указания, на каком XCD должны находиться данные. Это позволяет минимизировать кросс-XCD трафик и повысить hit rate локального L2.

Как LLC влияет на производительность в играх и графических приложениях и что можно настроить без доступа к CUDA

Для игровых GPU (GeForce RTX, AMD Radeon) прямого программного управления LLC нет. Игры и графические приложения работают через драйвер, который управляет кешем автоматически. Однако есть косвенные способы влиять на эффективность LLC.

Первый способ это управление качеством текстур и mipmapping. Текстуры с правильно сгенерированными mip-уровнями обращаются к меньшему объёму данных при каждом семплировании, что повышает hit rate в LLC. Использование BC-сжатия (BC7, BC6H) сокращает объём текстурных данных в 4 раза при незначительной потере качества и соответственно увеличивает эффективный размер LLC.

Второй способ это управление разрешением рендеринга и техниками масштабирования. DLSS, FSR и XeSS рендерят кадр в меньшем разрешении, что уменьшает объём данных фреймбуфера, проходящих через LLC. Это напрямую снижает нагрузку на кеш и увеличивает hit rate для геометрических и шейдерных данных.

Третий способ это настройка параметров драйвера NVIDIA через NVCP или реестр. Параметр PreferMaxPerformance в профиле управления питанием удерживает GPU на максимальных частотах, что косвенно снижает задержку на промахе в LLC, поскольку контроллер памяти работает на полной частоте постоянно.

Для разработчиков игр на DirectX 12 и Vulkan доступны явные подсказки через D3D12_RESOURCE_FLAGS и VkMemoryPropertyFlags. Ресурсы, помеченные как UPLOAD или READBACK, размещаются в некешируемой памяти. Ресурсы без явных флагов попадают в кешируемую память и задействуют LLC. Правильная маркировка буферов позволяет драйверу эффективнее управлять приоритетами вытеснения.