Оптимизация CUDA-ядер: coalesced memory access и occupancy

Введение

Написать корректное CUDA-ядро — относительно простая задача; написать эффективное ядро, использующее реальную пропускную способность GPU, требует понимания двух ключевых концепций: coalesced memory access (объединённый доступ к памяти) и occupancy (загрузка мультипроцессоров). В статье разберём оба понятия и покажем, как профилирование помогает находить узкие места.

Концепция

Coalesced access означает, что потоки одного warp обращаются к смежным адресам памяти, и аппаратное обеспечение GPU может объединить эти обращения в одну или несколько крупных транзакций памяти вместо множества мелких — это критично для производительности доступа к global memory, латентность которой высока, но пропускная способность при правильном паттерне доступа огромна. Доступ «построчно» (поток i обращается к элементу i) обычно coalesced; доступ «по столбцам» в построенной по строкам (row-major) матрице — нет, и может быть в десятки раз медленнее.

Occupancy — отношение числа активных warp на мультипроцессоре к максимально возможному числу. Низкая occupancy означает, что GPU недозагружен и не может скрывать латентность операций памяти за счёт переключения на другие готовые к выполнению warp. Occupancy ограничивается доступными ресурсами на SM: регистрами на поток, shared memory на блок, максимальным числом потоков и блоков на SM — увеличение использования одного ресурса (например, больше shared memory для лучшего переиспользования данных) может снизить другой показатель (меньше блоков влезает на SM одновременно).

Пример кода

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

#define N 1024

// НЕЭФФЕКТИВНО: каждый поток обращается к элементам через большой шаг (stride),
// что приводит к не-coalesced доступу к памяти
__global__ void transposeNaive(const float *input, float *output, int width)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < width) {
        // Чтение по строке (coalesced), запись по столбцу (НЕ coalesced)
        output[x * width + y] = input[y * width + x];
    }
}

// ЭФФЕКТИВНО: транспонирование с использованием shared memory
// для перегруппировки доступа так, чтобы и чтение, и запись были coalesced
#define TILE_DIM 32
__global__ void transposeShared(const float *input, float *output, int width)
{
    __shared__ float tile[TILE_DIM][TILE_DIM + 1]; // +1 избегает bank conflict

    int x = blockIdx.x * TILE_DIM + threadIdx.x;
    int y = blockIdx.y * TILE_DIM + threadIdx.y;

    if (x < width && y < width) {
        tile[threadIdx.y][threadIdx.x] = input[y * width + x]; // coalesced чтение
    }
    __syncthreads();

    int transposedX = blockIdx.y * TILE_DIM + threadIdx.x;
    int transposedY = blockIdx.x * TILE_DIM + threadIdx.y;

    if (transposedX < width && transposedY < width) {
        output[transposedY * width + transposedX] = tile[threadIdx.x][threadIdx.y]; // coalesced запись
    }
}

Пояснения к коду

В transposeNaive чтение input[y * width + x] для соседних потоков одного warp (различающихся по x) обращается к смежным адресам — это coalesced. Но запись output[x * width + y] для тех же соседних потоков (различающихся по x, который теперь множится на width) обращается к адресам, разнесённым на width элементов — это создаёт width отдельных транзакций памяти вместо одной объединённой, что резко снижает производительность.

transposeShared решает эту проблему: данные сначала coalesced считываются в shared memory как есть, а после барьера __syncthreads() записываются в выходной массив также coalesced, но с переставленными индексами — переупорядочивание происходит внутри быстрой shared memory, а не через медленные неэффективные обращения к global memory. Дополнительный +1 в объявлении tile[TILE_DIM][TILE_DIM + 1] — приём для избежания bank conflict при последующем чтении транспонированных данных по столбцам shared memory.

Подводные камни

Изменение размера блока без понимания влияния на occupancy. Увеличение threadsPerBlock не всегда улучшает производительность — если это приводит к превышению лимита регистров или shared memory на SM, число одновременно резидентных блоков на SM может уменьшиться, парадоксально снижая общую occupancy.
Профилирование «на глаз» без инструментов. Догадки о том, где находится узкое место (память или вычисления), часто неверны — Nsight Compute даёт точные метрики coalescing efficiency, occupancy и bottleneck analysis, и оптимизацию стоит начинать с измерения, а не с интуитивных предположений.
Излишняя оптимизация tile-размера под конкретный GPU. Параметры, оптимальные для одной архитектуры NVIDIA (например, Ampere), не обязательно оптимальны для другой (Turing, Ada Lovelace) — для кода, который должен работать на широком спектре GPU, иногда разумнее остановиться на универсально хорошем решении, чем переусложнять код под одну конкретную архитектуру.
Игнорирование register spilling. Если ядро использует слишком много локальных переменных, компилятор начинает «сбрасывать» часть из них в (медленную) local memory вместо регистров — это легко не заметить без явной проверки через флаг компиляции -Xptxas -v, который выводит статистику использования регистров и spill-памяти при компиляции.