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

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла

Лихогруд Николай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:Строчка, содержащая.

Характеристики

Тип файла PDF

PDF-формат наиболее широко используется для просмотра любого типа файлов на любом устройстве. В него можно сохранить документ, таблицы, презентацию, текст, чертежи, вычисления, графики и всё остальное, что можно показать на экране любого устройства. Именно его лучше всего использовать для печати.

Например, если Вам нужно распечатать чертёж из автокада, Вы сохраните чертёж на флешку, но будет ли автокад в пункте печати? А если будет, то нужная версия с нужными библиотеками? Именно для этого и нужен формат PDF - в нём точно будет показано верно вне зависимости от того, в какой программе создали PDF-файл и есть ли нужная программа для его просмотра.

Список файлов лекций