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

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

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

Лихогруд Николай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) Нити разделены на группы одинакового размера (блоки): В общем случае (есть исключение) , глобальнаясинхронизация всех нитей невозможна, нити из разныхблоков выполняются полностью независимо Есть локальная синхронизация внутри блока, нити из одногоблока могут взаимодействовать через специальную память Нити не мигрируют между блоками.

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

Тип файла PDF

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

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

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