Лекция 6. Pinned память. CUDA-потоки (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть шестаявремя выполнения задачи =время работы ядра + обмен даннымимежду CPU и GPU Как сократить накладные расходы на обмен даннымимежду CPU и GPU? Ускорение копирований Выполнение копирований параллельно с вычислениямиDMA & zero-copy “Zero-copy”- копирование памяти без участияцентрального процессора Копирование выполняется спец. контроллером,процессор переключается на другие задачи DMA (Direct memory access) – прямой доступ коперативной памяти, без участия ЦП Реализуется через zero-copy операции Скорость передачи увеличивается, так как данные непересылаются в ЦП и обратноDMA и виртуальная память Виртуальная память организованна в страницы,отображаемые на физические страницы ОП Виртуальные страницы могут быть Перемещены в оперативной памяти Отгружены на диск (swapping)В таких условиях реализовать DMA очень сложно!DMA и виртуальная память Запретим перемещение страниц по ОП и ихвыгрузку на диск Привяжем страницы виртуальной памяти к страницамфизической Эти физические страницы теперь недоступны ОС дляразмещения новых виртуальных страниц (paging)page-able  page-lockedДля page-locked памяти DMA реализуемо безсущественных накладных расходовPage-locked память & CUDA «Pinned» - синоним, «прикрепленный» В CUDA можно напрямую выделить page-locked(pinned) память на хосте или сделать таковойпамять, выделенную ранее Операции копирования Host  Device с нейпроисходят быстрее и могут выполнятьсяпараллельно с работой ядерcudaHostRegister “Залочить” память, предварительно выделеннуюобычным способом:float *ptr = malloc(n * sizeof(float))cudaHostRegister(ptr, n * sizeof(float),0);cudaMemcpy(devPtr,prt, n * sizeof(float),cudaMemcpyHostToDevice)…cudaHostUnregister(ptr);Mapped pinned-память Pinned-память можно отобразить в виртуальноеадресное пространство GPU Нити смогут обращаться к ней напрямую, безнеобходимости копирования в память GPU Необходимые копирования будут неявно выполнятьсяасинхронно, параллельно с работой ядра C хоста память будет так же доступнаMapped pinned-память Залочить память на хосте и получить указатель, покоторому к ней можно обращаться из ядер:cudaHostRegister(ptr,n * sizeof(float),cudaHostRegisterMapped);float *ptrForDevice = NULL;cudaHostGetDevicePointer(&ptrForDevice, ptr, 0);// не нужно выделять память на GPU и копировать в// неё входные данныеkernel<<<…>>>(ptrForDevice,…);Mapped pinned-память Для активации возможности маппирования pinned-памяти: До первого вызова функции из cuda-runtime (т.е.

доинициализации устройства) установить флаг инициализацииcudaDeviceMapHost:cudaSetDeviceFlags(cudaDeviceMapHost);cudaSetDevice(0);// инициализируется с флагами Проверить свойство устройства canMapHostMemory:cudaDeviceProp deviceProp;cudaGetDeviceProperties(0, &deviceProp);if (deviceProp.canMapHostMemory) {…}Прямое выделение pinned-памяти Самое простое:float *ptr = NULL;cudaMallocHost(&ptr, n * sizeof(float)); С флагами:cudaHostAlloc(&ptr, n * sizeof(float),cudaHostAllocDefault);Возможные флаги: cudaHostAllocDefault: эмулирование cudaMallocHost(). cudaHostAllocMapped: аналогично cudaHostRegisterMappedcudaMemcpy* и pinned-память Копировании обычной (pageable) памяти с хоста на GPUпроисходит через DMA к промежуточному буферу вpinned-памяти Управление хосту возвращается после выполнениякопирований в этот буфер, но необязательно дозавершения DMAcudaMemcpy* и pinned-память Копировании обычной (pageable) памяти с хоста на GPUпроисходит через DMA к промежуточному буферу вpinned-памяти Поэтому копирование сразу из pinned-памяти быстрее –не нужно выделять память под буфер и копировать в негоданныеТестfloat *hostPtr = (float *)malloc(numberOfBytes);cudaEventRecord(startPageable, 0);cudaMemcpy(devicePtr, hostPtr, numberOfBytes,cudaMemcpyHostToDevice);cudaEventRecord(stopPageable, 0);cudaHostRegister(hostPtr, numberOfBytes, 0);cudaEventRecord(startPinned, 0);cudaMemcpy(devicePtr, hostPtr, numberOfBytes,cudaMemcpyHostToDevice);cudaEventRecord(stopPinned, 0);cudaDeviceSynchronize();Тестfloat elapsedPinned, elapsedPageable;cudaEventElapsedTime(&elapsedPageable, startPageable,stopPageable);cudaEventElapsedTime(&elapsedPinned, startPinned,stopPinned);printf("Copy from pageable %f\n", elapsedPageable);printf("Copy from pinned %f\n", elapsedPinned);$./a.outCopy from pageable 555.893066Copy from pinned 339.375580Замечания Выделение pinned-памяти занимает больше времени,чем обычный malloc Доступный для выделения объем сильно ограничен Чрезмерное использование page-locked памятидеградирует систему Для освобождения использовать cudaFreeHost(),cudaHostUnregister()Unified Virtual Address (UVA) На 64-битной архитектуре, начиная с поколения Fermi(сс 2.0), используется единое виртуальное адресноепространство для памяти хоста и всех устройств Unified Virtual Address space, UVA Если UVA включено, тоcudaDeviceProp::unifiedAddressing == 1Unified Virtual Address (UVA) Без UVA для каждого указателя хранятся метаданныео том где реально расположена память, на которую онуказывает С UVA эта информация «вшита» в значение указателя Диапазоны адресов всех GPU и CPU не пересекаютсяUnified Virtual Address (UVA) Чтобы узнать где реально расположена память:float *ptr;cudaPointerAttributes pointerAttributes;cudaPointerGetAttribute(&pointerAttributes, ptr)Unified Virtual Address (UVA)struct cudaPointerAttributes {enum cudaMemoryType memoryType;int device;void *devicePointer;void *hostPointer;} memoryType - cudaMemoryTypeHost | cudaMemoryTypeDevice device - устройство, на котором расположена память devicePointer - NULL, если не доступна с текущегоустройства hostPointer - NULL, если не доступна с хостаPinned-память и UVA C UVA память, выделенная через сudaHostAlloc() Автоматически является mapped Доступна с хоста и с любого GPU по одному и тому жеуказателю (т.к.

