2016 Ответы на экзаменационные вопросы (1186037), страница 5

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

Текст из файла (страница 5)

Проблема Косвенной адресации

• Использование косвенной адресации нежелательно, поскольку требует двух чтений из памяти (сначала A[i], потом A[i][j])

• A[i] для разных нитей варпа скорее всего будут в одной кеш-линии -> одно обращение к кешу

• Но A[i,j] в общем случае могут быть «разбросаны»

• Принципиального решения нет!

• Скорее всего придется переработать алгоритм

1. Методы эффективной организации параллельных вычислений на графических процессорах.

Параллельная работа хоста и устройства

• Ядра выполняются асинхронно

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

Параллельное выполнение команд на GPU

• Команды из разных потоков, отличных от потока по умолчанию, могут исполняться параллельно

  • В зависимости от аппаратных возможностей

• Возможные случаи:

  • Параллельные копирование и выполнение ядра

  • Параллельные выполнение ядер

  • Параллельные копирования с хоста на устройство и с устройства на хост

Копирование и выполнение ядра

• Если cudaDeviceProp::asyncEngineCount > 0 устройство может выполнять параллельно копирование и счет ядра

• Хостовая память долна быть page-locked

cudaMallocHost(&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

  • Выполнение ядра

  • Копирование результатов обратно на хост

• Необходима синхронизация между командами, т.е. следующая выполняется после завершения предыдущей

Идеальный случай

• Выполнения копирований сопоставимы по времени с выполнением ядра

• Разобьем задачу на подзадачи

  • Разделим грид на части и запустим то же ядро на подгридах – результат не изменится

  • Подзадаче нужна только часть данных для старта

  • Запустим подзадачи в разных потоках

• Запустим подзадачи в разных потоках

• Выполнение первой подзадачи начнется сразу после копирования нужной ей части данных

• Выполнение ядер будет происходить параллельно с копированиями

• Параллельное копирование в обе стороны, если аппаратура позволяет

Параллельное выполнение ядер

• Ресурсы GPU ограничены – до 16 SM, до 1536 нитей на одном SM, до 8-ми блоков на SM

• Ядра запускаются на гридах из миллионов нитей = тысячи блоков

  • Блоки всех гридов попадают в одну общую очередь и выполняются по мере освобождения мультипроцессоров

  • Мультипроцессоры – «bottle neck» этой очереди

• При запуске двух ядер в разных потоках они могут выполняться параллельно только на границах

• Когда последняя волна блоков (хвост) первого из запущенных ядер не полностью загружает устройство

• Суммарное время выполнения одинаковых по сложности ядер, запущенных в разных потоках:

• Суммарное время отдельных выполнений ядра на подгридах в разных потоках равно времени выполнения на целом гриде

  • Суммарное число блоков не меняется

• Суммарное время отдельных выполнений ядра на подгридах в одном потоке >= времени выполнения на целом гриде

  • Из-за простоя на границах

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

• Параллельного выполнения ядер и копирований

• Параллельного выполнения копирований

=>При небольших копированиях не имеет смысла возиться с потоками для подзадач

Ядра запускаются на больших гридах

• Гриды большие => хвосты запусков ядер занимают малую долю в общем число блоков => доля параллельного выполнения невелика

  • Выигрываем в основном за счет параллельных копирований

Мало копирований – мало ускорения!

Ядра запускаются на малых гридах

• Пусть ядра запускаются на гридах такого размера, что ресурсов хватает для размещения всех блоков нескольких гридов

  • Например двух

• Пусть ядра примерно одинаковой сложности

Тогда два ядра, запущенные таких гридах, будут выполняться параллельно!

Можем получить ускорение даже при небольших копированиях!

Дополнительные проблемы

• Устройство не поддерживает параллельное копирование в обе стороны

• В рассматриваемой архитектуре одна аппаратная очередь блоков

Важен порядок отправки команд!

Нет параллельных копирований

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

• Ускорение только за счет параллельного копирования и выполнения ядер

Единственная аппаратная очередь

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

• В устройствах с Compute Capability <= 3.0 одна аппаратная очередь блоков, в которую попадают блоки всех запусков ядер во всех потоках

• Можно быть уверенным, что какой-либо запуск ядра полностью отработал, только когда в очереди больше нет блоков!

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

+

Можно быть уверенным, что какой-либо запуск ядра полностью отработал, только когда когда в очереди больше не блоков

=

Неявная синхронизация перед стартом выполнения зависимой от запуска ядра команды:

• Начало выполнения команды откладывается до момента, когда в очереди не останется блоков

• Добавление новых блоков в очередь приостанавливается, пока команда не начнет выполняться

• Начало выполнения зависимой от запуска ядра команды откладывается до момента, когда в очереди не останется блоков

  • Зависимая от запуска ядра команда может параллельно выполняться только с последней волной запуска ядра в другом потоке

• Добавление новых блоков в очередь приостанавливается, пока зависимая от запуска ядра команда не начнет выполняться

  • Запуски всех ядер во всех потоках приостанавливаются до момента, когда полностью отработает запуск-зависимость.

Неявная синхронизация

• Неявная синхронизация (ожидание завершения всех команд на устройстве) выполняется перед:

  • Выделением page-locked памяти / памяти на устройстве

  • cudaMemSet

  • Копированием между пересекающимися областями памяти на устройстве

  • Отправкой команды в поток по умолчанию

  • Переключением режима кэша L1

• Если между отправкой двух команд в разные потоки стоит что-то из этого списка – параллельного выполнения не будет

События (cudaEvent)

• Маркеры, приписываемые «точкам программы»

  • Можно проверить произошло событие или нет

  • Можно замерить время между двумя произошедшими событиями

  • Можно синхронизоваться по событию, т.е. заблокировать CPU-поток до момента его наступления

• «Точки программы» расположены между отправками команд на GPU

Запись события

cudaError_t cudaEventRecord (cudaEvent_t event, cudaStream_t stream = 0)

Приписывает событие точке программы, в которой вызывается

Совершение события

• Событие происходит когда выполнение команд на GPU реально доходит до точки, к которой в последний раз было приписано событие

• Событие происходит когда завершаются все команды, помещённые в поток, к которому приписано событие, до последнего вызова cudaEventRecord для него

• Если событие приписано потоку по умолчанию (stream = 0), то оно происходит в момент завершения всех команд, помещённых во все потоки до последнего вызова cudaEventRecord для него

Синхронизация по событию

cudaError_t cudaEventQuery(cudaEvent_t event)

• Возвращает cudaSuccess, если событие уже произошло (вся работа до последнего cudaEventRecord

выполнена): иначе cudaErrorNotReady

cudaError_t cudaEventSynchronize(cudaEvent_t event)

• Возвращает управление хостовой нити только после наступления события

Синхронизация на GPU

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags )

• Команды, отправленные в stream начнут выполняться после наступления события event

  • Синхронизация будет эффективно выполнена на GPU

  • При stream == NULL будут отложены все команды всех потоков

• Событие event может быть записано на другом GPU

  • Синхронизация между GPU

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

Список файлов ответов (шпаргалок)