Лихогруд Николай Часть седьмая

Когда их стоит использовать?

cudaStream Последовательность команд для GPU (запуски ядер, копирования памяти и т.д.), исполняемая строго последовательно следующая команда выполняется после полного завершения предыдущей

cudaStream Пользователь сам создает потоки и распределяет команды по ним По-умолчанию, все команды помещаются в «Default Stream», равный нулю

cudaStream Только команды из разных потоков, отличных от потока по- умолчанию, могут выполняться параллельно Пользователь сам задет необходимую синхронизацию между командами из разных потоков (при наличии зависимостей) В общем случае, порядок выполнения команд из разных потоков не определен

Асинхронное копирование Host->deviceDevice->host host-hostdev-dev pageablepinnedpageablepinned memcpy После копирования в буфер* После полного завершен ия сразу memcpyAsync После копирования в буфер сразу После полного завершен ия сразу После полного завершен ия сразу Когда возвращается управление хвостовой нити *В начале работы неявно вызывается cudaDeviceSynchronize

Параллельная работа хоста и устройства Ядра выполняются асинхронно Копирование между pinned-памятью и памятью устройства при помощи cudaMemcpyAsync также выполняется асинхронно => Добиться параллельной работы хоста и устройства достаточно просто!

Параллельная работа хоста и устройства Пример: cudaMallocHost(&aHost, size); cudaMemcpyAsync( aDev, aHost, size, cudaMemcpyHostToDevice); kernel >>(aDev, bDev); cudaMemcpyAsync( bHost, cudaMemcpyHostToDevice); doSomeWorkOnHost(); doSomeWorkOnHost будет выполняться параллельно с копированиями и выполнением ядра

Параллельное выполнение команд на GPU Команды из разных потоков, отличных от потока по- умолчанию, могут исполняться параллельно В зависимости от аппаратных возможностей Возможные случаи: Параллельные копирование и выполнение ядра Параллельные выполнение ядер Параллельные копирования с хоста на устройство и с устройства на хост

Если cudaDeviceProp::asyncEngineCount > 0 устройство может выполнять параллельно копирование и счет ядра Хостовая память должна быть page-locked cudaMallocHost(&aHost, size); cudaStreamCreate(&stream1); cudaStreamCreate(&stream2); cudaMemcpyAsync( aDev, aHost, size, cudaMemcpyHostToDevice, stream1); kernel >>(…); Копирование & выполнение ядра

Если cudaDeviceProp::concurrentKernels > 0 устройство может выполнять ядра параллельно cudaStreamCreate(&stream1); cudaStreamCreate(&stream2); kernel1(data_1); kernel2(data_2); Параллельное выполнение ядер

Если cudaDeviceProp::asyncEngineCount == 2 устройство может выполнять параллельно копирование в обе стороны и счет ядра cudaMallocHost(&aHost, size); cudaMallocHost(&bHost, size); // создать потоки cudaMemcpyAsync( aDev, aHost, size, cudaMemcpyHostToDevice, stream1); cudaMemcpyAsync( bHost, bDev, size, cudaMemcpyDeviceToHost, stream2); kernel >>(…); Копирование в обе стороны

Обозначения Ядро = код для GPU, __global void kernel(…) {} Размер ядра = длительность вычисления отдельной нитью Большое, сложное ядро = ядро большого размера Выполнение ядра = выполнение ядра на некотором гриде Запуск ядра = команда запуска ядра на гриде Длительный запуск ядра = запуск ядра, который долго выполняется Время вычисления ~ размер ядра * размер гряда

Примеры Рассмотрим классическую схему работы с GPU: Копирование входных данных на GPU Выполнение ядра Копирование результатов обратно на хост Необходима синхронизация между командами, т.е. следующая выполняется после завершения предыдущей Kernel HtoD DtoH

Идеальный случай Выполнения копирований сопоставимы по времени с выполнением ядра Разобьем задачу на подзадачи Разделим грид на части и запустим то же ядро на подгрядах – результат не изменится Подзадаче нужна только часть данных для старта Запустим подзадачи в разных потоках Kernel HtoD DtoH

Идеальный случай Ker HtoD DtoH Kernel HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Время Stream 0 Stream 1 Stream 2 Stream 3 Запустим подзадачи в разных потоках

Идеальный случай Выполнение первой подзадачи начнется сразу после копирования нужной ей части данных Выполнение ядер будет происходить параллельно с копированиями Параллельное копирование в обе стороны, если аппаратура позволяет Макимальное ускорение:

Небольшие копирования Kernel HtoD DtoH Kernel Получится ли? Kernel

Небольшие копирования Kernel HtoD DtoH Kernel Получится ли? Kernel Нет! Суммарное время выполнения ядра на подгрядах никак не сделать меньше времени выполнения на целом гриде

Параллельное выполнение ядер Ресурсы GPU ограничены – до 16 SM, до 1536 нитей на одном SM, до 8-ми блоков на SM Ядра запускаются на грядах из миллионов нитей = тысячи блоков Блоки всех грибов попадают в одну общую очередь и выполняются по мере освобождения мультипроцессоров Мультипроцессоры – «bottle neck» этой очереди

