Лекция. CUDA 3 (Колганов) (Электронные лекции)

Колганов Александрalexander.k.s@mail.ruчасть 3nvvpnvprofcuda-gdbУдаленное профилированиеVisual Profiler NVidia Visual Profiler, NVVP Поставляется вместе с toolkit-ом Linux/Windows timeline выполнения команд на GPU Показания счетчиков различных событий Расчёт метрик Автоматический анализ эффективностиFile -> new Session••••Путь к исполняемому файлуРабочая директорияПараметры запускаПеременные окруженияFile -> new Session -> next• Не выбирайте run analysis• Иначе программу сразувыполнят много разToolbar• Автоматический анализ• Рассчитать timeline• Собрать показания счетчиков и метрик• Выбрать интересующие счетчики иметрикиОсновное окноОсновное окноКогда и в какуюсторонукопировалиустройствоКакие ядрасколько работалипотокиОсновное окноРабота CUDAAPI на хостеКонсольный выводДанные по запуску:грид, ресурсы, времявыполненияСчетчики и метрикиСчетчики и метрикиВыборустройстваПереключение междусчетчиками /метрикамиДанные повыполнению командна вкладе 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 – число инструкций, запускаемых за цикл работы warpschedulerCommand-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Стандартные счетчики 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Несовместимость счетчиков Некоторые счетчики не могут быть собраны за один прогон Например, 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 это приходится делать вручнуюПоддерживаемые форматы 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 CSVVisual Profiler & Command-line Лог command-line профилировщика можно открыть в NVVP,если В конфигурации указаны gpustarttimestamp и streamid - дляпостроения timeline Формат CSV Можно открыть сразу несколько логов – их timeline будутсклеены Если несколько логов получилось в результате прогонов сразличными наборами счетчиков – лучше вручнуюобъединить столбцыПример$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 ProfilerCommand 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 driver API, но и CUDA runtimeCPU 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 cuDeviceGet126.74ms 285ns cuDeviceGetGPU Trace$nvprof --print-gpu-trace ./matmul 1024 1024 1024 Трасса GPU-операций – запуски ядер и копирования памяти Для каждой операции в отдельной строкеСтарт, продолжительностьДля запусков ядер:Для копирований памяти:ГридРесурсы (регистры, общая память)ОбъемСкорость GB/sУстройствоПотокGPU Trace$nvprof --print-gpu-trace ./matmul 1024 1024 1024==13406== NVPROF is profiling process 13406, command: ./a.out 32768 1 2 2==13406== Profiling application: ./a.out 32768 1 2 2==13406== Profiling result:Start DurationGrid SizeBlock SizeRegs*SSMem*DSMem*Size ThroughputDeviceContextStream Name289.04ms 13.981ms67.109MB 4.8001GB/s Tesla C2075 (2)113 [CUDA memcpy HtoD]303.02ms 15.931ms67.109MB 4.2124GB/s Tesla C2075 (2)114 [CUDA memcpy HtoD]303.03ms 1.8371ms(16384 1 1)(512 1 1)160B0B- Tesla C2075 (2)113 kernel(double*, int) [363]304.87ms 15.973ms67.109MB 4.2013GB/s Tesla C2075 (2)113 [CUDA memcpy DtoH]318.98ms 1.8346ms(16384 1 1)(512 1 1)160B0B- Tesla C2075 (2)114 kernel(double*, int) [369]320.85ms 11.025ms67.109MB 6.0872GB/s Tesla C2075 (2)114 [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:InvocationsMetric NameMetricDescriptionMinMaxAvgDevice "Tesla C2075 (1)"Kernel: matmul(cudaPitchedPtr, cudaPitchedPtr, cudaPitchedPtr)1ipcExecuted IPC1.4840291.4840291.4840291gld_efficiencyGlobal Memory LoadEfficiency100.20%100.20%100.20%1flop_sp_efficiencyFLOP Efficiency(PeakSingle)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Когда их стоит использовать?cudaStream Последовательность команд для GPU (запуски ядер,копирования памяти и т.д.), исполняемая строгопоследовательно следующая команда выполняется после полногозавершения предыдущейcudaStream Пользователь сам создает потоки и распределяеткоманды по ним По-умолчанию, все команды помещаются в «DefaultStream», равный нулюcudaStream Только команды из разных потоков, отличных от потока по-умолчанию, могут выполняться параллельно Пользователь сам задет необходимую синхронизациюмежду командами из разных потоков (при наличиизависимостей) В общем случае, порядок выполнения команд из разныхпотоков не определенСоздание и уничтожениеcudaStream_t stream;cudaStreamCreate(&stream);…cudaStreamDestroy(stream); Поток привязывается к текущему активному устройству Перед отправлением команды нужно переключаться наустройство, к которому привязан поток Если попробовать отправить в него команду при другомактивном устройстве, будет ошибкаАсинхронное копированиеКогда возвращается управление хостовой нитиHost->devicepageablememcpymemcpyAsyncpinnedDevice->hostpageableПослекопирования вбуфер*ПослеПослеполногополногозавершензавершенияияПослекопирования вбуферПослеполногозавершениясразуpinnedПослеполногозавершениясразуhost-hostdev-devПослеполногозавершенияПослеполногозавершения*В начале работы неявно вызывается cudaDeviceSynchronizeсразусразуПараллельная работа хоста и устройства Ядра выполняются асинхронно Копирование между pinned-памятью и памятью устройствапри помощи cudaMemcpyAsync также выполняетсяасинхронно=> Добиться параллельной работы хоста и устройствадостаточно просто!Параллельная работа хоста и устройстваПример:cudaMallocHost(&aHost, size);cudaMemcpyAsync( aDev, aHost, size,cudaMemcpyHostToDevice);kernel<<<grid, block>>>(aDev, bDev);cudaMemcpyAsync( bHost,cudaMemcpyHostToDevice);doSomeWorkOnHost();doSomeWorkOnHost будет выполняться параллельно скопированиями и выполнением ядраПараллельное выполнение команд на GPU Команды из разных потоков, отличных от потока по-умолчанию, могут исполняться параллельно В зависимости от аппаратных возможностей Возможные случаи: Параллельные копирование и выполнение ядра Параллельные выполнение ядер Параллельные копирования с хоста на устройство и сустройства на хостКопирование & выполнение ядра Если cudaDeviceProp::asyncEngineCount > 0устройство может выполнять параллельнокопирование и счет ядра Хостовая память долна быть page-lockedcudaMallocHost(&aHost, size);cudaStreamCreate(&stream1);cudaStreamCreate(&stream2); cudaMemcpyAsync(aDev, aHost, size,cudaMemcpyHostToDevice, stream1);kernel<<<grid, block, 0, stream2>>>(…);Параллельное выполнение ядер Если cudaDeviceProp::concurrentKernels > 0устройство может выполнять ядра параллельноcudaStreamCreate(&stream1);cudaStreamCreate(&stream2); kernel1<<<grid,block, 0,stream1>>>(data_1);kernel2<<<grid, block, 0,stream2>>>(data_2);Копирование в обе стороны & выполнение ядра Если cudaDeviceProp::asyncEngineCount == 2устройство может выполнять параллельно копирование вобе стороны и счет ядраcudaMallocHost(&aHost, size);cudaMallocHost(&bHost, size);// создать потокиcudaMemcpyAsync( aDev, aHost, size,cudaMemcpyHostToDevice, stream1);cudaMemcpyAsync( bHost, bDev, size,cudaMemcpyDeviceToHost, stream2);kernel<<<grid, block, 0, stream3>>>(…);Обозначения Ядро = код для GPU, __global void kernel(…) {} Размер ядра = длительность вычисления отдельной нитью Большое, сложное ядро = ядро большого размера Выполнение ядра = выполнение ядра на некотором гриде Запуск ядра = команда запуска ядра на гриде Длительный запуск ядра = запуск ядра, который долговыполняется Время вычисления ~ размер ядра * размер гридаПримеры Рассмотрим классическую схему работы с GPU: Копирование входных данных на GPU Выполнение ядра Копирование результатов обратно на хост Необходима синхронизация между командами, т.е.следующая выполняется после завершенияпредыдущейHtoDKernelDtoHИдеальный случай Выполнения копирований сопоставимы по времени свыполнением ядраHtoDKernelDtoH Разобьем задачу на подзадачи Разделим грид на части и запустим то же ядро наподгридах – результат не изменится Подзадаче нужна только часть данных для старта Запустим подзадачи в разных потокахИдеальный случай Запустим подзадачи в разных потокахHtoDStream 0Stream 1Stream 2Stream 3KernelHtoDDtoHKerDtoHHtoDKerDtoHHtoDKerDtoHHtoDKerВремяDtoHИдеальный случай Выполнение первой подзадачи начнется сразу послекопирования нужной ей части данных Выполнение ядер будет происходить параллельно скопированиями Параллельное копирование в обе стороны, еслиаппаратура позволяетМакимальное ускорение:321+число потоковНебольшие копированияHtoDKernelПолучится ли?KernelKernelKernelKernelDtoHНебольшие копированияHtoDKernelПолучится ли?KernelKernelKernelKernelНет!Суммарное время выполнения ядра наподгридах никак не сделать меньше временивыполнения на целом гридеDtoHПараллельное выполнение ядер Ресурсы GPU ограничены – до 16 SM, до 1536 нитей наодном SM, до 8-ми блоков на SM Ядра запускаются на гридах из миллионов нитей = тысячиблоков Блоки всех гридов попадают в одну общую очередь ивыполняются по мере освобождениямультипроцессоровМультипроцессоры – «bottle neck» этой очередиОчередь блоковХвостпервого ядраЕдинственное место, гдеядра могут выполнятьсяпараллельноВолны блоковKernel 1ГридыKernel 0SM 0SM 1SM 2А если синхронно?ХвостПервого запускаядраПервое ядро еще незавершилось – следующеене может начатьсяKernel 1Kernel 0Волны блоковГридыSM 0SM 1SM 2Пример в NVVP При большом приближении видно хвостыВывод При запуске двух ядер в разных потоках они могутвыполняться параллельно только на границах Когда последняя волна блоков (хвост) первого иззапущенных ядер не полностью загружает устройство Суммарное время выполнения одинаковых посложности ядер, запущенных в разных потоках:суммарное число блоков ∗ время выполнения блока1536число мультипроцессоров ∗ (∗ )размер блокаВывод Суммарное время отдельных выполнений ядра наподгридах в разных потоках равно времени выполнения нацелом гриде Суммарное число блоков не меняется Суммарное время отдельных выполнений ядра наподгридах в одном потоке >= времени выполнения нацелом гриде Из-за простоя на границахНебольшие копированияHtoDKernelDtoHKernelKernelKernelKernelВремяВывод При разбиении задачи на подзадачи выигрыш можемполучить только за счет Параллельного выполнения ядер и копирований Параллельного выполнения копированийПри небольших копированияхне имеет смысла возиться с потоками для подзадачМножество задачHtoDKerDtoH HtoDKerDtoH HtoDKerDtoH HtoDKer При распределении по потокам результат зависит от Соотношения копирования / счет Размера гридовDtoHЯдра запускаются на больших гридахHtoDKerDtoH HtoDKerHtoDKerDtoHHtoDKerDtoHHtoDKerDtoHHtoDKerDtoH HtoDKerDtoH HtoDKerDtoHDtoH Гриды большие => хвосты запусков ядер занимают малуюдолю в общем число блоков => доля параллельноговыполнения невелика Выигрываем в основном за счет параллельныхкопированийМало копирований – мало ускорения!Ядра запускаются на малых гридах Пусть ядра запускаются на гридах такого размера, чторесурсов хватает для размещения всех блоков несколькихгридов Например двух Пусть ядра примерно одинаковой сложностиТогда два ядра, запущенные таких гридах, будутвыполняться параллельно!Ядра запускаются на малых гридах Ядра запускаются на трех блоках по 1024 нити,устройство состоит из 6 SMОчередь блоковДва ядра выполняютсяпараллельно, в однойволне блоковSM 0SM 1SM 2SM 3SM 4SM 5Ядра запускаются на малых гридах Ядра запускаются на трех блоках по 1024 нити,устройство состоит из 6 SMKernelKernelKernelKernelKernelKernelKernelKernelМожем получить ускорение даже при небольшихкопированиях!Пример в NVVP Небольшие гридыДополнительные проблемы Устройство не поддерживает параллельноекопирование в обе стороны В рассматриваемой архитектуре одна аппаратнаяочередь блоковВажен порядок отправки команд!Нет параллельных копирований Отправка команды копирования блокирует стартвыполнения всех копирований в другую сторону,отправляемых после неё в любые потоки Ускорение только за счет параллельного копирования ивыполнения ядерМодельный примерfor (int i = 0; i < 3; ++i) {cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i *size, size, cudaMemcpyHostToDevice, stream[i]);MyKernel <<<100, 512, 0, stream[i]>>> (outputDevPtr+ i * size, inputDevPtr + i * size, size);cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i* size, size, cudaMemcpyDeviceToHost, stream[i]);}Отправка командHtoDKerDtoH HtoDStream[0]KerDtoH HtoDStream[1]KerDtoHStream[2]Нет параллельных копированийОтправка командHtoDKerDtoH HtoDStream[0]KerDtoH HtoDKerDtoHStream[2]Stream[1]Не будут выполняться параллельноHtoDKerDtoHHtoDKerDtoHHtoDKerDtoHПример в NVVPЕсли поменять порядок запусковОчередь командHtoDKer[0][0]HtoD DtoHHtoD[1][0]KerDtoHHtoDKerKerHtoD DtoH[1][2][1]DtoHHtoDKerDtoHKerDtoH[2][2]Пример в NVVPЕдинственная аппаратная очередь Если в поток отправлен запуск ядра, то все последующиезапуски команд в тот же поток должны начать выполнятьсятолько после его полного завершения В устройствах с Compute Capability <= 3.0 одна аппаратнаяочередь блоков, в которую попадают блоки всех запусков ядерво всех потоках Можно быть уверенным, что какой-либо запуск ядраполностью отработал, только когда в очереди большенет блоков!Единственная аппаратная очередьЕсли в поток отправлен запуск ядра, то все последующие запускикоманд в тот же поток должны начать выполняться только послеего полного завершения+Можно быть уверенным, что какой-либо запуск ядра полностьюотработал, только когда когда в очереди больше не блоков=Неявная синхронизация перед стартом выполнения зависимой отзапуска ядра команды: Начало выполнения команды откладывается до момента,когда в очереди не останется блоков Добавление новых блоков в очередь приостанавливается,пока команда не начнет выполнятьсяЕдинственная аппаратная очередь Начало выполнения зависимой от запуска ядра командыоткладывается до момента, когда в очереди не останется блоков Зависимая от запуска ядра команда может параллельновыполняться только с последней волной запуска ядра вдругом потоке Добавление новых блоков в очередь приостанавливается, показависимая от запуска ядра команда не начнет выполняться Запуски всех ядер во всех потоках приостанавливаются домомента, когда полностью отработает запуск-зависимость.Единственная аппаратная очередьОтправка командHtoDKerDtoH HtoDStream[0]KerDtoH HtoDKerDtoHStream[2]Stream[1]DtoH в блокирует все последующиезапуски ядерОжиданиеHtoDРеальностьKerDtoHHtoDKerDtoHHtoDKerHtoDDtoHKerDtoHHtoDKerDtoHHtoDKerDtoHЕдинственная аппаратная очередьОтправка командKerKerKerStream[0]Stream[1]Stream[2]DtoH в блокирует все последующиезапуски ядерОжиданиеРеальностьKerKerKerKerKerKerЕдинственная аппаратная очередьОтправка командHtoD HtoD HtoD[0][1][2]KerKer[0][1]Ker DtoH DtoH DtoH[2][0][2][1]Начнет выполняться когда опустееточередь блоковОжиданиеHtoDРеальностьKerDtoHHtoDKerDtoHHtoDKerHtoDHtoDDtoHDtoHKerDtoHKerHtoDKerDtoHПример в NVVPВывод Мало копирований с запусками ядер на большихгридах – с потоками лучше не возиться Аккуратно отправлять команды на GPUНеявная синхронизация Неявная синхронизация (ожидание завершения всехкоманд на устройтве ) выполняется перед: Выделением page-locked памяти / памяти на устройстве cudaMemSet Копированием между пересекающимися областямипамяти на устройстве Отправкой команды в поток по-умолчанию Переключением режима кеша L1 Если между отправкой двух команд в разные потоки стоитчто-то из этого списка – параллельного выполнения небудетСобытия (cudaEvent) Маркеры, приписываемые «точкам программы» Можно проверить произошло событие или нет Можно замерить время между двумя произошедшимисобытиями Можно синхронизоваться по событию, т.е.заблокировать CPU-поток до момента его наступления• «Точки программы» расположены между отправкамикоманд на GPUЗапись событияcudaError_t cudaEventRecord (cudaEvent_t event,cudaStream_t stream = 0) Приписывает событие точке программы, в которойвызываетсяcudaMemcpyAsync(…)…;//код без запуска командkernel<<<….>>> (…);cudaEventRecord(event, stream)cudaMemcpyAsync(…)ТочкипрограммыСовершение события Событие происходит когда выполнение команд на GPUреально доходит до точки, к которой в последний раз былоприписано событиеСовершение события Событие происходит когда завершаются все команды,помещённые в поток, к которому приписано событие, допоследнего вызова cudaEventRecord для него Если событие приписано потоку по умолчанию (stream = 0),то оно происходит в момент завершения всех команд,помещённых во все потоки до последнего вызоваcudaEventRecord для негоСинхронизация по событиюcudaError_t cudaEventQuery(cudaEvent_t event) Возвращает cudaSuccess, если событие уже произошло(вся работа до последнего cudaEventRecordвыполнена): иначе cudaErrorNotReadycudaError_t cudaEventSynchronize(cudaEvent_tevent) Возвращает управление хостовой нити только посленаступления событияСинхронизация на GPUcudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event,unsigned int flags ) Команды, отправленные в stream начнут выполнятьсяпосле наступления события event Синхронизация будет эффективно выполнена на GPU При stream == NULL будут отложены все команды всехпотоков Событие event может быть записано на другом GPU Синхронизация между GPUСинхронизация на GPUk_A1<<<1,1,0,streamA>>>(d); // A1cudaEventRecord(halfA,streamA);cudaStreamWaitEvent(streamB,halfA,0);k_B1<<<1,1,0,streamB>>>(d); // B1 начнется послезавершения A1cudaEventRecord(halfB,streamB);cudaStreamWaitEvent(streamA,halfB,0);k_A2<<<1,1,0,streamA>>>(d); // A2 начнется послезавершения B1k_B2<<<1,1,0,streamB>>>(d); // B2Синхронизация на GPUСинхронизация по потокуcudaError_t cudaStreamQuery(cudaStream_t stream)cudaError_t cudaStreamSynchronize(cudaStream_t event)cudaStreamCallbacktypedef void (*cudaStreamCallback_t)(cudaStream_t stream, cudaError_t status,void *userData )cudaError_t cudaStreamAddCallback (cudaStream_t stream, cudaStreamCallback_t callbac,void *userData, unsigned int flags ) callback будет вызван когда выполнятся всепредшествующие команды, отправленные в поток В callback запрещены обращения к CUDA APIИспользование нескольких GPUCUDA+openmpCUDA+MPIP2P обмены между GPUCUDA Context Аналог процесса CPU Выделения памяти, выполнение операций происходит врамках некоторого контекста (=процесса) Отдельное адресное пространство Выделенная память неявноосвобождается приудалении контекста Операции из разныхконтекстов не могутвыполнятся параллельноCUDA context• Адресноепространство• Ресурсы• ОперацииCUDA Context Контексты устройств неявно создаются при инициализацииCUDA-runtime На каждом устройстве создается по одному контексту –«primary-контекст» Все нити программы совместно их используют Инициализация CUDA-runtime происходит неявно, припервом вызове любой функции, не относящейся к Device /Version Management (см.