Лекция. CUDA 1 (Колганов) (1186105)

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

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

Колганов Александр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); // переслать на устройствовходные данные//запустить ядро.

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

Тип файла PDF

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

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

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