Память GPU: шесть этажей

1. Две разные памяти и мост между ними

Оперативная память компьютера принадлежит хосту. У видеокарты — своя собственная DRAM (на коробке её называют «видеопамять», 8–24 ГБ у современных карт). Процессор не умеет напрямую адресовать память видеокарты, а видеокарта — память хоста; между ними — шина PCIe, по которой данные копируются. Отсюда протокол любой CUDA-программы, который вы уже видели в схеме и теперь выучите руками:

Три функции моста:

// выделить size байт в глобальной памяти устройства
cudaMalloc((void**)&dev_ptr, size);

// скопировать size байт; последний аргумент — направление:
// cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice
cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);

// освободить память устройства
cudaFree(dev_ptr);

Все три возвращают код cudaError_t: cudaSuccess при удаче, иначе — код ошибки (неверный указатель, перепутанное направление, нехватка памяти). Привычка проверять эти коды спасает часы отладки; систематически займёмся этим в C06, а пока запомните функцию-переводчик cudaGetErrorString(err) — она превращает код в человекочитаемую строку.

2. Шесть этажей: от регистров до текстур

Теперь заглянем внутрь устройства. Память видеокарты делится на шесть типов, и ключ к пониманию — где физически находится каждый: прямо на кристалле мультипроцессора (быстро) или в DRAM за его пределами (сотни тактов латентности).

Три уточнения, на которых часто спотыкаются старые учебники:

• Разделяемая память — быстрая. Она стоит на кристалле рядом с регистрами, и обращение к ней занимает единицы тактов (если нет конфликтов банков — о них ниже). Медленной её называют только по ошибке.
• Локальная память — медленная, несмотря на уютное название: «локальная» она по видимости (личная для треда), а физически это та же DRAM. Компилятор сгружает туда переменные, когда треду не хватило регистров, и большие массивы, объявленные внутри ядра. Если профилировщик показывает «register spilling» — программа внезапно ходит в DRAM на каждой итерации.
• Глобальная память находится на видеокарте, а не «на центральном процессоре»: её латентность — плата за DRAM вне кристалла, а не за обращение к хосту. С хостом (через PCIe) обменивается данными только cudaMemcpy, и эта передача ещё на порядок дороже.

Константная память заслуживает одного абзаца практики: если все треды читают одни и те же неизменные данные — коэффициенты формулы, параметры модели, — их объявляют глобальной переменной со спецификатором __constant__ и заполняют с хоста функцией cudaMemcpyToSymbol. Чтение через константный кэш при таком сценарии почти бесплатно:

__constant__ float coeffs[64];                 // объявление на уровне файла

// на хосте: заполнить из обычного массива
cudaMemcpyToSymbol(coeffs, host_coeffs, 64 * sizeof(float));

// в ядре coeffs читается как обычный массив

Текстурная память — специализированный инструмент из графического прошлого GPU: кэш, оптимизированный под чтение «соседних» точек двумерных данных, плюс аппаратная интерполяция. В вычислительных программах её ниша узкая (обработка изображений), и в этой траектории мы ограничимся знакомством.

3. Разделяемая память: банки и конфликты

Разделяемая память объявляется внутри ядра спецификатором __shared__ и видна всем тредам блока — «общий склад бригады» из аналогии C02. Чтобы 32 треда варпа могли читать её одновременно, память нарезана на 32 банка: подряд идущие 32-битные слова лежат в подряд идущих банках. Если каждый тред варпа обращается к своему банку — все 32 обращения происходят за раз. Если два треда попали в разные слова одного банка — обращения выстраиваются в очередь: это конфликт банков, и обращение варпа замедляется во столько раз, какова глубина очереди.

__shared__ float a[256];
float x = a[threadIdx.x];        // конфликтов нет: тред k читает слово k — банк k % 32

__shared__ float b[32][32];
float y = b[threadIdx.x][0];     // конфликт 32-го порядка: столбец матрицы 32×32 —
                                 // все элементы в одном банке, варп идёт в очередь

Единственное исключение: если все треды читают одно и то же слово, конфликта нет — работает широковещательная рассылка. Стандартный рабочий цикл с разделяемой памятью выглядит так (подробно проживём его в C05):

4. Тренажёр: куда положить данные

Контрольные вопросы

• Перечислите пять шагов жизненного цикла данных в CUDA-программе.
  cudaMalloc — выделить память устройства; cudaMemcpy (HostToDevice) — скопировать вход; запуск ядра — посчитать; cudaMemcpy (DeviceToHost) — забрать результат; cudaFree — освободить. Хост и устройство не видят память друг друга напрямую, поэтому без копирований не обойтись.
• Почему «локальная» память — медленная, хотя звучит уютно?
  Локальная она только по видимости: принадлежит одному треду. Физически это область в DRAM видеокарты — туда компилятор сгружает переменные при нехватке регистров и большие массивы, объявленные в ядре. Латентность у неё как у глобальной памяти — сотни тактов.
• Чем разделяемая память отличается от глобальной по скорости и по видимости?
  Разделяемая стоит на кристалле SM: латентность в единицы тактов, близко к регистрам; видна только тредам одного блока и живёт, пока работает блок. Глобальная — в DRAM: сотни тактов; видна всем тредам грида и хосту, живёт до cudaFree.
• Когда данные стоит класть в константную память и как это делается?
  Когда все треды читают одни и те же неизменяемые данные — коэффициенты, параметры. Объявляется глобальная переменная __constant__, заполняется с хоста через cudaMemcpyToSymbol; чтение идёт через отдельный кэш и при общем чтении почти бесплатно. Из ядра эта память доступна только на чтение.
• Что такое конфликт банков и как выглядит классический пример?
  Разделяемая память разбита на 32 банка по 32-битным словам. Если треды варпа обращаются к разным словам одного банка, обращения выполняются по очереди — варп замедляется кратно глубине очереди. Классика: чтение столбца матрицы 32×32 float — весь столбец лежит в одном банке, конфликт 32-го порядка. Чтение одного и того же слова всеми тредами конфликтом не является.
• Программа передала в ядро указатель от new. Что произойдёт и почему?
  Ядро получит адрес из адресного пространства хоста, который на устройстве ничему не соответствует: чтение вернёт мусор либо ядро упадёт с illegal memory access. Компилятор ошибку не поймает — типы совпадают. В ядро можно передавать только указатели, выделенные cudaMalloc.

Источники

1. CUDA C++ Programming Guide. Memory Hierarchy // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-programming-guide/ (дата обращения: 09.07.2026).
2. CUDA C++ Best Practices Guide. Memory Optimizations // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/ (дата обращения: 09.07.2026).
3. Using Shared Memory in CUDA C/C++ // NVIDIA Developer Blog : [сайт]. — URL: https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/ (дата обращения: 09.07.2026).