Лекция 7. CUDA-потоки применение (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть седьмаяКогда их стоит использовать?cudaStream Последовательность команд для GPU (запуски ядер,копирования памяти и т.д.), исполняемая строгопоследовательно следующая команда выполняется после полногозавершения предыдущейcudaStream Пользователь сам создает потоки и распределяеткоманды по ним По-умолчанию, все команды помещаются в «DefaultStream», равный нулюcudaStream Только команды из разных потоков, отличных от потока по-умолчанию, могут выполняться параллельно Пользователь сам задет необходимую синхронизациюмежду командами из разных потоков (при наличиизависимостей) В общем случае, порядок выполнения команд из разныхпотоков не определенАсинхронное копированиеКогда возвращается управление хостовой нитиHost->devicepageablememcpymemcpyAsyncpinnedDevice->hostpageableПослекопирования вбуфер*ПослеПослеполногополногозавершензавершенияияПослекопирования вбуферПослеполногозавершениясразуpinnedПослеполногозавершениясразуhost-hostПослеполногозавершенияПослеполногозавершенияdev-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 При распределении по потокам результат зависит от Соотношения копирования / cсчет Размера гридов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Вывод Мало копирований с запусками ядер на большихгридах – не стоит пытаться ускорить программу за счетиспользования потоков Аккуратно отправлять команды на GPUThe end.