Основы технологии CUDA (Compute Unified Device Architecture) Powerpoint

Основы технологии CUDA (Compute Unified Device Architecture) Powerpoint Templates

Операции с плавающей точкой в секунду для центрального и графического процессоров Powerpoint Templates

Сравнение архитектур CPU и GPU Параллельная обработка данных распределяет элементы данных на параллельно обрабатываемых потоках. GPU особенно хорошо подходит для решения проблем, которые могут быть выражены как вычислений данными параллельно - та же программа выполняется на многих элементов данных параллельно - с высокой интенсивностью - арифметическое отношение арифметических операций к операциям с Powerpoint Templates памятью.

CUDA предназначен для поддержки различных языков и интерфейсы прикладного программирования Powerpoint Templates

Модель программирования Ключевые абстракции: - hierarchy of thread groups (иерархия групп потоков (нитей)), - shared memories (разделяемая память), - barrier synchronization (барьерная синхронизация. Powerpoint Templates

Вычислительная модель GPU Powerpoint Templates

Структура блоков Powerpoint Templates

Масштабируемость Каждый блок потоков может быть распределён на любом количестве процессорных ядер, в любом порядке, параллельно или последовательно, так что скомпилированная CUDA-программа может выполняться на любом количестве процессорных ядер Powerpoint Templates

Модель программирования Powerpoint Templates

Ядра (Kernel ) CUDA Cи основан на Си, позволяет программисту использовать функции Си, вызывать программные ядра, которые при запуске выполняются N раз параллельно на N разных CUDA потоках. // Определение ядра __global__ void Vec. Add(float* A, float* B, float* C) { int i = thread. Idx. x; C[i] = A[i] + B[i]; Каждому потоку, который выполняется на ядре, дан уникальный } идентификационный номер, к которому есть доступ в ядре с помощью int main() встроенной thread. Idx переменной. { . . . // Вызов ядра с N нитями Vec. Add<<<1, N>>>(A, B, C); } Powerpoint Templates

Иерархия потоков Тhread. Idx является 3 -компонентным вектором, поэтому потоки могут быть идентифицированы с помощью одномерного, двумерного или трёхмерного индекса потока, формируя одномерный, двумерный или трёхмерный блок потока. Это обеспечивает способ запуска вычислений через элементы в домене, такие как вектор, матрица или блок. Индекс потока и его ID имеют прямую зависимость друг от друга: Для одномерного блока они одинаковы; для двумерного блока размера (Dx, Dy), для потока с индексом (x, y), ID потока – (x+y. Dx); для трёхмерного блока размера (Dx, Dy, Dz), для потока с индексом (x, y, z), ID потока – (x+y. Dx+z. Dx. Dy). Существует лимит количества потоков в одном блоке, так как все потоки в блоке могут находиться на одном процессорном ядре и должны совместно использовать ограниченные ресурсы памяти этого ядра. На современных GPU блок потоков может содержать до 1024 потоков. Однако ядро может выполняться на многих эквивалентных блоках потоков, поэтому общее число потоков равно числу потоков на Powerpoint Templates блок, умноженному на число блоков.

Сложение матриц А и В размера Nx. N // Определение ядра __global__ void Mat. Add(float A[N][N], float B[N][N], float C[N][N]) { int i = thread. Idx. x; int j = thread. Idx. y; C[i][j] = A[i][j] + B[i][j]; } int main() { . . . // Вызов ядра с блоком N * 1 нитями int num. Blocks = 1; dim 3 threads. Per. Block(N, N); Mat. Add<< >>(A, B, C); } Powerpoint Templates

Пример обработки несколькими блоками // Kernel definition __global__ void Mat. Add(float A[N][N], float B[N][N], float C[N][N]) { int i = block. Idx. x * block. Dim. x + thread. Idx. x; int j = block. Idx. y * block. Dim. y + thread. Idx. y; if (i < N && j < N) C[i][j] = A[i][j] + B[i][j]; } int main() { . . . // Kernel invocation dim 3 threads. Per. Block(16, 16); dim 3 num. Blocks(N / threads. Per. Block. x, N / threads. Per. Block. y); Mat. Add<<>>(A, B, C); . . . } Powerpoint Templates

Иерархия памяти Потоки CUDA могут получить доступ к данным из нескольких областей памяти. Powerpoint Templates

Гетерогенное программирование Модель программирования CUDA предполагает, что потоки CUDA можно выполнить на физически отдельных устройствах. Это имеет место, например, когда ядра выполняются на GPU, а остальные программы выполняет центральный процессор. Это включает в себя распределение памяти устройства и освобождения, а также передачу данных между хостом и памяти устройства. Powerpoint Templates

Типы памяти Powerpoint Templates

Типы памяти 1. Регистры. 2. Локальная память. 3. Глобальная память. 4. Разделяемая память. 5. Константная память. 6. Текстурная память Тип памяти Доступ Уровень выделения Скорость работы регистры R/W per-thread высокая (on chip) local R/W per-thread низкая (DRAM) shared R/W per-block высокая (on-chip) global R/W per-grid низкая(DRAM) высокая(on chip L 1 constant R/O per-grid cache) высокая(on chip L 1 texture R/O per-grid cache) Powerpoint Templates

Регистры. Компилятор по возможности размещает все локальные переменные функций в регистрах. Доступ к таким переменным осуществляется с максимальной скоростью. Для того чтобы рассчитать количество регистров, доступных одной нити GPU, необходимо разделить общее число регистров на произведение количества нитей в блоке и количества блоков в гриде. Все регистры GPU 32 разрядные. Компилятор старается максимизировать число одновременно работающих блоков. Powerpoint Templates

Локальная память В случаях, когда локальные данные процедур занимают слишком большой размер, или компилятор не может вычислить для них некоторый постоянный шаг при обращении, он может поместить их в локальную память. По скоростным характеристикам локальная память работает намного медленнее , чем регистровая. В какой-либо функции она может быть использована компилятором при большем количестве локальных переменных. Ее рекомендуется использовать при необходимости. Физически локальная память является аналогом глобальной памяти. Powerpoint Templates

Глобальная память Одним из достижений технологии является возможность произвольность адресации глобальной памяти, но одновременно с тем же является самым медленным типом памяти. Глобальная память не кэшируется. В основном она необходима для хранения результатов работы программы. Для размещения данных в глобальной памяти используется спецификатор __global__. Распределением памяти должна заниматься программа-хост, работающая на CPU при помощи cuda. Malloc(void* mem, int size). Powerpoint Templates

Разделяемая память Относиться к быстрому типу памяти. Разделяемую память рекомендуется использовать для минимизации обращение к глобальной памяти, а так же для хранения локальных переменных функций. Адресация разделяемой памяти между нитями потока одинакова в пределах одного блока, что может быть использовано для обмена данными между потоками в пределах одного блока. Для размещения данных в разделяемой памяти используется спецификатор __shared__. Гарантируется, что во время исполнения блока на мультипроцессоре содержимое разделяемой памяти будет сохраняться. Однако после того как на мультипроцессоре сменился блок, не гарантируется, что содержимое старого блока сохранилось. Разделяемая память — это некэшируемая, но быстрая память. Ее и рекомендуется использовать как управляемый кэш. На один мультипроцессор доступно всего 16 KB разделяемой памяти. Разделив это число на количество задач в блоке, получим максимальное количество разделяемой памяти, доступной на один Powerpoint Templates поток.

Константная память Является достаточно быстрым типом памяти. Она кэшируется. Кэш существует в единственном экземпляре для одного мультипроцессора, следовательно, общий для всех задач внутри блока. В константной памяти можно размещать данные любого типа и читать их при помощи простого присваивания. Для размещения данных в константной памяти предусмотрен спецификатор __constant__. Динамическое выделение в отличие от глобальной памяти в константной не поддерживается, если необходимо использовать массив в константной памяти, то его размер необходимо указать заранее. Ее размер составляет всего 64 Kбайт (на все устройство). Powerpoint Templates

Текстурная память Предназначена главным образом для работы с текстурами. Она оптимизирована под выборку 2 D данных и имеет следующие возможности: • быстрая выборка значений фиксированного размера из одномерного или двухмерного массива; • нормализованная адресация числами типа float в интервале [0, 1). • аппаратная линейная или билинейная интерполяция соседних значений в случае нормализованной адресации; • аппаратная обработка выхода за границу массива с использованием двух режимов: clamp и wrap. Размер текстурной памяти ограничивается только максимальным размером памяти, которую может выделить устройство. Но так же из текстурной памяти можно читать данные только встроенных в nvcc типов, имеющих размер 1, 2, 4, 8 или 16 байт, и только с помощью специальных функций — tex 1 D, tex 2 D или tex 1 Dfetch, tex 2 Dfetch. Другими словами, нельзя сделать указатель на текстурную память и переименовать его произвольным образом. Powerpoint Templates

Организация памяти устройства Powerpoint Templates

Расширение языка Си Powerpoint Templates

Расширение языка Си • Спецификаторы функций, которые показывают, как и откуда будут выполняться функции. • Спецификаторы переменных, которые служат для указания типа используемой памяти GPU. • Спецификаторы запуска ядра GPU. • Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU. • Дополнительные типы переменных. Powerpoint Templates

Спецификаторы функций определяют, как и откуда буду вызываться функции: • __host__ — выполнятся на CPU, вызывается с CPU. • __global__ — выполняется на GPU, вызывается с CPU. • __device__ — выполняется на GPU, вызывается с GPU. Powerpoint Templates

Добавлены: Добавленные переменные В язык добавлены следующие специальные переменные grid. Dim – размер grid’а (имеет тип dim 3) block. Dim – размер блока (имеет тип dim 3) block. Idx – индекс текущего блока в grid’е (имеет тип uint 3) thread. Idx – индекс текущей нити в блоке (имеет тип uint 3) warp. Size – размер warp’а (имеет тип int) Добавленные типы В язык добавляются 1/2/3/4 -мерные вектора из базовых типов — char 1, char 2, char 3, char 4, uchar 1, uchar 2, uchar 3, uchar 4, short 1, short 2, short 3, short 4, ushort 1, ushort 2, ushort 3, ushort 4, int 1, int 2, int 3, int 4, uint 1, uint 2, uint 3, uint 4, long 1, long 2, long 3, long 4, ulong 1, ulong 2, ulong 3, ulong 4, float 1, float 2, float 3, float 2, и double 2. Обращение к компонентам вектора идет по именам - x, y, z и w. Для создания значений-векторов заданного типа служит конструкция вида make_. Powerpoint Templates

Добавленные функции Для GPU операции с double-числами выполняются медленнее, чем операции с float-числами, то предпочтительнее там, где это возможно, использовать float-аналоги стандартных функций. Так, float-аналогом функции sin является функция sinf. Кроме того, CUDA предоставляет также специальный набор функций пониженной точности, но обеспечивающих еще большее быстродействие. Таким аналогом для функции синусы является функция __sinf. Для ряда функций можно задать требуемый способ округления. Используемый способ задается при помощи одного из следующих суффиксов: rn – округление к ближайшему; rz - округление к нулю; ru – округление вверх; rd – округление вниз. Powerpoint Templates

Powerpoint Templates

Целочисленные функции в CUDA Функция Значение lstinline__[u]mul 24(x, y) Вычисляет произведение младших 24 бит целочисленных параметров x и y, возвращает младшие 32 бита результата. Старшие 8 бит аргументов игнорируются lstinline__[u]mulhi(x, y) Возвращает старшие 32 бита произведения целочисленных операндов x и y lstinline__[u]mul 64 hi(x, y) Вычисляет произведение 64 -битовых целых чисел и возвращает младшие 64 бита этого произведения lstinline__[u]sad(x, y, z) Возвращает z + |x - y| lstinline__clz(x) Возвращает целое число от 0 до 32 включительно последовательных нулевых битов для целочисленного параметра x, начиная со старших бит lstinline__clzll(x) Возвращает целое число от 0 до 64 включительно последовательных нулевых битов для целочисленного 64 -битового параметра x, начиная со старших бит lstinline__ffs(x) Возвращает позицию первого (наименее значимого) единичного бита для аргумента x. Если x равен нулю, то возвращается нуль lstinline__ffsll(x) Возвращает позицию первого (наименее значимого) единичного бита для целочисленного 64 -битового аргумента x. Если x равен нулю, то возвращается нуль lstinline__popc(c) Возвращает число бит, которые равны единице в двоичном представлении 32 – битового целочисленного аргумента x lstinline__popcll(x) Возвращает число бит, которые равны единице в двоичном представлении 64 - битового целочисленного аргумента x lstinline__brev(x) Возвращает число, полученное перестановкой (то есть биты в позициях k и 31 -k меняются местами для всех k от 0 до 31) битов исходного 32 -битового целочисленного аргумента x lstinline__brevll(x) Возвращает число, полученное перестановкой (то есть биты в позициях k и 63 -k Powerpoint Templates меняются местами для всех k от 0 до 63) битов исходного 64 -битового целочисленного аргумента x

Последовательность выполнения программы с использованием CUDA: • Получить данные для расчетов. • Скопировать эти данные в GPU память. • Произвести вычисление в GPU • Скопировать вычисленные данные из GPU памяти в ОЗУ. • Вывести результаты. • Высвободить используемые ресурсы. Powerpoint Templates

Копирование данных в память GPU 1. Выделение памяти - функция cuda. Malloc: cuda. Error_t cuda. Malloc( void** dev. Ptr, size_t count ), где dev. Ptr – указатель, в который записывается адрес выделенной памяти, count – размер выделяемой памяти в байтах. Возвращает: cuda. Success – при удачном выделении памяти cuda. Error. Memory. Allocation – при ошибке выделения памяти Пример: cuda. Malloc((void**)&dev. Vec 1, sizeof(float) * SIZE); Powerpoint Templates

Копирование в память GPU и обратно 2. Для копирования - функция cuda. Memcpy: cuda. Error_t cuda. Memcpy(void* dst, const void* src , size_t count, enum cuda. Memcpy. Kind), где dst – указатель, содержащий адрес места-назначения копирования, src – указатель, содержащий адрес источника копирования, count – размер копируемого ресурса в байтах, cuda. Memcpy. Kind – перечисление, указывающее направление копирования (может быть cuda. Memcpy. Host. To. Device, cuda. Memcpy. Device. To. Host, cuda. Memcpy. Host. To. Host, cuda. Memcpy. Device. To. Device). Возвращает: cuda. Success – при удачном копировании cuda. Error. Invalid. Value – неверные параметры аргумента (например, размер копирования отрицателен) cuda. Error. Invalid. Device. Pointer – неверный указатель памяти в видеокарте cuda. Error. Invalid. Memcpy. Direction – неверное направление (например, перепутан источник и место-назначение копирования) Пример: cuda. Memcpy(dev. Vec 1, vec 1, sizeof(float) * SIZE, cuda. Memcpy. Host. To. Device); Powerpoint Templates

Вычисления в GPU Синтаксис запуска ядра имеет следующий вид: my. Kernel. Func<<>>(float* param 1, float * param 2), где my. Kernel. Func — функция ядра (спецификатор __global__) grid. Size — размерность сетки блоков (dim 3), выделенную для расчетов, block. Size — размер блока (dim 3), выделенного для расчетов, shared. Mem. Size — размер дополнительной памяти, выделяемой при запуске ядра, cuda. Stream — переменная cuda. Stream_t, задающая поток, в котором будет произведен вызов. Некоторые переменные при вызове ядра можно опускать( shared. Mem. Size и cuda. Stream). cuda. Device. Synchronize() – функция используется для синхронизации потоков. Пример: add. Kernel<<<1, size>>>(dev_c, dev_a, dev_b); cuda. Device. Synchronize(); Powerpoint Templates

Освобождение используемых ресурсов Прототип: cuda. Error_t cuda. Free(void *dev. Ptr), где *dev. Ptr — указатель, в который записывается адрес выделенной памяти Пример: cuda. Free(dev_c); Powerpoint Templates

Пример программы 1 Powerpoint Templates

#include " src="https://present5.com/presentation/3/-128262303_438454940.pdf-img/-128262303_438454940.pdf-38.jpg" alt=" Пример 1 #include "cuda_runtime. h" #include "device_launch_parameters. h" #include #include " /> Пример 1 #include "cuda_runtime. h" #include "device_launch_parameters. h" #include #include #include int main() { Set. Console. CP(1251); Set. Console. Output. CP(1251); int device. Count; cuda. Get. Device. Count(&device. Count); for(int device = 0; device < device. Count; device++) { cuda. Device. Prop device. Prop; cuda. Get. Device. Properties(&device. Prop, device); Powerpoint Templates

printf("Номер устройства: %dn", device); printf("Имя устройства: %sn", device. Prop. name); printf("Объем глобальной памяти: %d Мбайтn", device. Prop. total. Global. Mem/1024); printf("Объем shared-памяти в блоке : %dn", device. Prop. shared. Mem. Per. Block); printf("Объем регистровой памяти: %dn", device. Prop. regs. Per. Block); printf("Размер warp'a: %dn", device. Prop. warp. Size); printf("Размер шага памяти: %dn", device. Prop. mem. Pitch); printf("Макс количество потоков в блоке: %dn", device. Prop. max. Threads. Per. Block); printf("Максимальная размерность потока: x = %d, y = %d, z = %dn", device. Prop. max. Threads. Dim[0], device. Prop. max. Threads. Dim[1], device. Prop. max. Threads. Dim[2]); printf("Максимальный размер сетки: x = %d, y = %d, z = %dn", device. Prop. max. Grid. Size[0], device. Prop. max. Grid. Size[1], device. Prop. max. Grid. Size[2]); Powerpoint Templates

printf("Тактовая частота: %d к. Гцn", device. Prop. clock. Rate); printf("Общий объем константной памяти: %dn", device. Prop. total. Const. Mem); printf("Вычислительная мощность: %d. %dn", device. Prop. major, device. Prop. minor); printf("Величина текстурного выравнивания : %dn", device. Prop. texture. Alignment); printf("Количество процессоров: %dn", device. Prop. multi. Processor. Count); system("Pause"); return 0; } } Powerpoint Templates

