Память в CUDA: global, shared, constant и registers — где что использовать

Введение

Производительность CUDA-кода в значительной степени определяется тем, как эффективно используется память — у GPU есть несколько физически разных видов памяти с очень разной латентностью, объёмом и областью видимости. В статье разберём global, shared, constant память и регистры, и дадим практические рекомендации по их выбору.

Концепция

Global memory — основная память GPU (видеопамять), большая по объёму (гигабайты), но с относительно высокой латентностью доступа; доступна всем потокам всех блоков. Shared memory — быстрая память, физически расположенная на самом мультипроцессоре (SM), общая для потоков одного блока, на порядок быстрее global memory, но крайне ограничена по объёму (обычно десятки-сотни килобайт на блок). Constant memory — специальная область для данных, не меняющихся во время выполнения ядра, с аппаратным кешированием, эффективная, когда все потоки warp читают одно и то же значение. Registers — самая быстрая память, выделяется компилятором автоматически для локальных переменных потока, но крайне ограничена в количестве на поток.

Типичная оптимизация — загрузка часто используемых данных из global memory в shared memory в начале работы блока (tiling/blocking pattern), чтобы дальнейшие обращения шли к быстрой памяти, а не повторяли медленные обращения к global memory.

Пример кода

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

#define TILE_SIZE 16

// Использование shared memory для умножения матриц (tiled matrix multiplication)
__global__ void matrixMulShared(const float *A, const float *B, float *C, int n)
{
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[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 tile = 0; tile < n / TILE_SIZE; ++tile) {
        // Каждый поток загружает один элемент из global в shared memory
        tileA[threadIdx.y][threadIdx.x] = A[row * n + (tile * TILE_SIZE + threadIdx.x)];
        tileB[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * n + col];

        __syncthreads(); // ждём, пока все потоки блока загрузят свои данные

        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }

        __syncthreads(); // ждём, пока все потоки закончат использовать tile перед перезаписью
    }

    C[row * n + col] = sum;
}

// Constant memory для коэффициентов фильтра, одинаковых для всех потоков
__constant__ float filterCoefficients[9];

__global__ void applyFilter(const float *input, float *output, 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) return;

    float sum = 0.0f;
    for (int i = 0; i < 9; ++i) {
        sum += input[y * width + x] * filterCoefficients[i];
    }
    output[y * width + x] = sum;
}

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

В matrixMulShared каждый блок потоков загружает «плитку» (tile) исходных матриц A и B размером TILE_SIZE x TILE_SIZE из медленной global memory в быструю shared memory один раз, а затем все потоки блока многократно переиспользуют эти данные для вычисления частичной суммы — это резко снижает число обращений к global memory по сравнению с наивной реализацией, где каждый элемент читался бы напрямую из global memory на каждой итерации внутреннего цикла.

Вызовы __syncthreads() — это барьер синхронизации внутри блока: первый гарантирует, что все потоки завершили загрузку tile перед началом вычислений, второй — что все потоки закончили использовать текущий tile перед тем, как он будет перезаписан данными следующей итерации.

__constant__ float filterCoefficients[9] демонстрирует использование constant memory для коэффициентов свёрточного фильтра — поскольку все потоки в applyFilter читают одни и те же 9 значений, аппаратное кеширование constant memory делает эти обращения очень быстрыми.

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

Превышение лимита shared memory на блок. Если запросить больше shared memory, чем доступно на одном SM (для текущего размера блока и числа одновременно выполняемых блоков), запуск ядра завершится ошибкой или резко снизит occupancy — нужно балансировать размер tile с доступной памятью конкретного GPU.
Забытый __syncthreads() или его вызов внутри условного блока, выполняемого не всеми потоками (divergent branch) — это undefined behavior и одна из самых частых причин трудноуловимых багов в CUDA-коде: если часть потоков блока не доходит до барьера синхронизации, остальные потоки зависают в ожидании навечно.
Constant memory эффективна только при равномерном доступе (broadcast). Если разные потоки warp обращаются к разным адресам constant memory, преимущество кеширования теряется, и производительность падает до уровня обычного обращения к памяти.
Bank conflicts в shared memory. Shared memory физически разделена на банки, и если несколько потоков одного warp обращаются к разным адресам в одном банке одновременно, обращения сериализуются — для двумерных массивов в shared memory это частая проблема, решаемая добавлением «padding» к размеру второй размерности массива.