CUDA Unified Memory: упрощённое управление памятью CPU/GPU

Введение

Традиционная модель CUDA требует явного управления двумя отдельными областями памяти — host (CPU) и device (GPU) — с явным копированием данных между ними. Unified Memory (управляемая память) упрощает это: единый указатель, выделенный через cudaMallocManaged, доступен и из CPU-кода, и из GPU-ядер, а система (драйвер и аппаратное обеспечение, начиная с архитектуры Pascal) автоматически переносит страницы памяти между CPU и GPU по мере необходимости (page migration on demand). В статье разберём, когда Unified Memory упрощает разработку и какие у неё ограничения по производительности.

Концепция

С Unified Memory код, обращающийся к данным как с CPU, так и с GPU, не требует ручных вызовов cudaMemcpy — обращение к странице памяти, физически находящейся не на том устройстве, которое к ней обращается, вызывает автоматическую миграцию страницы (page fault, обрабатываемый драйвером). Это значительно упрощает написание прототипов и алгоритмов со сложным паттерном доступа к данным с обеих сторон, но добавляет накладные расходы на сами миграции, которые при неудачном паттерне использования могут «пинговать» одну и ту же страницу между CPU и GPU многократно (так называемый thrashing).

Для повышения производительности Unified Memory предоставляет подсказки через cudaMemPrefetchAsync() (заранее перенести данные на конкретное устройство) и cudaMemAdvise() (указать паттерн доступа, например «только чтение с нескольких устройств»).

Пример кода

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

__global__ void incrementKernel(float *data, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] += 1.0f;
    }
}

int main()
{
    const int n = 1'000'000;
    float *data;

    // Единый указатель, доступный и CPU, и GPU
    cudaMallocManaged(&data, n * sizeof(float));

    // Инициализация на CPU — обычный код, без cudaMemcpy
    for (int i = 0; i < n; ++i) {
        data[i] = static_cast<float>(i);
    }

    int deviceId;
    cudaGetDevice(&deviceId);
    // Подсказка драйверу: заранее перенести данные на GPU перед запуском ядра
    cudaMemPrefetchAsync(data, n * sizeof(float), deviceId);

    const int threadsPerBlock = 256;
    const int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    incrementKernel<<<blocksPerGrid, threadsPerBlock>>>(data, n);
    cudaDeviceSynchronize();

    // Подсказка: вернуть данные на CPU перед последующим доступом с host
    cudaMemPrefetchAsync(data, n * sizeof(float), cudaCpuDeviceId);
    cudaDeviceSynchronize();

    // Чтение результата на CPU — снова обычный код, без cudaMemcpy
    std::cout << "Первый элемент: " << data[0] << ", последний: " << data[n - 1] << std::endl;

    cudaFree(data);
    return 0;
}

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

cudaMallocManaged выделяет память, доступную из одного и того же указателя data как при обычной работе с CPU (циклы инициализации и чтения результата в примере), так и внутри CUDA-ядра incrementKernel — никакого cudaMemcpy не требуется, миграция страниц происходит автоматически «под капотом» по мере фактического обращения к данным с того или другого устройства.

Явные вызовы cudaMemPrefetchAsync перед запуском ядра и перед последующим чтением на CPU — это оптимизация: без них первое обращение GPU к каждой странице данных вызвало бы page fault и миграцию «по требованию» постранично во время выполнения ядра, что медленнее, чем заранее перенести весь блок данных одним массовым вызовом.

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

• Неявные page fault при доступе без prefetch замедляют выполнение ядра, особенно при первом обращении к большим объёмам данных — без cudaMemPrefetchAsync Unified Memory часто медленнее явного управления памятью именно из-за накладных расходов миграции по требованию вместо одной крупной операции копирования.
• Thrashing при попеременном интенсивном доступе CPU и GPU к одним и тем же страницам в коротких циклах — данные могут многократно мигрировать туда и обратно, что значительно хуже по производительности, чем единократное явное копирование в начале и в конце вычислений.
• Ограниченная поддержка на старых архитектурах GPU. Полная поддержка Unified Memory с автоматической миграцией страниц по требованию доступна начиная с архитектуры Pascal — на более старых GPU (Kepler, Maxwell) Unified Memory работает с существенными ограничениями (фактически копирование всего блока при каждом запуске ядра).
• Иллюзия «бесплатной» простоты для production-кода. Unified Memory отлично подходит для прототипирования и алгоритмов со сложным, труднопредсказуемым паттерном доступа, но для production-кода с предсказуемым потоком данных явное управление памятью (cudaMalloc/cudaMemcpyAsync с потоками) почти всегда даёт более предсказуемую и часто более высокую производительность.