Лекция 5. Профилирование и отладка (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть пятаяnvvpnvprofУдаленное профилированиеcuda-gdbVisual Profiler NVidia Visual Profiler, NVVP Поставляется вместе с toolkit-ом Linux/Windows timeline выполнения команд на GPU Показания счетчиков различных событий Расчёт метрик Автоматический анализ эффективностиFile -> new Session••••Путь к исполняемому файлуРабочая директорияПараметры запускаПеременные окруженияFile -> new Session -> next• Не выбирайте run analysis• Иначе программу сразувыполнят много разToolbar• Автоматический анализ• Рассчитать timeline• Собрать показания счетчиков и метрик• Выбрать интересующие счетчики иметрикиОсновное окноОсновное окноКогда и в какуюсторонукопировалиустройствоКакие ядрасколько работалипотокиОсновное окноРабота CUDA APIна хостеКонсольный выводДанные по запуску:грид, ресурсы, времявыполненияСчетчики и метрикиСчетчики и метрикиВыборустройстваПереключение междусчетчиками /метрикамиДанные повыполнению командна вкладе DetailsЗначения счетчикови метрик длязапусков ядерСчетчики и метрики Event Некоторое низкоуровневое событие, например кешпромах, банк-конфликт, бранчинг, завершениевыполнения варпа и т.д. В специальных регистрах-счетчиках намультипроцессорах накапливаются количества такихсобытий Metric Более высокоуровневая информация, рассчитываемаяна основе счетчиков, например, % кеш промахов Примеры счетчиков и метрик в следующий разПолезные метрики Метрики рассчитываются на основе значений аппаратныхсчетчиков Описания метрик - http://docs.nvidia.com/cuda/profiler-users-guide/index.html#metrics-reference Точные формулы расчета можно найти вCUDA_Profiler_User_Guide.pdfПолезные метрики achieved_occupancy branch_efficiency – доля варпов, выполненных без бранчей warp_execution_efficiency – средняя доля активных нитей вварпах inst_replay_overhead – повторы инструкций (из-засериализации) / общее число выполненных устройствоминструкций shared_replay_overhead, global_cache_replay_overhead,local_replay_overhead – разделение предыдущего попричинамПолезные метрики dram_read_throughput– объем памяти, отданной ядру /время выполненияgld_throughput - объем глобальной памяти, отданной ядру/ время выполненияgld_requested_throughput - объем глобальной памяти,запрошенный ядром / время выполненияgld_efficiency – gld_requested_throughput / gld_throughputdram_write_throughput, gst_throughput,gst_requested_throughput, gst_efficiency - аналогично длязаписиПолезные метрики l1_cache_global_hit_rate – доля кеш-попаданий приобращении к глобальной памятиl1_cache_local_hit_rate – доля кеш-попаданий приобращении к локальной памятиl2_l1_read_hit_rate –доля кеш попаданий при обращениикеша l1 к кешу l2local_memory_overhead - доля локальной памяти вобменах между l1 и l2IPC – число инструкций, запускаемых за цикл работы warpschedulerПример: перемножение матрицCommand-line profilerВключается/отключается через переменную окруженияCOMPUTE_PROFILE:$export COMPUTE_PROFILE=1$./a.outИли$COMPUTE_PROFILE=1 ./a.outФайл лога По умолчанию пишет лог профилировки в файлыcuda_profile_%d.log в текущей директории (pwd), где %dзаменяется на номер контекста Для каждого GPU свой лог профилировки cuda_profile_0.log, cuda_profile_1.log и т.д. Можно изменить через переменнуюCOMPUTE_PROFILE_LOG:$export COMPUTE_PROFILE_LOG=example.%d.log$COMPUTE_PROFILE=1 ./a.outПоддерживаемые форматы Key-Value-Pairs (по-умолчанию)method,gputime,cputime,occupancy_Z6kernelPfi,1724792.500,17.270,1.000_Z6kernelPfi,1724766.000,5.320,1.000_Z6kernelPfi,1724765.625,8.675,1.000 Comma-Separated Valuesmethod,gputime,cputime,occupancymethod=[ _Z6kernelPfi ] gputime=[ 1724745.250 ] cputime=[ 17.638 ] occupancy=[ 1.000 ]method=[ _Z6kernelPfi ] gputime=[ 1724793.750 ] cputime=[ 5.287 ] occupancy=[ 1.000 ]method=[ _Z6kernelPfi ] gputime=[ 1724740.500 ] cputime=[ 8.878 ] occupancy=[ 1.000 ]Поддерживаемые форматы Comma-Separated Values включается через переменнуюCOMPUTE_PROFILE_CSV:$export COMPUTE_PROFILE_CSV=1$COMPUTE_PROFILE=1 ./a.out Или строчкой в файле конфигурации:profilelogformat CSVСтандартные счетчики Command-line profiler по умолчанию для каждой команды,выполненной на GPU, выводит в отдельной строке: method – имя ядра | memcpyHtoD | memcpyDtoH gputime – сколько времени команда выполнялась на GPU cputime – сколько времени команда выполнялась на хостеДля неблокирующих команд – время запуска, затраченноехостом на запуск Для блокирующих – полное время, на котороезаблокировалась хост-нить occupancyФайл конфигурации Файл, в котором указаны идентификаторы интересующихопций/счетчиков, по одному в строке + строки скомментариями, начинающиеся с # Доступные опции:http://docs.nvidia.com/cuda/profiler-users-guide/index.html#commandline-profiler-options Доступные счетчики:$nvprof --query-eventsФайл конфигурации Указываем файл в переменной окруженияCOMPUTE_PROFILE_CONFIG:$ cat profiler.config#some commentgpustarttimestampstreamid$export COMPUTE_PROFILE_CONFIG=profiler.config$COMPUTE_PROFILE=1 ./a.outПример$cat profiler.configgpustarttimestampstreamidconckerneltracegridsize$export COMPUTE_PROFILE_CSV=1$export COMPUTE_PROFILE_CONFIG=profiler.config$export COMPUTE_PROFILE_LOG=example.%d.log$export COMPUTE_PROFILE=1$./a.outПример$cat example.0.log# CUDA_PROFILE_LOG_VERSION 2.0# CUDA_DEVICE 3 Tesla C2075# CUDA_CONTEXT 1# CUDA_PROFILE_CSV 1# TIMESTAMPFACTOR 13a27e35a487845egpustarttimestamp,method,gputime,cputime,gridsizeX,gridsizeY,occupancy,streamid13a71b99010f2020,memcpyHtoD,397217.531,1501.800,,,,113a71b9930a33d40,memcpyHtoD,393306.625,1338.747,,,,113a71b99483ee600,_Z6matmul14cudaPitchedPtrS_S_,12667.392,5.871,32,32,0.667,113a71b99490c0bc0,memcpyDtoH,2406.752,4200.187,,,,1• Файл example.0.log можно открыть в Visual ProfilerVisual Profiler & Command-line Лог command-line профилировщика можно открыть в NVVP,если В конфигурации указаны gpustarttimestamp и streamid - дляпостроения timeline Формат CSV Можно открыть сразу несколько логов – их timeline будутсклеены Если несколько логов получилось в результате прогонов сразличными наборами счетчиков – лучше вручнуюобъединить столбцыНесовместимость счетчиков Некоторые счетчики не могут быть собраны за один прогон Например, gld_inst_32bit и sm_cta_launched В этом случае профилировщик выберет совместимый наборсчетчиков, а в лог профилировки добавит «NV_Warning:Counter 'sm_cta_launched' is not compatible with otherselected counters and it cannot be profiled in this run» Visual Profiler автоматически делает нужное число прогонов,но с command-line profiler это приходится делать вручнуюCommand Line Profiler & метрикиCommand-line profiler не поддерживает метрики!nvprof «Бэкенд» nvvp Позволяет быстро посмотреть что-сколько считается,собрать события и метрики Данные профилировки nvprof могут бытьэкспортированы в nvvpИспользование Компиляция остается без изменений Профилировка:$nvprof [опции] [исполняемый файл] [параметры] Режимы: Summary Trace Events/metricsРежимы SummaryБыстрый способ узнать на что ушла большая частьвремени исполнения TraceТрасса вызовов CUDA API и GPU-операций Events/metricsСобытия и метрикиSummary$nvprof ./matmul 1024 1024 1024 Вывести список использованных функций тулкита, ядер,пересылок данных. Для каждой операции: Доля от общего времени выполнения, затраченное на всевыполнения операции Суммарное время всех выполнений операции Число выполнений Время выполнения (среднее, минимальное, максимальное)Пример Summary$nvprof ./matmul 1024 1024 1024Elapsed time: 5.10317CUDA error: no error==12718== NVPROF is profiling process 12718, command: ./matmul 1024 1024 1024==12718== Profiling application: ./matmul 1024 1024 1024==12718== Profiling result:Time(%)TimeCallsAvgMinMax Name58.87% 5.0759ms1 5.0759ms 5.0759ms 5.0759ms matmul(cudaPitchedPtr, cudaPitchedPtr, cudaPitchedPtr)21.73% 1.8737ms2 936.83us 923.18us 950.47us [CUDA memcpy HtoD]19.40% 1.6731ms1 1.6731ms 1.6731ms 1.6731ms [CUDA memcpy DtoH]==12718== API calls:Time(%)Time59.33% 104.54ms31.37% 55.274ms2.95% 5.2055ms2.84% 5.0000ms1.55% 2.7308ms1.29% 2.2730ms0.33% 577.77us0.17% 308.30us0.13% 221.07us…Calls31133324344Avg34.847ms55.274ms5.2055ms1.6667ms8.2250us568.24us192.59us77.076us55.267usMin153.32us55.274ms5.2055ms1.0558ms270ns549.16us161.06us64.779us52.939usMax104.23ms55.274ms5.2055ms2.7587ms425.65us579.66us252.17us90.876us59.994usNamecudaMallocPitchcudaDeviceResetcudaEventSynchronizecudaMemcpy2DcuDeviceGetAttributecudaGetDevicePropertiescudaFreecuDeviceTotalMemcuDeviceGetNameПример Summary$nvprof ./matmul 1024 1024 1024Elapsed time: 5.10317CUDA error: no error==12718== NVPROF is profiling process 12718, command: ./matmul 1024 1024 1024==12718== Profiling application: ./matmul 1024 1024 1024==12718== Profiling result:Time(%)TimeCallsAvgMinMax Name58.87% 5.0759ms1 5.0759ms 5.0759ms 5.0759ms matmul(cudaPitchedPtr, cudaPitchedPtr, cudaPitchedPtr)21.73% 1.8737ms2 936.83us 923.18us 950.47us [CUDA memcpy HtoD]19.40% 1.6731ms1 1.6731ms 1.6731ms 1.6731ms [CUDA memcpy DtoH]==12718== API calls:Time(%)Time59.33% 104.54ms31.37% 55.274ms2.95% 5.2055ms2.84% 5.0000ms1.55% 2.7308ms1.29% 2.2730ms0.33% 577.77us0.17% 308.30us0.13% 221.07us…Calls31133324344Avg34.847ms55.274ms5.2055ms1.6667ms8.2250us568.24us192.59us77.076us55.267usMin153.32us55.274ms5.2055ms1.0558ms270ns549.16us161.06us64.779us52.939usMax104.23ms55.274ms5.2055ms2.7587ms425.65us579.66us252.17us90.876us59.994usNamecudaMallocPitchcudaDeviceResetcudaEventSynchronizecudaMemcpy2DcuDeviceGetAttributecudaGetDevicePropertiescudaFreecuDeviceTotalMemcuDeviceGetNameПример Summary$nvprof ./matmul 1024 1024 1024Elapsed time: 5.10317CUDA error: no error==12718== NVPROF is profiling process 12718, command: ./matmul 1024 1024 1024==12718== Profiling application: ./matmul 1024 1024 1024==12718== Profiling result:Time(%)TimeCallsAvgMinMax Name58.87% 5.0759ms1 5.0759ms 5.0759ms 5.0759ms matmul(cudaPitchedPtr, cudaPitchedPtr, cudaPitchedPtr)21.73% 1.8737ms2 936.83us 923.18us 950.47us [CUDA memcpy HtoD]19.40% 1.6731ms1 1.6731ms 1.6731ms 1.6731ms [CUDA memcpy DtoH]==12718== API calls:Time(%)Time59.33% 104.54ms31.37% 55.274ms2.95% 5.2055ms2.84% 5.0000ms1.55% 2.7308ms1.29% 2.2730ms0.33% 577.77us0.17% 308.30us0.13% 221.07us…Calls31133324344Avg34.847ms55.274ms5.2055ms1.6667ms8.2250us568.24us192.59us77.076us55.267usMin153.32us55.274ms5.2055ms1.0558ms270ns549.16us161.06us64.779us52.939usMax104.23ms55.274ms5.2055ms2.7587ms425.65us579.66us252.17us90.876us59.994usNamecudaMallocPitchcudaDeviceResetcudaEventSynchronizecudaMemcpy2DcuDeviceGetAttributecudaGetDevicePropertiescudaFreecuDeviceTotalMemcuDeviceGetNameCPU Trace$nvprof --print-api-trace ./matmul 1024 1024 1024 Трасса вызовов CUDA API Для каждого вызова в отдельной строке время старта и времявыполнения Не только CUDA runtime API, но и CUDA driverCUDA driver API – низкоуровневое API, на котором построен CUDAruntimeCPU Trace$nvprof --print-api-trace ./matmul 1024 1024 1024==13358== NVPROF is profiling process 13358, command: ./a.out 32768 1 2 2==13358== Profiling application: ./a.out 32768 1 2 2==13358== Profiling result:Start Duration Name126.71ms 1.3140us cuDeviceGetCount126.72ms 402ns cuDeviceGet126.73ms 404ns cuDeviceGetОчень много неявных вызовов...cuDeviceGet для полученияпараметров устройства450.57ms 206.81us cudaMallocPitch450.78ms 1.2243ms cudaMemcpy2D452.01ms 1.4435ms cudaMemcpy2D453.46ms 9.1320us cudaEventCreateGPU Trace$nvprof --print-gpu-trace ./matmul 1024 1024 1024 Трасса GPU-операций – запуски ядер и копирования памяти Для каждой операции в отдельной строкеСтарт, продолжительностьДля запусков ядер:Для копирований памяти:ГридРесурсы (регистры, общая память)ОбъемПропускная способность GB/sУстройствоПотокGPU Trace$nvprof --print-gpu-trace ./matmul 1024 1024 1024==13406====13406====13406==Start289.04ms303.02ms303.03ms304.87ms318.98ms320.85msNVPROF is profiling process 13406, command: ./a.out 32768 1 2 2Profiling application: ./a.out 32768 1 2 2Profiling result:DurationGrid SizeBlock SizeRegs*SSMem*13.981ms15.931ms1.8371ms(16384 1 1)(512 1 1)160B15.973ms1.8346ms(16384 1 1)(512 1 1)160B11.025ms-Первая половина столбцовDSMem*0B0B-GPU Trace$nvprof --print-gpu-trace ./matmul 1024 1024 1024Size67.109MB67.109MB67.109MB67.109MBThroughput4.8001GB/s4.2124GB/s4.2013GB/s6.0872GB/sTeslaTeslaTeslaTeslaTeslaTeslaDeviceC2075 (2)C2075 (2)C2075 (2)C2075 (2)C2075 (2)C2075 (2)Context111111Stream131413131414Вторая половина столбцовName[CUDA memcpy HtoD][CUDA memcpy HtoD]kernel(double*, int) [363][CUDA memcpy DtoH]kernel(double*, int) [369][CUDA memcpy DtoH]События/метрики$nvprof --devices <индекс> --query-metrics$nvprof --devices <индекс> --query-events Запросить доступные для устройства метрики/события$nvprof -m <метрика>,[<метрика>] ./matmul 1024 1024 1024$nvprof -e <cобытие>,[<событие>] ./matmul 1024 1024 1024 Посчитать значения метрик/событий (среднее, максимальное,минимальное) для всех запусков каждого ядраСобытия/метрики$nvprof -m ipc,gld_efficiency,flop_sp_efficiency ./matmul 1024 1024 1024==13638== NVPROF is profiling process 13638, command: ./matmul 1024 1024 1024==13638== Warning: Some kernel(s) will be replayed on device 1 in order to collect allevents/metrics.==13638== Profiling application: ./matmul 1024 1024 1024==13638== Profiling result:==13638== Metric result:Invocations Metric NameMetric DescriptionMinMaxAvgDevice "Tesla C2075 (1)"Kernel: matmul(cudaPitchedPtr, cudaPitchedPtr, cudaPitchedPtr)1ipcExecuted IPC1.484029 1.484029 1.4840291gld_efficiencyGlobal Memory Load Efficiency 100.20%100.20%100.20%1flop_sp_efficiency FLOP Efficiency(Peak Single)5.87%5.87%5.87%Перенаправление вывода$nvprof --log-file nvprof.log ./matmul 1024 1024 1024 Перенаправить вывод в файл$nvprof --profile-api-trace none ./matmul 1024 1024 512 Не профилировать вызовы CUDA API методов на хосте$nvprof -o nvprof.trace ./matmul 1024 1024 512 Записать результаты профилирования в файл со специальнымформатом, которые можно экспортировать в nvvpУдаленное профилирование Если на локальной машине стоит тулкит: Экспорт логов Command Line Profiler / nvprof в nvvp Удаленное профилирование прямо из nvvp Если на локальной машине не стоит тулкит: Используем nvprofЭкспорт логов Получаем экспортируемые логи Для command line profilerВ конфигурации, как минимум, указаныgpustarttimestamp и streamid - для построения timeline Формат CSV Для nvvp $nvprof -o nvprof.log ./matmul 1024 1024 512 Открываем логи в nvvp File->OpenУдаленное профилирование в nvvp Добавление удаленного сервераНажать для просмотрасписка серверовНажать для добавлениянового сервераЗаполнить поляУдаленное профилирование в nvvp Выбор сервера, соединениеВыбрать серверПользователь ипарольУстановитьсоединениеУдаленное профилирование в nvvp Указание пути до тулкита на удаленном сервереРедактироватьПуть до бинарников ибиблиотек тулкитаОпределитьавтоматическиУдаленное профилирование в nvvp Исполняемый файл, параметры запускаДоступ без пароля Если нет возможности установить пароль Положить приватный ключ в$HOME\.ssh\id_rsa на Unix %UserProfile%\.ssh\id_rsa на Windows При подключении к серверу из nvvp оставить поле пароляпустым nvvp не позволяет указать приватный ключ, поэтому кладем сименем id_rsaКомпиляция и запуск Компилируем с флагами –g –G:$nvcc –arch=sm_20 -c -g -G kernel.cu -o kernel.o -g – для отладки хост-кода, -G для device-кода Не проводит оптимизации, добавляет в объектные файлыотладочные символы Запуск:$cuda-gdb ./matmulNVIDIA (R) CUDA Debugger6.5 release...(cuda-gdb) r 1024 1024 1024Запустить программус заданнымипараметрамиНовые команды в cuda-gdb сuda-gdb содержит расширенный набор команд Новые команды разделены на три группы:(cuda-gdb) cuda <команда> Переключение между нитями, ядрами, варпами и т.д., просмотртекущего местоположения (focus)(cuda-gdb) info cuda <команда> Состояние системы: устройства, работающие ядра, нити и т.д.(cuda-gdb) set cuda <команда> Различные настройки Для всех команд работает автодополнениеCправка по командам Справка по отдельной команде выводится командой help:(cuda-gdb) help cuda <команда>(cuda-gdb) help info cuda <команда>(cuda-gdb) help set cuda <команда> Если команда не указана, то выведется информация о всей группе(cuda-gdb) help set cuda(cuda-gdb) help cuda(cuda-gdb) help set cudaСтандартные команды gdb cuda-gdb поддерживает все стандартные команды gdbприменительно к CUDA-коду(cuda-gdb)(cuda-gdb)(cuda-gdb)(cuda-gdb)(cuda-gdb)(cuda-gdb)...b[reak] kernel.cu:15b[reak] simpleKernelcond[ition] 3 threadIdx.x == 1 && i < 5listn[ext]c[ontinue]Остановка в начале ядра Вручную поставить breakpoint на ядро:(cuda-gdb) break <имя ядра> Автоматически поставить метки в начала всех ядер:(cuda-gdb) set cuda break_on_launch application(cuda-gdb) set cuda break_on_launch all Справка:(cuda-gdb) help set cuda break_on_launchФокус Фокус – полное определение остановленной нитиЯдро, грид, индексы блока (blockIdx), индексы нити (threadIdx),устройство, мультипроцессор, варп, номер нити в варпе (lane) Автоматически выводится при остановке:[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,21,0),device 1, sm 1, warp 21, lane 0]Breakpoint 1, simple<<<(64,64,1),(32,32,1)>>> (A=..., B=..., C=...) atkernel.cu:1515int blockY = blockIdx.y * blockDim.y;Фокус Вывести текущий фокус: Частично:(cuda-gdb) cuda block threadblock (0,0,0), thread (0,21,0) Полностью:(cuda-gdb) cuda kernel grid block thread device sm warp lanekernel 0, grid 1, block (0,0,0), thread (0,21,0), device 1, sm 1, warp21, lane 0Переключение фокуса Переключение фокуса:(cuda-gdb) cuda <параметр> <значение>, <параметр> <значение>, ... Фокус полностью задается Физическими параметрами: устройство, мультипроцессор, варп,нить в варпе (lane)(cuda-gdb) cuda device 1 sm 3 warp 0 lane 3[Switching focus to CUDA kernel 0, grid 1, block (2,0,0), thread (3,0,0),device 1, sm 3, warp 0, lane 3]12 __global__ void simple(cudaPitchedPtr A, cudaPitchedPtr B) {Переключение фокуса Переключение фокуса:(cuda-gdb) cuda <параметр> <значение>, <параметр> <значение>, ... Фокус полностью задается Виртуальными параметрами: ядро, грид, блок, нить(cuda-gdb) cuda kernel 0 grid 1 block (2,0) thread (23,6,0)[Switching focus to CUDA kernel 0, grid 1, block (2,0,0), thread (23,6,0),device 1, sm 3, warp 6, lane 23]12 __global__ void simple(cudaPitchedPtr A, cudaPitchedPtr B) {Переключение фокуса Если указаны не все параметры, то остальным будут неявно присвоенызначения из текущего фокуса,(cuda-gdb) cuda threadthread (0,3,0)(cuda-gdb) cuda thread 2[Switching focus to CUDA kernel 0, grid 1, block (5,0,0), thread (2,3,0),device 1, sm 2, warp 3, lane 2]14int blockX = blockIdx.x * blockDim.x;Поиск ошибок использования API Если cudaGetLastError() вернул что-то отличное от cudaSuccess, то1.2.Включаем опцию остановки при ошибке APIВыводим стек вызовов в точке остановки$cuda-gdb ./matmul(cuda-gdb) set cuda api_failures stop(cuda-gdb) rStarting program:Строчка, содержащая.