Лекция 2. Программирования GPU CPU. Компиляция. Иерархия_ доступ к памяти (1265183)

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

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

Лихогруд Николайn.lihogrud@gmail.comЧасть втораяFermi StreamingMultiprocessor (SM) Потоковый мультипроцессор «Единица» построения устройства (как ядров CPU):32 скалярных ядра CUDA Core, ~1.5ГГц2 Warp Scheduler-аФайл регистров, 128KB3 Кэша – текстурный, глобальный (L1),константный(uniform)PolyMorphEngine – графический конвейерТекстурные юниты16 x Special Function Unit (SFU) –интерполяция и трансцендентная математикаодинарной точности16 x Load/StoreМодель выполненияБлокипрограммыGigaThread engineМаксимальное числоварпов на SM– 48 = 1536нитейРезидентные блокина мультипроцесореВиртуальныйблок нитейВиртуальныйблок нитейВиртуальныйблок нитейМаксимальное числоблоков на SM– 8ВарпWarp SchedulerпрограммаблокSIMTMIMDварпSIMDнитьNvidia SIMT-всенити из одноговарпа одновременновыполняют однуинструкцию, варпывыполняютсянезависимоSIMD – все нитиодновременно выполняютодну инструкциюMIMD – каждая нитьвыполняется независимо отдругих, SMP – все нитиимеют равные возможностидля доступа к памятиУтилизация латентности памяти GPU: Много нитей, покрывать обращения одних нитей впамять вычислениями в других за счёт быстрогопереключения контекста За счёт наличия сотен ядер и поддержки миллионов нитей(потребителей) на GPU легче утилизировать всю полосупропусканияCUDA Kernel («Ядро») Специальная функция, являющая входной точкой для кодана GPUНет возвращаемого значения (void)Выделена атрибутом __global__ Объявления параметров и их использование такое же, каки для обычных функций__global__ void kernel (int * ptr) {ptr = ptr + 1;ptr[0] = 100;….; //other code for GPU} Хост запускает именно «ядра», устройство их выполняетCUDA Grid Двумерный грид из трёхмерных блоков Логический индекс по переменной z у всех блоков равеннулю Каждый блок состоит из трёх «слоёв» нитей,соответствующих z=0,1,2Пример ядра__global__ void sum_kernel( int *A, int *B, int *C ){int threadLinearIdx =blockIdx.x * blockDim.x + threadIdx.x; //определить свой индексint elemA = A[threadLinearIdx]; //считать нужный элемент Aint elemB = B[threadLinearIdx]; // считать нужный элемент BC[threadLinearIdx] = elemA + elemB; //записать результат суммирования}На хостеcudaMalloc( (void**)& resultOnDevice, nb) ; // выделить памятьcudaMemcpy(inputDataOnDevice, inputDataOnHost , …); // переслать наустройство входные данныеdim3 blockDim = dim3(512);dim3 gridDim = dim3((n – 1) / 512 + 1 ); // рассчитать гридkernel <<< gridDim, blockDim >>> (inputDataOnDevice,…); // запустить ядроВыбор устройства, обработка ошибок, вычисление временивыполненияВыбор устройства struct cudaDeviceProp cтруктура с параметрами устройства.

Полный список параметром см.в документации cudaError_tcudaGetDeviceCount ( int* count ) записывает в *count число доступных устройств в системе cudaError_t cudaGetDeviceProperties (* prop, int device )cudaDeviceProp записывает параметры устройства с индексом device в *prop cudaError_t cudaSetDevice ( int device ) выбрать устройство c индексом device для проведения вычисленийВыбор устройстваint deviceCount=0, suitableDevice=-1;cudaDeviceProp devProp; // структура с параметрами устройстваcudaGetDeviceCount( &deviceCount ); // число доступных устройствprintf ( "Found %d devices\n", deviceCount );for ( int device = 0; device < deviceCount; device++ ) {cudaGetDeviceProperties ( &devProp, device ); // получить параметры устройства сзаданным номеромprintf( "Device %d\n", device );printf( "Compute capability: %d.%d\n", devProp.major, devProp.minor);pintf( "Name: %s\n", devProp.name );printf("Total Global Memory: %d\n", devProp.totalGlobalMem );if (ourRequirementsPassed(devProp)) // ищем устройство с нужными параметрамиsuitableDevice = device ;}assert(suitableDevice != -1);cudaSetDevice(suitableDevice); // Выбрать для работы заданное устройствоАсинхронность в CUDA Чтобы GPU больше времени работало в фоновом режиме,параллельно с CPU, некоторые вызовы являются асинхронными Отправляют команду на устройство и сразу возвращаютуправление хосту К таким вызовам относятся: Запуски ядер ( если CUDA_LAUNCH_BLOCKING не установлена на 1) Копирование между двумя областями памяти на устройстве Копирование с хоста на устройство менее 64KB Копирования, выполняемые функциями с окончанием *Async cudaMemSet – присваивает всем байтам области памяти наустройстве одинаковое значение (чаще всего используется дляобнуления)Асинхронность в CUDA Почему тогда верно работает код?//запуск ядра (асинхронно)sum_kernel<<< blocks, threads >>>(aDev, bDev, cDev);//переслать результаты обратно на хостcudaMemcpy(cHost, cDev, nb, cudaMemcpyDeviceToHost); Ведь хост вызывает cudaMemcpy до завершениявыполнения ядра!cudaStream Последовательность команд для GPU (запуски ядер, копированияпамяти и т.д.), исполняемая строго последовательно, следующаявыполняется после завершения предыдущей Команды из разных потоков могут выполняться параллельно,независимо от исполнения команд в других потоках Пользователь сам объединяет команды в потоки. Результаты взаимодействия команд из разных потоковнепредсказуемы По умолчанию, все команды помещаются в «Default Stream», равныйнулюАсинхронность в CUDA Почему тогда верно работает код?// запуск ядра (асинхронно)sum_kernel<<< blocks, threads >>>(aDev, bDev, cDev);//переслать результаты обратно на хостcudaMemcpy(cHost, cDev, nb, cudaMemcpyDeviceToHost); Вызов ядра и cudaMemcpy попадают в один поток (потокпо умолчанию) Устройство гарантирует их последовательноевыполнениеОбработка ошибок Коды всех возникающих ошибок автоматическизаписываются в единственную специальную хостовуюпеременную типа enum cudaError_t Эта переменная в каждый момент времени равна коду последнейошибки, произошедшей в системеcudaPeekAtLastError() – возвращает текущеезначение этой переменной cudaError_tcudaGetLastError() - возвращает текущеезначение этой переменной и присваивает ей cudaSuccess cudaError_t const char* cudaGetErrorString (cudaError_t–по коду ошибки возвращает её текстовое описаниеerror )Обработка ошибок Простейший способ быть уверенным, что в программе не произошлоCUDA-ошибки: В конце main вставитьstd::cout << cudaGetErrorString(cudaGetLastError());События Маркеры, приписываемые точкам программы Можно проверить произошло событие или нет Можно замерить время между двумя произошедшими событиями Событие происходит когда завершаются все команды, помещённые впоток, к которому приписано событие, до последнего вызоваcudaEventRecord для него Если событие приписано потоку по умолчанию (stream = 0), то онопроисходит в момент завершения всех команд, помещённых во всепотоки до последнего вызова cudaEventRecord для негоИзмерение времени выполнения Расчёт времени выполнения ядра Записать одно событие до вызова ядра, другое – сразу после:cudaEvent_t start, stop;float time;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord( start, 0 ); // В потоке могут быть некоторые команды// приписываем этой точке событие startkernel<<<grid,threads>>> ( params);cudaEventRecord( stop, 0 ); // К данному моменту в поток отправлена командавызова ядра, приписываем этой точке событие stopcudaEventSynchronize( stop ); // Дождаться событие stop, т.е.

завершениявыполнения всех команд, в т.ч. запуска ядраcudaEventElapsedTime( &time, start, stop ); // Время между двумя событиямиcudaEventDestroy( start );cudaEventDestroy( stop );Как скомпилировать и запустить программу сиспользованием CUDAОсобое отношение к .cuПри работе с CUDA используются расширения Си++: Конструкция запуска ядра <<< …. >>> Встроенные переменные threadIdx, blockIdx Квалификаторы __global__ __device__ и т.д. …. Эти расширения могут быть обработаны только в *.cu файлах! cudafe не запускается для файлов с другим расширением В этих файлах можно не делать #include <cuda_runtime.h> Вызовы библиотечных функций вида cuda* можно располагать в*.cpp файлах Они будут слинкованы обычным линковщиком из библиотекиlibcudart.soКомпиляция хост-кодаВ файле test.cpp : Основной хост-код.

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

Тип файла PDF

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

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

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