Профилирование через rcprof

rcprof/Radeon Compute Profiler – это инструмент анализа производительности, который собирает данные из среды выполнения API и графического процессора для приложений OpenCL™ и ROCm/HSA.

CodeXL использует rcprof при запуске анализа производительности (профайлинга), а также позволяет просматривать отчёты, сгенерированные rcprof.

На сервере расположен обновлённый профайлер с новыми, исправленными и более информативными параметрами. Пример запуска:

Для 32-битных программ:

rcprof.exe -o output.csv -p prog.exe

Для 64-битных программ:

rcprof-x64.exe -o output.csv -p prog.exe

Расширенный отчёт можно получить запуском:

rcprof-x64.exe --outputfile report.csv --sessionname <exe>_rcprof --perfcounter --occupancy --kerneloutput all --workingdirectory . prog.exe <program input args>

Из CodeXL можно запускать профилирование по нажатию GPU Profile. Важно дождаться конца работы программы, даже если консоль просто повиснет.

Спустя некоторое время в случае успешной отработки программы будет показан результат замеров времени и значения счетчиков.

Если же что-то пошло не так, то может появиться сообщение об ошибке, например, как показано ниже.

Если такое произошло, то сначала нужно проверить, что ваш exe действительно верно отрабатывает, а не падает с ошибками вида “Не удалось открыть входной файл” или “Файл с device кодом не найден”.

Вернемся к успешному завершению профилирования. Перед вами таблица с результатами замеров по каждому запуску кернела. Каждый запуск вынесен на отдельную строку, поэтому, если у вас был один кернел и вы запускали его один раз, то и результат у вас будет представлен на одной строке. Если же запускалось несколько кернелом и/или несколько раз кернел, то и строк будет соответствующее кол-во.

Из общих параметров можно выделить следующее:

• Название кернела.

• Время работы (в миллисекундах).

• Global worksize – размер глобальной работы (тот же параметр, что и global_work_size для clEnqueueNDRangeKernel).

• Local worksize – размер локальной работы (тот же параметр, что и local_work_size для clEnqueueNDRangeKernel).

С точки зрения производительности нас интересуют:

• ScratchRegs – кол-во приватных переменных, которые не уместились в регистрах и сохраняются в глобальную память. Если этот параметр не 0, то, скорее всего, это главная причина потери производительности.

• LDBankConflict – процент обращений к локальной памяти, приводящих к конфликтам. Большое значение параметра говорит о неэффективном использовании многобанковой архитектуры локальной памяти.

• KernelOccupancy – процент запущенных тредов на каждом исполнителе от максимально возможного. Большие значения позволяют лучше скрывать долгий отклик глобальной памяти. Достаточное значение зависит от соотношения вычислений к обращениям к глобальной памяти (чем больше вычислений на одно обращение к глобальной памяти, тем меньше требуется значение Occupancy для того, чтобы не было снижения производительности).

• VGPRs – потребление векторных регистров приватными данными кернела. Большие значения снижают Occupancy, а переполнение максимума приводит к появлению ScratchRegs.

• VALUUtilization – процент загруженности вычислительных блоков. Низкое значение связано с некратностью размера локальный группы к размеру wavefront’а и расхождением тредов в конструкциях ветвления (if-else) внутри wavefront’а.

• VALUBusy – процент времени, когда активны вычислительные блоки. Значения, далекие от 100%, говорят о большом кол-ве простоев, которые могут быть вызваны неэффективным обращением к памяти, накладными расходами на синхронизацию тредов и т.д.

• MemUnitStalled / WriteUnitStalled – процент времени, когда возникают простои, связанные с обращениями к глобальной памяти. Уменьшить значение можно, в частности, увеличением occupancy и оптимизацией порядка обращения к памяти (read/write coalescing).

Полный список счетчиков с их описанием (наведение курсора на название) можно посмотреть в File – Project Setting.

По нажатию на значение KernelOccupancy можно перейти на вкладку Kernel Occupancy Viewer. Там показаны графики заполняемости по конкретному запуску кернела.

На каждом графике фактическое использование конкретного ресурса выделено оранжевым квадратом. Если вы наведите указатель мыши на точку на графике, отобразится всплывающая подсказка, показывающая текущие значения X и Y в этом месте.

Таблица под графиками предоставляет информацию об устройстве, кернеле и загруженности кернела. В разделе Limiting factors вы можете увидеть ограничения, налагаемые каждым ресурсом.

Также можно посмотреть свой device код по нажатию на название кернела:

CL – код кернела, который непосредственно подается программе

ISA – скомпилированный код вашей программы на ассемблере для выбранного девайса.

Для представления результатов на защите можно дополнительно показывать различные параметры результатов профилирования, помимо обычных графиков. Исходная таблица с данными счетчиков лежит в файле <session_name>.csv в C:\Users\<user_name>\AppData\Roaming\CodeXL\<project_name>_ProfilerOutput\<session_name>