Меряем время: события CUDA и пропускная способность

1. Часы для асинхронного мира: cudaEvent

Первый порыв — обложить запуск ядра системными часами. Ловушка в асинхронности (C01): хост «пробегает» строку запуска, не дожидаясь выполнения, и часы намерят время постановки задачи в очередь — микросекунды вместо настоящих миллисекунд. Можно вставить cudaDeviceSynchronize перед остановкой часов, но есть инструмент точнее — события CUDA: метки, которые устройство само проставляет, проходя очередь работы.

cudaEvent_t start, stop;
float ms = 0;

cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);            // метка «до» — в очередь устройства
kernel<<<grid, block>>>(…);
cudaEventRecord(stop);             // метка «после» — туда же

cudaEventSynchronize(stop);        // подождать, пока устройство дойдёт до метки
cudaEventElapsedTime(&ms, start, stop);   // миллисекунды между метками

printf("ядро: %.3f мс\n", ms);
cudaEventDestroy(start);
cudaEventDestroy(stop);

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

kernel<<<grid, block>>>(…);
cudaError_t err = cudaGetLastError();          // ошибка самого запуска
if (err != cudaSuccess)
    printf("CUDA: %s\n", cudaGetErrorString(err));

2. SAXPY: эталонный подопытный

SAXPY — Single-precision A·X Plus Y, операция y = a·x + y из библиотеки BLAS: вектор x умножается на скаляр a и прибавляется к вектору y. Одна строка математики, map-паттерн, ровно по два чтения и одной записи на элемент — идеальная линейка для измерения памяти. Ядро — близнец vecAdd из C04:

__global__ void saxpy(int n, float a, const float* x, float* y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        y[i] = a * x[i] + y[i];
}

// запуск на 4 миллиона элементов, замер событиями — по скелету раздела 1
const int n = 4 * 1024 * 1024;
saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, dx, dy);

Пусть события намерили t миллисекунд. Что это значит? Считаем байты. На каждый элемент ядро читает x[i] и y[i] (8 байт) и пишет y[i] (4 байта) — всего 12 байт; арифметики — две операции. Соотношение «12 байт на 2 флопа» означает, что задача упирается не в счёт, а в память, и мерить её надо байтами:

Например, n = 4 194 304 элемента, t = 0.19 мс: 12 · 4 194 304 / 0.00019 с ≈ 265 ГБ/с. Много это или мало — зависит от карты: у каждой в паспорте есть теоретическая пропускная способность памяти (например, 360 ГБ/с у GeForce RTX 3060, 320 у Tesla T4 из Colab). 265 из 360 — это 74% потолка: для одной строки кода отлично. Хорошо написанные memory-bound ядра выжимают 70–90% паспорта; если у вас 15% — ядро что-то делает не так (чаще всего — вразнобой читает память), и наоборот: ядро на 85% пика бессмысленно оптимизировать дальше — байтам физически некуда ускоряться.

3. Честный замер: а теперь с копированиями

Замер раздела 2 — время ядра. Но данные на устройство надо скопировать, а результат — скопировать обратно (C03), а PCIe даёт в лучшем случае десятки ГБ/с — на порядок меньше памяти видеокарты. Рассмотрим показательный антипример: «увеличить каждый элемент массива на единицу» — одно чтение, одна запись, одно сложение:

cudaEventRecord(start);
cudaMemcpy(dev, a, bytes, cudaMemcpyHostToDevice);   // туда
inc<<<grid, block>>>(dev, n);                        // +1 каждому
cudaMemcpy(a, dev, bytes, cudaMemcpyDeviceToHost);   // обратно
cudaEventRecord(stop);

Само ядро пролетит за доли миллисекунды, но копирования займут в десятки раз больше: каждый байт передан по PCIe дважды ради одного сложения. Процессор сделает ту же работу быстрее, вообще никуда ничего не возя. Это не приговор видеокарте — это арифметика, из которой следуют три правила рентабельности:

• Больше вычислений на каждый переданный байт. Выгодны задачи, где над скопированными данными выполняется много вычислений: матрицы (C07), фракталы (C08), обучение сетей.
• Данные живут на устройстве. Если над массивом выполняется цепочка ядер — копируйте его один раз, а не вокруг каждого ядра.
• Маленьким задачам — CPU. На сотне элементов GPU не успеет даже проснуться: константные накладные расходы запуска ядра (микросекунды) превысят всю работу.

4. Тренажёр: калькулятор пропускной способности

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

• Почему системные часы хоста не годятся для замера ядра и что использовать вместо них?
  Запуск ядра асинхронный: хост проходит строку запуска мгновенно, и часы намерят постановку в очередь, а не выполнение. Правильный инструмент — события cudaEvent: метки в очереди устройства, между которыми cudaEventElapsedTime считает миллисекунды настоящей работы.
• Что такое прогрев и почему первый запуск ядра не учитывают?
  Первый запуск включает разовые накладные расходы — инициализацию контекста, загрузку кода ядра на устройство — и может быть в разы дольше. Для честного замера ядро запускают вхолостую, а мерят последующие запуски, усредняя несколько повторов.
• Как считается эффективная пропускная способность и зачем сравнивать её с паспортной?
  Все прочитанные и записанные ядром байты делятся на время работы: для SAXPY это 12 байт на элемент. Сравнение с паспортным потолком карты превращает число в диагноз: 70–90% пика — ядро выжимает память, оптимизировать почти нечего; 10–20% — есть системная проблема, обычно в характере доступа к памяти.
• Почему «прибавить всем по единице» невыгодно на GPU, хотя это идеальный map?
  Вычислений на каждый переданный байт ничтожно мало: байт пересылается по PCIe туда и обратно ради одного сложения, а PCIe на порядок медленнее памяти видеокарты. Копирования съедают выигрыш целиком — CPU сделает быстрее вообще без пересылки данных. Параллельность задачи — необходимое условие, но не достаточное: нужна ещё плотность вычислений или данные, уже живущие на устройстве.
• Ядро «отработало за 0.003 мс». Какие две проверки сделать до радости?
  Спросить cudaGetLastError — не был ли запуск отвергнут (неверная конфигурация, превышен лимит тредов): события честно намерят пустую очередь. И проверить результат на хосте — что данные действительно посчитаны. Подозрительно малое время чаще означает «ничего не выполнилось», чем «феноменально быстро».

Источники

1. How to Implement Performance Metrics in CUDA C/C++ // NVIDIA Developer Blog : [сайт]. — URL: https://developer.nvidia.com/blog/how-implement-performance-metrics-cuda-cc/ (дата обращения: 09.07.2026).
2. Six Ways to SAXPY // NVIDIA Developer Blog : [сайт]. — URL: https://developer.nvidia.com/blog/six-ways-saxpy/ (дата обращения: 09.07.2026).
3. CUDA C++ Best Practices Guide. Performance Metrics // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/ (дата обращения: 09.07.2026).