Как наблюдать за событиями и метриками CUDA для подраздела исполняемого файла (например, только во время выполнения ядра)?

Question

Я знаком с использованием nvprof для доступа к событиям и метрикам бенчмарка, например,

nvprof --system-profiling on --print-gpu-trace -o (file name) --events inst_issued1 ./benchmarkname

system-profiling on --print-gpu-trace -o (filename)    

Команда дает метки времени начала, окончания ядра, мощности, температуры и сохраняет информацию в файлах nvvp, чтобы мы могли просматривать ее в визуальном профилировщике. Это позволяет нам видеть, что происходит в любом разделе кода, в частности, когда работает конкретное ядро. У меня вопрос такой-

Есть ли способ изолировать события, подсчитанные только для части выполнения теста, например, во время выполнения ядра? В приведенной выше команде

--events inst_issued1    

просто дает инструкции для всего исполняемого файла. Спасибо!

Answer 1

После более подробного изучения выясняется, что информация об уровне ядра также предоставляется для всех ядер (без использования --kernels и указав их конкретно) с помощью

nvprof --events <event names> --metrics <metric names> ./<cuda benchmark>   

Фактически, это дает вывод вида

"Устройство", "Ядро", "Вызовы", "Имя события", "Мин", "Макс","Avg"

Если ядро ​​вызывается несколько раз в тесте, это позволяет вам увидеть Min, Max, Avg желаемых событий для этих запусков ядра. Очевидно, что --kernels опция в Cuda 7.5 Profiler позволяет указывать каждый запуск каждого ядра.

Answer 2

Вы можете прочитать документацию по профилировщику.

Вы можете включать и выключать профилирование в исполняемом файле. API времени выполнения CUDA для этого:

cudaProfilerStart() 
cudaProfilerStop() 

Итак, если вы хотите собрать информацию профиля только для конкретного ядра, вы можете сделать:

#include <cuda_profiler_api.h>
...

cudaProfilerStart();
myKernel<<<...>>>(...);
cudaProfilerStop();

и выписка из документации:

При использовании функций запуска и остановки также необходимо указать инструменту профилирования отключить профилирование при запуске приложения. Для nvprof вы делаете это с флагом --profile-from-start off. Для визуального профилировщика вы используете флажок "Запустить выполнение с включенным профилированием" в представлении "Настройки".

Также из документации по nvprof в частности, вы можете ограничить табуляцию событий / метрик одним ядром с помощью переключателя командной строки:

 --kernels <kernel name>

Документация дает дополнительные возможности использования.