Пример 2 Сложение элементов массивов. Powerpoint Templates

cuda. Error_t add." src="https://present5.com/presentation/3/-128262303_438454940.pdf-img/-128262303_438454940.pdf-42.jpg" alt=" Часть1 #include "cuda_runtime. h" #include "device_launch_parameters. h" #include cuda. Error_t add." /> Часть1 #include "cuda_runtime. h" #include "device_launch_parameters. h" #include cuda. Error_t add. With. Cuda(int *c, const int *a, const int *b, unsigned int size); __global__ void add. Kernel(int *c, const int *a, const int *b) { int i = thread. Idx. x; c[i] = a[i] + b[i]; } Powerpoint Templates

Часть 2 int main() { const int array. Size = 5; const int a[array. Size] = { 1, 2, 3, 4, 5 }; const int b[array. Size] = { 10, 20, 30, 40, 50 }; int c[array. Size] = { 0 }; cuda. Error_t add. With. Cuda(c, a, b, array. Size); printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d}n", c[0], c[1], c[2], c[3], c[4]); . cuda. Device. Reset(); return 0; } Powerpoint Templates

Часть 3 cuda. Error_t add. With. Cuda(int *c, const int *a, const int *b, unsigned int size) { int *dev_a = 0; int *dev_b = 0; int *dev_c = 0; cuda. Set. Device(0); cuda. Malloc((void**)&dev_c, size * sizeof(int)); cuda. Malloc((void**)&dev_a, size * sizeof(int)); cuda. Malloc((void**)&dev_b, size * sizeof(int)); cuda. Memcpy(dev_a, a, size * sizeof(int), cuda. Memcpy. Host. To. Device); cuda. Memcpy(dev_b, b, size * sizeof(int), cuda. Memcpy. Host. To. Device); add. Kernel<<<1, size>>>(dev_c, dev_a, dev_b); cuda. Device. Synchronize(); cuda. Memcpy(c, dev_c, size * sizeof(int), cuda. Memcpy. Device. To. Host); cuda. Free(dev_c); cuda. Free(dev_a); cuda. Free(dev_b); } Powerpoint Templates

Синхронизация Powerpoint Templates

Функция __syncthreads() обеспечивает синхронизацию потоков в блоке, которая будет ждать до тех пор, пока все запущенные потоки отработают до этой точки. Функция необходима для данных обрабатываемых одним потоком, затем будут использоваться другими потоками. __syncthreads() позволяет устанавливать точку синхронизации как барьер, и до того момента как все нити блока не закончат выполнение всех предыдущих инструкций, ни один последующий код не будет выполняться. Ее не рекомендуется использовать внутри условий. Powerpoint Templates