Профилирование CUDA-кода в связке с Qt-приложением (Nsight Systems/Compute)

Введение

Оптимизация CUDA-кода «на глаз» — ненадёжный подход: интуиция о том, где находится узкое место, часто ошибочна. NVIDIA Nsight Systems и Nsight Compute — инструменты профилирования, которые дают точные данные о времени выполнения операций, утилизации GPU, паттернах доступа к памяти и многом другом. В статье разберём, как профилировать CUDA-код внутри Qt-приложения и на что обращать внимание в отчётах.

Концепция

Nsight Systems даёт общую картину временной шкалы выполнения приложения: когда выполняются CPU-операции, когда — GPU-ядра, копирования памяти, и как они перекрываются друг с другом и с событийным циклом Qt. Это инструмент верхнего уровня — помогает понять, простаивает ли GPU в ожидании CPU, есть ли излишняя сериализация операций, эффективно ли используются потоки CUDA.

Nsight Compute даёт детальный анализ отдельного запуска ядра: occupancy, эффективность использования памяти (coalescing), узкие места (memory-bound vs compute-bound), использование регистров и shared memory. Это инструмент для глубокой оптимизации конкретного «горячего» ядра, уже выявленного через Nsight Systems или собственные измерения.

Для Qt-приложений важно учитывать, что профилирование захватывает весь процесс, включая накладные расходы самого Qt event loop, рендеринга виджетов и так далее — нужно уметь визуально отделять «шум» GUI-фреймворка от интересующих CUDA-операций на временной шкале.

Пример кода

#include <cuda_runtime.h>
#include <nvtx3/nvToolsExt.h>
#include <QElapsedTimer>
#include <QDebug>

// Использование NVTX-меток для аннотирования участков кода в Nsight Systems
void processImageWithProfiling(unsigned char *deviceBuffer, int width, int height)
{
    nvtxRangePushA("ImageProcessing::Upload"); // именованный диапазон в профиле
    // ... cudaMemcpy на устройство ...
    nvtxRangePop();

    nvtxRangePushA("ImageProcessing::KernelExecution");
    dim3 blockSize(16, 16);
    dim3 gridSize((width + 15) / 16, (height + 15) / 16);
    // invertColorsKernel<<<gridSize, blockSize>>>(deviceBuffer, width, height);
    cudaDeviceSynchronize();
    nvtxRangePop();

    nvtxRangePushA("ImageProcessing::Download");
    // ... cudaMemcpy обратно на host ...
    nvtxRangePop();
}

// Простое измерение времени средствами Qt для предварительной оценки
void measureWithQtTimer()
{
    QElapsedTimer timer;
    timer.start();

    // ... вызов CUDA-операций ...
    cudaDeviceSynchronize(); // важно дождаться завершения перед измерением

    qDebug() << "Время выполнения GPU-операций:" << timer.elapsed() << "мс";
}

# Запуск Nsight Systems для профилирования Qt-приложения с CUDA
nsys profile --output=app_profile --trace=cuda,nvtx,osrt ./MyQtCudaApp

# Запуск Nsight Compute для детального анализа конкретного ядра
ncu --kernel-name invertColorsKernel --set full -o kernel_profile ./MyQtCudaApp

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

Функции nvtxRangePushA/nvtxRangePop из библиотеки NVTX (NVIDIA Tools Extension) позволяют размечать произвольные участки кода именованными диапазонами, которые отображаются на временной шкале Nsight Systems — это особенно полезно для разделения логических этапов обработки (загрузка, вычисление, выгрузка) визуально на графике профилировщика, а не только видеть голые вызовы CUDA API.

measureWithQtTimer показывает более простой, «ручной» способ измерения времени через QElapsedTimer — пригоден для быстрой предварительной оценки без полноценного профилировщика, но требует обязательного cudaDeviceSynchronize() перед остановкой таймера, поскольку без синхронизации измеряется только время постановки асинхронных операций в очередь, а не их реального выполнения на GPU.

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

Измерение времени без cudaDeviceSynchronize() перед остановкой таймера — частая ошибка: поскольку запуск ядра и cudaMemcpyAsync асинхронны, время, «измеренное» таким способом, отражает лишь накладные расходы на постановку команды в очередь, а не реальное время выполнения на GPU.
Профилирование Debug-сборки вместо Release. Debug-сборки CUDA-кода компилируются без оптимизаций (-O0, дополнительные проверки), и их профиль производительности может радикально отличаться от Release-сборки — выводы, сделанные на Debug-профиле, часто неприменимы к реальной производительности.
Игнорирование «холодного старта» GPU. Первый запуск CUDA-кода в приложении включает время инициализации контекста CUDA-драйвера, JIT-компиляции PTX-кода и прочие одноразовые накладные расходы — при замере производительности их стоит исключать из усреднённых показателей, выполняя несколько «прогревочных» итераций перед измерением.
Несовместимость версии Nsight с версией драйвера/CUDA Toolkit. Слишком старая версия Nsight Systems/Compute может не поддерживать профилирование новейших архитектур GPU или новых возможностей CUDA Toolkit — стоит сверять совместимость версий перед началом серьёзной оптимизации.