Очередь блоков SM 0SM 1SM 2 Единственное место, где ядра могут выполняться параллельно Kernel 0 Kernel 1 Волны блоков Хвост первого ядра Гриды

А если синхронно? SM 0SM 1SM 2 Волны блоков Первое ядро еще не завершилось – следующее не может начаться Хвост Первого запуска ядра Kernel 1 Гриды Kernel 0

При большом приближении видно хвосты Пример в NVVP

Вывод При запуске двух ядер в разных потоках они могут выполняться параллельно только на границах Когда последняя волна блоков (хвост) первого из запущенных ядер не полностью загружает устройство Суммарное время выполнения одинаковых по сложности ядер, запущенных в разных потоках:

Вывод Суммарное время отдельных выполнений ядра на подгрядах в разных потоках равно времени выполнения на целом гриде Суммарное число блоков не меняется Суммарное время отдельных выполнений ядра на подгрядах в одном потоке >= времени выполнения на целом гриде Из-за простоя на границах

Небольшие копирования Kernel HtoD DtoH Kernel Время

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

Множество задач При распределении по потокам результат зависит от Соотношения копирования / cсчет Размера грибов Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH

Ядра запускаются на больших грядах Гриды большие => хвосты запусков ядер занимают малую долю в общем число блоков => доля параллельного выполнения невелика Выигрываем в основном за счет параллельных копирований Мало копирований – мало ускорения! Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH

Ядра запускаются на малых грядах Пусть ядра запускаются на грядах такого размера, что ресурсов хватает для размещения всех блоков нескольких грибов Например двух Пусть ядра примерно одинаковой сложности Тогда два ядра, запущенные таких грядах, будут выполняться параллельно!

Ядра запускаются на малых грядах Ядра запускаются на трех блоках по 1024 нити, устройство состоит из 6 SM SM 2SM 4SM 0SM 3SM 5SM 1 Очередь блоков Два ядра выполняются параллельно, в одной волне блоков

Ядра запускаются на малых грядах Можем получить ускорение даже при небольших копированиях! Kernel Ядра запускаются на трех блоках по 1024 нити, устройство состоит из 6 SM

Небольшие гряды Пример в NVVP

Дополнительные проблемы Устройство не поддерживает параллельное копирование в обе стороны В рассматриваемой архитектуре одна аппаратная очередь блоков Важен порядок отправки команд!

Нет параллельных копирований Отправка команды копирования блокирует старт выполнения всех копирований в другую сторону, отправляемых после неё в любые потоки Ускорение только за счет параллельного копирования и выполнения ядер

Модельный пример for (int i = 0; i < 3; ++i) { cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]); MyKernel >> (outputDevPtr + i * size, inputDevPtr + i * size, size); cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]); } Отправка команд Stream[0] Stream[1] Stream[2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH

Нет параллельных копирований Отправка команд Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Не будут выполняться параллельно Stream[0] Stream[1] Stream[2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH

Пример в NVVP

Очередь команд [0] [1] [2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH [0] [1] [0] [1] [2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Если поменять порядок запусков

Пример в NVVP

Единственная аппаратная очередь Если в поток отправлен запуск ядра, то все последующие запуски команд в тот же поток должны начать выполняться только после его полного завершения В устройствах с Compute Capability

Единственная аппаратная очередь Если в поток отправлен запуск ядра, то все последующие запуски команд в тот же поток должны начать выполняться только после его полного завершения + Можно быть уверенным, что какой-либо запуск ядра полностью отработал, только когда когда в очереди больше не блоков = Неявная синхронизация перед стартом выполнения зависимой от запуска ядра команды: Начало выполнения команды откладывается до момента, когда в очереди не останется блоков Добавление новых блоков в очередь приостанавливается, пока команда не начнет выполняться

Единственная аппаратная очередь Начало выполнения зависимой от запуска ядра команды откладывается до момента, когда в очереди не останется блоков Зависимая от запуска ядра команда может параллельно выполняться только с последней волной запуска ядра в другом потоке Добавление новых блоков в очередь приостанавливается, пока зависимая от запуска ядра команда не начнет выполняться Запуски всех ядер во всех потоках приостанавливаются до момента, когда полностью отработает запуск-зависимость.

Единственная аппаратная очередь Отправка команд Stream[0] Stream[1] Stream[2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH DtoH в блокирует все последующие запуски ядер Ожидание Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Реальность

Единственная аппаратная очередь Отправка команд Stream[0] Stream[1] Stream[2] Ker DtoH в блокирует все последующие запуски ядер Ожидание Реальность Ker

Единственная аппаратная очередь Отправка команд [0] [1] [2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH [0] [1] [0] [1] [2] Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Ожидание Ker HtoD DtoH Ker HtoD DtoH Ker HtoD DtoH Реальность Начнет выполняться когда опустеет очередь блоков

Пример в NVVP

Вывод Мало копирований с запусками ядер на больших грядах – не стоит пытаться ускорить программу за счет использования потоков Аккуратно отправлять команды на GPU

The end