Лекция. CUDA 2 (Колганов) (Электронные лекции)

Колганов Александрalexander.k.s@mail.ruчасть 2Чип в максимальной конфигурации• 16 SM• 512 ядер CUDA Core• Кеш L2 758KB• GigaThreadEngine• Контроллеры памятиDDR5• Интерфейс PCIGPU 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программаблокMIMDNvidia SIMTварпSIMDнитьNvidia SIMT-всенити из одноговарпа одновременновыполняют однуинструкцию, варпывыполняютсянезависимоSIMD – все нитиодновременно выполняютодну инструкциюMIMD – каждая нитьвыполняется независимо отдругих, SMP – все нитиимеют равные возможностидля доступа к памятиВетвление (branching)Утилизация латентности памяти Цель: эффективно загружать ЯдраПроблема: латентность памятиРешение: CPU: Сложная иерархия кешей GPU: Много нитей, покрывать обращения одних нитей впамять вычислениями в других за счёт быстрогопереключения контекстаУтилизация латентности памяти GPU: Много нитей, покрывать обращения одних нитей впамять вычислениями в других за счёт быстрогопереключения контекста За счёт наличия сотен ядер и поддержки миллионов нитей(потребителей) на GPU легче утилизировать всю полосупропусканияSIMT и масштабирование Виртуальная GPU может поддерживатьмиллионы виртуальных нитей Виртуальные блокинезависимыПрограмму можно запустить налюбом количестве SM Аппаратная МультипроцессорынезависимыМожно «нарезать» GPU cразличным количеством SMCUDA 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Ориентация нити в гриде Осуществляется за счёт встроенных переменных: threaIdx.x threaIdx.y threaIdx.z - индексы нити в блоке blockIdx.x blockIdx.y blockIdx.z – индексты блока в гриде blockDim.x blockDim.y blockDim.z – размеры блоков внитях gridDim.x gridDim.y gridDim.z – размеры грида в блоках Линейный индекс нити в гриде:intintintintgridSizeX = blockDim.x * gridDim.x;gridSizeZ = … ; gridSizeY = …;gridSizeAll = gridSizeX * gridSizeY * gridSizeZ;threadLinearIdx =threaIdx.z * gridSizeY + threaIdx.y) * gridSizeX +threadIdx.x;Пример: сложение векторов__global__ void sum_kernel( int *A, int *B, int *C, int N){int threadLinearIdx =blockIdx.x * blockDim.x + threadIdx.x; //определить свой индексif (threadLinearIdx < N) { // все нити [0 ..

N – 1]int elemA = A[threadLinearIdx]; //считать нужный элемент Aint elemB = B[threadLinearIdx]; // считать нужный элемент BC[threadLinearIdx] = elemA + elemB; //записать результат}} Каждая нить Получает копию параметров (В данном случае, это адресавектором на GPU); Определяет своё положение в гриде threadLinearIdx ; Cчитывает из входных векторов элементы с индексомthreadLinearIdx и записывает их сумму в выходной вектор поиндексу threadLinearIdx ; рассчитывает один элемент выходного массива.Host Code Выбрать устройство По умолчанию, устройство с номером 0 Выделить память на устройстве Переслать на устройство входные данные Рассчитать грид Размер грида зависит от размера задачи Запустить ядро Переслать с устройства на хост результатВыбор устройства struct cudaDeviceProp cтруктура с параметрами устройства. Полный список параметром см.в документации cudaError_tcudaGetDeviceCount ( int* count ) записывает в *count число доступных устройств в системе cudaError_t cudaGetDeviceProperties (cudaDeviceProp* prop, int device ) записывает параметры устройства с индексом device в *prop cudaError_tcudaSetDevice ( 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); // Выбрать для работы заданное устройствоПример: Копирование памятиint n = getSize(); // размер задачиint nb = n * sizeof (float); // размер размер задачи в байтахfloat * inputDataOnHost = (float *)malloc( nb ) ;// память на хосте для входных данныхfloat * resultOnHost = (float *)malloc( nb ); // память на хосте для результатаfloat * inputDataOnDevice= NULL , *resultOnDevice = NULL; // память на устройствеgetInputData(inputDataOnHost); // получить входные данныеcudaMalloc( (void**)& inputDataOnDevice, nb) ; // выделить память на устройстве для входных данныхcudaMalloc( (void**)& resultOnDevice, nb) ; // выделить память на устройстве для хранения результатаcudaMemcpy(inputDataOnDevice, inputDataOnHost , nb, cudaMemcpyHostToDevice); // переслать на устройствовходные данные//запустить ядро.

Выходные данные получим в 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 )) ; // освободить память на устройствеАсинхронность в 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()); Рекомендуемый:#define SAFE_CALL(err) do \{ if (err != 0) \{ printf("ERROR [%s] in line %d: %s\n", __FILE__,__LINE__, cudaGetErrorString(err)); \exit(1); \}\} while(0)События Маркеры, приписываемые точкам программы Можно проверить произошло событие или нет Можно замерить время между двумя произошедшими событиями Событие происходит когда завершаются все команды, помещённые впоток, к которому приписано событие, до последнего вызова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, т.е.