Лекция. CUDA 1 (Колганов) (1186105), страница 2

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

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

Выходные данные получим в resultOnDevicecudaMemcpy(resultOnHost , resultOnDevice , nb, cudaMemcpyDeviceToHost); // переслать результаты на хостcudaFree(inputDataOnDevice) ; // освободить память на устройствеcudaFree(resultOnDevice )) ; // освободить память на устройствеПример: Запуск ядраint n = getSize(); // размер задачи//определения указателей, получение входных данных на хостеcudaMalloc( (void**)& inputDataOnDevice, nb) ; // выделить память на устройстве для входных данныхcudaMalloc( (void**)& resultOnDevice, nb) ; // выделить память на устройстве для хранения результатаcudaMemcpy(inputDataOnDevice, inputDataOnHost , nb, cudaMemcpyHostToDevice); // переслать на устройствовходные данныеdim3 blockDim = dim3(512), gridDim = dim3( (n – 1) / 512 + 1 ); // рассчитать конфигурацию запускаkernel <<< gridDim, blockDim >>> (inputDataOnDevice, resultOnDevice, n); // запустить ядро с рассчитаннойконфигурацией и параметрамиcudaMemcpy(resultOnHost , resultOnDevice , nb, cudaMemcpyDeviceToHost); // переслать результаты на хостcudaFree(inputDataOnDevice) ; // освободить память на устройствеcudaFree(resultOnDevice )) ; // освободить память на устройствеМодель исполнения SIMTCUDA и классификация ФлиннаАрхитектура ЭВМSIMD – всепроцессыодновременновыполняют однуинструкциюSMP – все процессыимеют равныеправа на доступ кпамятиMIMD – каждыйпроцессвыполняетсянезависимо отдругих,MPPMISDNUMASISDcc-NUMACUDA и классификация Флинна У Nvidia собственная модель исполнения,имеющая черты как SIMD, так и MIMD: Nvidia SIMT: Single Instruction – Multiple ThreadSIMDNvidiaSIMTMIMD(SMP)SIMT: виртуальные нити, блоки Виртуально все нити: выполняются параллельно (MIMD) Имеют одинаковые права на доступ к памяти (MIMD :SMP) Нити разделены на группы одинакового размера (блоки): В общем случае (есть исключение) , глобальнаясинхронизация всех нитей невозможна, нити из разныхблоков выполняются полностью независимо Есть локальная синхронизация внутри блока, нити из одногоблока могут взаимодействовать через специальную память Нити не мигрируют между блоками.

Каждая нитьнаходится в своём блоке с начала выполнения и до конца.SIMT: аппаратное выполнение Все нити из одного блока выполняются на одноммультипроцессоре (SM)Максимальное число нитей в блоке – 1024Блоки не мигрируют между SMРаспределение блоков по мультироцесссорам непредсказуемоКаждый SM работает независимо от другихБлокипрограммыВиртуальныйблок нитейGigaThread engineБлоки и варпы Блоки нитей по фиксированному правилу разделяются нагруппы по 32 нити, называемые варпами (warp) Все нити варпа одновременно выполняют одну общуюинструкцию (в точности SIMD-выполнение) Warp Scheduler на каждом цикле работы выбирает варп,все нити которого готовы к выполнению следующейинструкции и запускает весь варпВиртуальныйблок нитейВарпWarp SchedulerВетвление (branching) Все нити варпа одновременно выполняют одну и туже инструкцию. Как быть, если часть нитей эту инструкцию выполнятьне должна? if(<условие>), где значение условия различается длянитей одного варпаЭти нити «замаскируются» нулями в специальномнаборе регистров и не будут её выполнять, т.е.

будутпростаиватьВетвление (branching)Несколько блоков на одном SM SM может работать с варпами нескольких блоководновременно Максимальное число резидентных блоков на одноммультипроцессоре – 8 Максимальное число резидентных варпов – 48 = 1536 нитейВиртуальный блокВиртуальнынитейй блокнитейВиртуальный блокнитейВарпWarp SchedulerЗагруженность (Occupancy) Чем больше нитей активно на мультипроцессоре, темэффективнее используется оборудование Блоки по 1024 нити – 1 блок на SM, 1024 нити, 66% от максимума Блоки по 100 нитей – 8 блоков на SM, 800 нитей, 52% Блоки по 512 нитей – 3 блока на SM, 1536 нитей, 100%SIMT и глобальная синхронизация В общем случае, из-за ограничений по числу нитей иблоков на одном SM, не удаётся разместить сразу всеблоки программы на GPU Часть блоков ожидает выполненияПоэтому в общем случае невозможна глобальнаясинхронизация Блоки выполняются по мере освобождения ресурсовНельзя предсказать порядок выполнения блоков Если все блоки программы удалось разместить, товозможна глобальная синхронизация через атомарныеоперацииВручную, специальная техника «Persistent Threads»SIMT и масштабирование Виртуальное GPU может поддерживатьмиллионы виртуальных нитей Виртуальные блокинезависимыПрограмму можно запустить налюбом количестве SM Аппаратное МультипроцессорынезависимыМожно «нарезать» GPU cразличным количеством SMпрограммаблокSIMTMIMDварпSIMDнитьNvidia SIMT-всенити из одноговарпа одновременновыполняют однуинструкцию, варпывыполняютсянезависимоSIMD – все нитиодновременно выполняютодну инструкциюMIMD – каждая нитьвыполняется независимо отдругих, SMP – все нитиимеют равные возможностидля доступа к памятиВыводыХорошо распараллеливаются на GPU задачи, которые: Имеют параллелизм по данным Одна и та же последовательность вычислений, применяемая к разнымданным Могут быть разбиты на подзадачи одинаковой сложности подзадача будет решаться блоком нитей Каждая подзадача может быть выполнена независимо от всехостальных нет потребности в глобальной синхронизации Число арифметических операций велико по сравнению с операциямидоступа в память для покрытия латентности памяти вычислениями Если алгоритм итерационный, то его выполнение может бытьорганизовано без пересылок памяти между хостом и GPU послекаждой итерации Пересылки данных между хостом и GPU накладны.

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

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