адресное пространство единое)Не нужно использовать cudaHostGetDevicePointer()Исключение – cudaHostAllocWriteCombinedPinned-память и UVA Для памяти, залоченной через cudaHostRegister идля write-combined памяти указатели для хоста и дляустройства являются разными Нужен cudaHostGetDevicePointer()Пример Без UVA и mapped памяти:float *ptr = NULL;cudaHostAlloc(&ptr, 1024, 0);float *ptrForDevice = NULL;cudaMalloc(&ptrForDevice, 1024);cudaMemcpy(ptrForDevice, ptr, 1024,cudaMemcpyHostToDevice)kernel<<<…>>>(ptrForDevice,…);Пример С mapped памятью:cudaSetDeviceFlags(cudaDeviceMapHost);cudaDeviceProp deviceProp;cudaGetDeviceProperties(device, &deviceProp);if (deviceProp.canMapHostMemory ) {float *ptr = NULL;cudaHostAlloc(&ptr, 1024, cudaHostAllocMapped)float *ptrForDevice = NULL;cudaHostGetDevicePointer(&ptrForDevice, ptr, 0);kernel<<<…>>>(ptrForDevice,…);}Пример С mapped памятью и UVA:cudaSetDeviceFlags(cudaDeviceMapHost)cudaDeviceProp deviceProp;cudaGetDeviceProperties(device, &deviceProp);if (deviceProp.unifiedAddressing ) {float *ptr = NULL;cudaHostAlloc(&ptr, 1024, cudaHostAllocMapped)kernel<<<…>>>(ptr,…);}Примерfloat *ptrForDevice = NULL;if (deviceProp.unifiedAddressing ) {ptrForDevice = ptr} else if (deviceProp.canMapHostMemory ) {cudaHostGetDevicePointer(&ptrForDevice, ptr, 0);} else {cudaMalloc(&ptrForDevice, 1024);cudaMemcpy(ptrForDevice, ptr, 1024,cudaMemcpyHostToDevice)}kernel<<<…>>>(ptrForDevice,…);cudaMemcpy* и UVA С UVA система в состоянии сама определить гденаходится память Можно указывать cudaMemcpyDefault вcudaMemcpyKind:float *dstPtr, *srcPtr;cudaMemcpy(dstPtr, srcPtr,n*sizeof(float), cudaMemcpyDefault)Выводывремя выполнения задачи =время работы ядра + обмен данными между CPU и GPU page-locked (pinned) память позволяетУменьшить время обмена данными2.

