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

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

Текст из файла

Лихогруд Николай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.

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

Тип файла PDF

PDF-формат наиболее широко используется для просмотра любого типа файлов на любом устройстве. В него можно сохранить документ, таблицы, презентацию, текст, чертежи, вычисления, графики и всё остальное, что можно показать на экране любого устройства. Именно его лучше всего использовать для печати.

Например, если Вам нужно распечатать чертёж из автокада, Вы сохраните чертёж на флешку, но будет ли автокад в пункте печати? А если будет, то нужная версия с нужными библиотеками? Именно для этого и нужен формат PDF - в нём точно будет показано верно вне зависимости от того, в какой программе создали PDF-файл и есть ли нужная программа для его просмотра.

Список файлов лекций