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

Колганов Александрalexander.k.s@mail.ruчасть 1Введение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,позволяющая производить вычисления сиспользованием графических процессоровРост CUDA10x growth in GPU computing!Ускорители в рейтинге TOP500Основные преимущества GPUпо сравнению с CPU Высокое соотношение цена / производительность Высокое соотношение производительность /энергопотреблениеРейтинг Green500Эффективность растетПрограммно-аппаратная модель:архитектура GPU NVidiaСемейства GPU устройствCPU Intel Core i7 Небольшое число мощныхнезависимых ядер;До 22 ядер, ~3.0 ГГц каждое;Поддержка виртуальных потоков(Hyper-Threading)3х уровневый кеш, большой кешL3 до 50 МБ;На каждое ядро L1=32KB (data) +32KB ( Instructions), L2=256KB;Обращения в памятьобрабатываются отдельно длякаждого процесса\нитиCore i7 3960x,6 ядер, 15MB L3GPU StreamingMultiprocessor (SM) Потоковый мультипроцессор «Единица» построенияустройства (как ядро в CPU): 32 скалярных ядра CUDA Core,~1.5ГГц2 Warp SchedulerФайл регистров, 128KB2х уровневый кэшТекстурные юниты16 x Special Function Unit (SFU) –интерполяция итрансцендентная математикаодинарной точности16 x Load/StoreGPC - Graphics Processing Cluster Объединение потоковых мультипроцессоров в блокиЧип поколения Fermi вмаксимальной конфигурации 16 SM 512 ядер CUDACore Кеш L2 758KB Контроллерыпамяти GDDR5 Интерфейс PCI2.0Вычислительная мощностьПропускная способность памятиСравнение GPU и CPU Сотни упрощённых вычислительных ядер, работающих нанебольшой тактовой частоте ~1.8 ГГц; Небольшие кеши на GPU 32 CUDA-ядра разделяют 64 КБ L1 L2 общий для всех CUDA ядер 2 MB, L3 отсутствует Оперативная память с высокой пропускной способностью ивысокой латентностью, оптимизированная дляколлективного доступа; Поддержка миллионов виртуальных нитей, быстроепереключение контекста для групп нитей.Утилизация латентности памяти Цель: эффективно загружать CUDA-ядра Проблема: латентность памяти Решение: CPU: Сложная иерархия кешей; 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 Шаблоны для работы с текстурами… Компилируется специальным компилятором nvccКод для CPU (host-code) Код для CPU дополняется вызовами специальныхфункций для работы с устройством; Код для CPU компилируется обычным компиляторомgcc/icc/cl; Кроме конструкции запуска ядра <<<...>>> ! Функции для GPU линкуются из динамическихбиблиотекСложение векторовСложение векторов Без GPU:for(int i=0; i<N; ++i)c[i]= a[i] + b[i]; С GPU:{ //на CPU<Переслать данные с CPU на GPU>;<Запустить вычисления на N GPU-нитях>;<Скопировать результат с GPU на CPU>;}{//в нити с номером IDXc[IDX] = a[IDX] + b[IDX];}CUDA Grid Хост может запускать на GPU множества виртуальныхнитей; Каждая нить приписана некоторому виртуальномублоку; Грид (от англ.

Grid-сетка ) – множество блоководинакового размера; Положение нити в блоке и блока в гридеиндексируются по трём измерениям (x, y, z).CUDA Grid Грид задаётся количеством блоков по [X, Y, Z] (размергрида в блоках) и размерами каждого блока по [X, Y, Z]; Например, еcли по Z размер грида и блоков равенединице, то получаем плоскую прямоугольную сеткунитей.CUDA Grid пример Двумерный грид из трёхмерных блоков Логический индекс по переменной z у всех блоковравен нулю; Каждый блок состоит из трёх «слоёв» нитей,соответствующих z=0, 1, 2.CUDA Kernel («Ядро») Каждая нить выполняет копию специальнооформленных функций «ядер», компилируемых дляGPU. Нет возвращаемого значения (void); Обязательный атрибут __global__.__global__ void kernel (int * ptr){ptr = ptr + 1;ptr[0] = 100; ….; //other code for GPU}CUDA Kernel («Ядро») Терминология: Хост запускает вычисление ядра на гриде нитей (либопросто хост запускает ядро на GPU). Одно и то же ядро может быть запущено на разныхгридах «Ядро» – что делать «Грид» – сколько делатьЗапуск ядра kernel<<< execution configuration >>>(params); “kernel” – имя ядра, “params” – параметры ядра, копию которых получит каждая нить execution configuration - Dg, Db, Ns, S dim3 Dg - размеры грида в блоках, Dg.x * Dg.y * Dg.z число блоков dim3 Db - размер каждого блока, Db.x * Db.y * Db.z - число нитей в блоке size_t Ns – размер динамически выделяемой общей памяти (опционально) cudaStream_t S - поток, в котором следует запустить ядро (опционально) struct dim3 – стуктура, определённая в CUDA Toolkit, Три поля: unsigned x,y,z Конструктор dim3(unsigned x=1, unsigned y=1, unsigned z=1)Ориентация нити в гриде Осуществляется за счёт встроенных переменных: 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 threadLinearIdx =blockIdx.x * blockDim.x + threadIdx.x; //определить свой индексint elemA = A[threadLinearIdx]; //считать нужный элемент Aint elemB = B[threadLinearIdx]; // считать нужный элемент BC[threadLinearIdx] = elemA + elemB; //записать результат суммирования} Каждая нить Получает копию параметров (В данном случае, это адресавектором на GPU); Определяет своё положение в гриде threadLinearIdx ; Cчитывает из входных векторов элементы с индексомthreadLinearIdx и записывает их сумму в выходной вектор поиндексу threadLinearIdx ; рассчитывает один элемент выходного массива.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, constvoid* 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); // размер размер задачи в байтах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); // переслать на устройствовходные данные//запустить ядро.