Упростить хост-код при использовании mapped pinnedпамяти и доступе к ней напрямую из ядер Не нужно возиться с пересылкой данных на GPU иобратно С UVA обращаемся к памяти с хоста и с устройствапо одному указателю1.cudaStream Последовательность команд для GPU (запуски ядер,копирования памяти и т.д.), исполняемая строгопоследовательно следующая команда выполняется после полногозавершения предыдущейcudaStream Пользователь сам создает потоки и распределяеткоманды по ним По умолчанию, все команды помещаются в «DefaultStream», равный нулюcudaStream Только команды из разных потоков, отличных от потока по-умолчанию, могут выполняться параллельно Пользователь сам задет необходимую синхронизациюмежду командами из разных потоков (при наличиизависимостей) В общем случае, порядок выполнения команд из разныхпотоков неопределенСоздание и уничтожениеcudaStream_t stream;cudaStreamCreate(&stream);…cudaStreamDestroy(stream); Поток привязывается к текущему активному устройству Перед отправлением команды нужно переключаться наустройство, к которому привязан поток Если попробовать отправить в него команду при другомактивном устройстве, будет ошибкаСоздание и уничтожениеcudaStream_t stream;cudaStreamCreate(&stream);…cudaStreamDestroy(stream); cudaStreamDestroy не выполняет синхронизацию Управление возвращается хостовому процессу сразу,реальное освобождение ресурсов произойдет послезавершения всех команд потокаАсинхронное копированиеcudaMemcpyAsync ( void* dst, const void* src,size_t count,cudaMemcpyKind kind,cudaStream_t stream = 0)cudaMemcpy2DAsync ( void* dst, size_t dpitch,const void* src, size_t spitch,size_t width, size_t height,cudaMemcpyKind kind,cudaStream_t stream = 0 )Асинхронное копированиеКогда возвращается управление хостовой нитиHost->devicepageablepinnedDevice->hostpageablememcpyПослекопирования вбуфер*ПослеПослеполногополногозавершен завершенияияmemcpyAsyncПослекопирования вбуферПослеполногозавершениясразуhost-hostdev-devПослеполногозавершенияПослеполногозавершениясразусразуПослеполногозавершениясразуpinned*В начале работы неявно вызывается cudaDeviceSynchronizeАсинхронное копированиеКогда возвращается управление хостовой нитиHost->devicepageablepinnedDevice->hostpageablememcpyПослекопирования вбуфер*ПослеПослеполногополногозавершен завершенияияmemcpyAsyncПослекопирования вбуферПослеполногозавершениясразуhost-hostdev-devПослеполногозавершенияПослеполногозавершениясразусразуПослеполногозавершениясразуpinned*В начале работы неявно вызывается cudaDeviceSynchronizeПараллельное выполнение команд Команды из разных потоков, отличных от потока поумолчанию, могут исполняться параллельно В зависимости от аппаратных возможностей Возможные случаи: Параллельные копирование и выполнение ядра Параллельные выполнение ядер Параллельные копирования с хоста на устройство и сустройства на хостКопирование & выполнение ядра Если 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>>>(…);Неявная синхронизация Неявная синхронизация (ожидание завершения всех команд наустройтве ) выполняется перед: Выделением page-locked памяти / памяти на устройстве cudaMemSet Копированием между пересекающимися областями памятина устройстве Отправкой команды в поток по-умолчанию Переключением режима кеша L1 Если между отправкой двух команд в разные потоки стоит что-тоиз этого списка – параллельного выполнения не будетСобытия (cudaEvent) Маркеры, приписываемые «точкам программы» Можно проверить произошло событие или нет Можно замерить время между двумя произошедшимисобытиями Можно синхронизоваться по событию, т.е.заблокировать CPU-поток до момента его наступления• «Точки программы» расположены между отправкамикоманд на GPUЗапись событияcudaError_t cudaEventRecord (cudaEvent_t event,cudaStream_t stream = 0) Приписывает событие к точке программы в потоке stream, вкоторой вызываетсяkernel<<<…, stream>>> (…);…;//нет запусков команд в потоке streamcudaEventRecord(event, stream);…;//нет запусков команд в потоке streamcudaMemcpyAsync(…, stream);Точка программы впотоке stream междувызовом ядра иасинхроннымкопированиемСовершение события Событие происходит, когда выполнение команд на GPUреально доходит до точки, к которой в последний раз былоприписано событиeСовершение события Событие происходит когда завершаются все команды,помещённые в поток, к которому приписано событие, допоследнего вызова 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Синхронизация на GPUA1<<<1, 1 , 0, streamA>>>(d); // A1cudaEventRecord(halfA, streamA);cudaStreamWaitEvent(streamB, halfA, 0);B1<<<1, 1, 0, streamB>>>(d); // B1 начнется послезавершения A1cudaEventRecord(halfB, streamB);cudaStreamWaitEvent(streamA, halfB, 0);A2<<<1, 1, 0, streamA>>>(d); // A2 начнется послезавершения B1B2<<<1, 1, 0, streamB>>>(d); // B2Синхронизация на GPUСинхронизация по потокуcudaError_t cudaStreamQuery (cudaStream_t stream);Возвращает cudaSuccess, если выполнены все команды в потокеstream, иначе cudaErrorNotReadycudaError_t cudaStreamSynchronize (cudaStream_t stream); Возвращает управление хостовой нити, когда завершится выполнениевсех команд, отправленных в поток streamcudaStreamCallbacktypedef void (*cudaStreamCallback_t)(cudaStream_t stream,cudaError_t status, void *userData );cudaError_t cudaStreamAddCallback (cudaStream_t stream, cudaStreamCallback_t callback,void *userData, unsigned int flags ); callback будет вызван, когда выполнятся все предшествующиекоманды, отправленные в поток В callback запрещены обращения к CUDA APIВыводыВ следующий раз, после детального разбораслучаев применения потоковThe end.