Лекция 1. GPU аппаратно-программная модель (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть перваяGPGPU & CUDA GPU - Graphics Processing Unit GPGPU - General-Purpose computing on GPU, вычисленияобщего вида на GPU Первые GPU от NVIDIA с поддержкой GPGPU – GeForce восьмогопоколения, G80 (2006 г) CUDA - Compute Unified Device Architecture Программно-аппаратная архитектура от Nvidia, позволяющаяпроизводить вычисления с использованием графических процессоровCUDA растетУскорители в top500Преимущества GPGPU Соотношения Цена\производительность Производительность\энергопотреблениеGreen500Эффективность энергопотребленияСемейства GPU NvidiaВысокопроизводительныевычисленияПрофессиональнаяграфикаРазвлеченияКак устроено GPU?Compute Capability CUDA развивалась постепенно, многие возможности APIнедоступны на старых архитектурах Возможности устройства определяются его Compute Capability<номер поколения>.<номер модификации> Для эффективного программирования с использованием GPUнужно учитывать Compute Capability используемого устройстваCompute Capability Поколение Tesla ( не путать с линией продуктов для HPC ) 1.1 – базовые возможности CUDA, атомарные операции с глобальной памятью 1.2 - атомарные операции с общей памятью, warp vote-функции 1.3 – вычисления с двойной точностью Поколение Fermi 2.0 - новая архитектура чипа, ECC, кеши L1 и L2, асинхронное выполнениеядер, UVA и др. 2.1 - новая архитектура warp scheduler-ов Поколение Kepler 3.0, 3.2 – Новая архитектура чипа, Unified memory programming, warp shfl и др. 3.0 - Динамический параллелизм, Hyper Queqe и др. Поколение Maxwell sm_50 and sm_52 – Новая архитектура чипа ….CPU Intel Core I-7 Небольшое число мощныхнезависимых ядер 2,4,6,8 ядер, 2,66—3,6ГГц каждое Каждое физическое ядроопределяется системой как 2логических и может параллельновыполнять два потока (HyperThreading) 3 уровня кешей, большой кеш L3 На каждое ядро L1=32KB (data) +32KB ( Instructions), L2=256KB Разделяемый L3 до 20 mb Обращения в памятьобрабатываются отдельно длякаждого процесса\нитиCore I7-3960x,6 ядер, 15MB L3Fermi: StreamingMultiprocessor (SM) Потоковый мультипроцессор «Единица» построения устройства (как ядров CPU):32 скалярных ядра CUDA Core, ~1.5ГГц2 Warp Scheduler-аФайл регистров, 128KB3 Кэша – текстурный, глобальный (L1),константный(uniform)PolyMorphEngine – графический конвейерТекстурные юниты16 x Special Function Unit (SFU) –интерполяция и трансцендентная математикаодинарной точности16 x Load/StoreFermi: Чип в максимальной конфигурации• 16 SM• 512 ядер CUDA Core• Кеш L2 758KB• GigaThreadEngine• Контроллеры памятиDDR5• Интерфейс PCIKepler: SMX 192 cuda core 64 x DP Unit 32 x SFU 32x load/store Unit 4 x warp scheduler 256KB регистровKepler: Чип в максимальной конфигурации 15 SXM = 2880 cuda coreВычислительная мощностьПропускная способность памятиСравнение GPU и CPU Сотни упрощённых вычислительных ядер, работающих нанебольшой тактовой частоте ~1.5ГГц (вместо 2-8 на CPU) Небольшие кеши 32 ядра разделяют L1, с двумя режимами: 16KB или 48KB L2 общий для всех ядер, 768 KB, L3 отсутствует Оперативная память с высокой пропускной способностью ивысокой латентностью Оптимизирована для коллективного доступа Поддержка миллионов виртуальных нитей, быстроепереключение контекста для групп нитейУтилизация латентности памяти Цель: эффективно загружать ЯдраПроблема: латентность памятиРешение: CPU: Сложная иерархия кешей GPU: Много нитей, покрывать обращения одних нитей впамять вычислениями в других за счёт быстрогопереключения контекстаУтилизация латентности памяти GPU: Много нитей, покрывать обращения одних нитей впамять вычислениями в других за счёт быстрогопереключения контекста За счёт наличия сотен ядер и поддержки миллионов нитей(потребителей) на GPU легче утилизировать всю полосупропусканияИз чего состоит программа с использованием CUDA?Вычисления с использованием GPU Программа, использующая GPU, состоит из: Кода для GPU, описывающего необходимые вычисленияи работу с памятью устройства Кода для CPU, в котором осуществляетсяУправление памятью GPU – выделение / освобождениеОбмен данными между GPU/CPUЗапуск кода для GPUОбработка результатов и прочий последовательный кодВычисления с использованием GPU GPU рассматривается как периферийное устройство,управляемое центральным процессором GPU «пассивно», т.е.

не может само загрузить себяработой Код для GPU можно запускать из любого места программыкак обычную функцию «Точечная», «инкрементная» оптимизация программТерминология CPU Будем далее называть «хостом»(от англ. host ) код для CPU - код для хоста, «хост-код» (host-code ) GPU будем далее называть«устройством» или «девайсом»(от англ.device) код для GPU – «код для устройства»,«девайс-код» ( device-code ) Хост выполняет последовательный хост-код, в котором содержатся вызовыфункций, побочный эффект которых –манипуляции с устройством.Код для GPU (device-code) Код для GPU пишется на C++ с некоторыми надстройками: Атрибуты функций, переменных и структур Встроенные функцииМатематика, реализованная на GPUСинхронизации, коллективные операции Векторные типы данных Встроенные переменныеthreadIdx, blockIdx, gridDim, blockDim Шаблоны для работы с текстурами … Компилируется специальным компилятором ciccКод для CPU (host-code) Код для CPU дополняется вызовами специальных функцийдля работы с устройством Код для CPU компилируется обычным компилятором Кроме конструкции запуска ядра <<<...>>> Функции линкуются из динамических библиотекСложение векторовВектор AВектор BldldldldldРезультатldldldldldldstststldldstldldstldldstldldstldstststСложение векторов Без GPU:for (int i = 0; i < N; i++) {c[i] = a[i] + b[i];} С GPU{// на CPU:<Переслать данные с CPU на GPU>;<Запустить вычисления на N GPU-нитях>;<Скопировать результат с GPU на CPU>;}{// на GPU в нити с номером threadIndex:c[threadIndex] = a[theadIndex] + b[threadIndex];}SPMD & CUDA GPU работает по методу SPMD - единая программа,множество данных Задается программа (CUDA kernel) Запускается множество нитей (CUDA grid) Каждая нить выполняет копию программы над своимиданнымиCUDA Grid Хост может запускать на GPU множества виртуальных нитей Каждая нить приписана некоторому виртуальному блоку Грид (от англ.

Grid-сетка ) – множество блоков одинаковогоразмера Положение нити в блоке и блока в гриде индексируются потрём измерениям (x,y,z)CUDA Grid Грид задаётся количествомблоков по x,y,z (размер грида вблоках) и размерами каждогоблока по x,y,z Ели по z размер грида и блоковравен единице, то получаемплоскую прямоугольную сеткунитейCUDA Grid пример Двумерный грид из трёхмерных блоков Логический индекс по переменной z у всех блоковравен нулю Каждый блок состоит из трёх «слоёв» нитей,соответствующих z=0,1,2CUDA Kernel («Ядро») Нити выполняют копии т.н.

«ядер» - специальнооформленных функций, компилируемых под GPUНет возвращаемого значения (void)Атрибут __global____global__ void kernel (int * ptr) {ptr = ptr + 1;ptr[0] = 100;….; //other code for GPU}Терминология Хост запускает вычисление ядра на гриде нитейИногда «на гриде нитей» опускается Одно и то же ядро может быть запущено на разныхгридахЗапуск ядра kernel<<< execution configuration >>>(params); “kernel” – имя ядра, “params” – параметры ядра, копию которых получит каждая нить execution configuration:<<< dim3 gridDim, dim3 blockDim >>> dim3 - структура, определённая в CUDA Toolkitstruct dim3 {unsigned x,y,z;dim3(unsigned vx=1, unsigned vy=1, unsigned vz=1);}Запуск ядра kernel<<< execution configuration >>>(params); “kernel” – имя ядра, “params” – параметры ядра, копию которых получит каждая нить execution configuration:<<< dim3 gridDim, dim3 blockDim >>> dim3 gridDim - размеры грида в блокахчисло блоков = gridDim.x * gridDim.y * gridDim.z dim3 blockDim - размер каждого блокачисло нитей в блоке = blockDim.x * blockDim.y *blockDim.zЗапуск ядра• Рассчитать грид:dim3 blockDim = dim3(512);gridDim = dim3( (n – 1) /512 + 1• Запустить ядро с именем “kernel”kernel <<< gridDim, blockDim>>>(…);Ориентация нити в гриде Осуществляется за счёт встроенных переменных:dim3dim3dim3dim3threadIdxblockIdxblockDimgridDim- индексы нити в блоке- индексты блока в гриде- размеры блоков в нитях- размеры грида в блоках Линейный индекс нити в гриде:int gridSizeX = blockDim.x*gridDim.x;int gridSizeAll = gridSizeX * gridSizeY * gridSizeZint threadLinearIdx =(threaIdx.z * gridSizeY + threadIdx.y) * gridSizeX +threadIdx.xПример: ядро сложения__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; //записать результат суммирования} Каждая нить Получает копию параметров Рассчитывает свой элемент выходного массиваHost Code Выделить память на устройстве Переслать на устройство входные данные Рассчитать грид Размер грида зависит от размера задачи Запустить вычисления на гриде В конфигурации запуска указываем грид Переслать с устройства на хост результатВыделение памяти на устройстве cudaError_t cudaMalloc ( void** devPtr, size_t size ) Выделяет size байтов линейной памяти на устройстве и возвращаетуказатель на выделенную память в *devPtr.

Память не обнуляется.Адрес памяти выровнен по 512 байт cudaError_tcudaFree ( void* devPtr ) Освобождает память устройства на которую указывает devPtr. Вызов cudaMalloc(&p, N*sizeof(float)) соответствует вызову p= malloc(N*sizeof(float));Копирование памяти cudaError_tcudaMemcpy ( void* dst, const void* src,size_t count, cudaMemcpyKind kind ) Копирует count байтов из памяти, на которую указывает src в память,на которую указывает dst, kind указывает направление передачиcudaMemcpyHostToHost– копирование между двумя областямипамяти на хостеcudaMemcpyHostToDevice – копирование с хоста на устройствоcudaMemcpyDeviceToHost – копирование с устройства на хостcudaMemcpyDeviceToDevice – между двумя областями памятина устройстве Вызов cudaMemcpy() с kind, не соответствующим dst и src , приводит кнепредсказуемому поведениюПример кода хостаint n = getSize(); // размер задачиint nb = n * sizeof (float); // размер размер задачи в байтахПриходится дублировать указатели для хоста и GPUfloat *inputDataOnHost = getInputData(); // входные данные на хостеfloat *resultOnHost = (float *)malloc( nb );float *inputDataOnDevice = NULL, *resultOnDevice = NULL;cudaMalloc( (void**)& inputDataOnDevice, nb);cudaMalloc( (void**)& resultOnDevice, nb);Выделение памятина GPUПример кода хостаcudaMemcpy(inputDataOnDevice, inputDataOnHost,nb, cudaMemcpyHostToDevice);Копирование входныхданных на GPUdim3 blockDim = dim3(512);dim3 gridDim = dim3((n – 1) / 512 + 1 );kernel <<< gridDim, blockDim >>> (inputDataOnDevice,resultOnDevice, n);cudaMemcpy(resultOnHost, resultOnDevice,nb, cudaMemcpyDeviceToHost);cudaFree(inputDataOnDevice);cudaFree(resultOnDevice);ЗапускядраКопирование результатана хостОсвободить памятьКак реализовать выполнение миллионов нитей наимеющейся архитектуре?CUDA и классификация ФлиннаАрхитектура ЭВМSIMD – всепроцессыодновременновыполняют однуинструкциюSMP – все процессыимеют равныеправа на доступ кпамятиMIMD – каждыйпроцессвыполняетсянезависимо отдругих,MPPMISDNUMASISDcc-NUMACUDA и классификация Флинна У Nvidia собственная модель исполнения,имеющая черты как SIMD, так и MIMD: Nvidia SIMT: Single Instruction – Multiple ThreadSIMDNvidiaSIMTMIMD(SMP)SIMT: виртуальные нити, блоки Виртуально все нити: выполняются параллельно (MIMD) Имеют одинаковые права на доступ к памяти (MIMD :SMP) Нити разделены на группы одинакового размера (блоки): В общем случае (есть исключение) , глобальнаясинхронизация всех нитей невозможна, нити из разныхблоков выполняются полностью независимо Есть локальная синхронизация внутри блока, нити из одногоблока могут взаимодействовать через специальную память Нити не мигрируют между блоками.