Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах Аксёнов Сергей Владимирович к.т.н., доцент каф.ОСУ ТПУ Томский политехнический университет

Преимущества и недостатки SISD архитектуры 2 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Преимущества Недостатки Последовательное исполнение инструкций Применение конвейеров Общая память данных Кэширование обращений к памяти Применение отработанных оптимизирующих компиляторов Блочное чтение памяти Программирование е требует особых знаний об архитектуре Невозможность наращивания тактовой частоты современных CPU

GPU (Graphics Processing Unit) архитектура VS SMP Кэш и много- функциональные АЛУ 3 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Простые АЛУ, имеющие общую память Преимущества GPU: Однократная систематическая задержка между потоком исходных данных и результатов. Недостаток GPU: Тщательная разработка алгоритма на управляющих элементах

GPU: Основные положения 4 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Является сопроцессором к CPU Обладает собственной памятью Обладает возможностью параллельного выполнения огромного числа нитей Отличия между нитями на CPU и GPU: 1.Создание, управление и удаление нитей на GPU – низкозатратный процесс. Те же самые операции с нитями на CPU достаточно ресурсоёмки. 2. Число нитей на CPU зависит от числа ядер и крайне мало, количество нитей на GPU достаточно большое (несколько тысяч или десятков тысяч)

Работа с нитями GPU 5 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 GPU обладает рядом потоковых мультипроцессоров (Streaming Multiprocessor), каждый из которых способен одновременно выполнять 768 (1024) нитей. Все мультипроцессоры работают независимо друг от друга. Нити разбиваются на группы по 32 нити (warp ). Нити в пределах каждой группы выполняются физически одновременно. Нити в разных группах могут находится на разных стадиях выполнения программ. Управление группой выполняет GPU. Обычно каждой нити соответствует один элемент вычисляемых данных.

Иерархия нитей в CUDA 6 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Ядро – функция, которая выполняется каждой нитью Верхний уровень иерархии – сетка (все нити, выполняющие ядро). Верхний уровень – одномерный или двухмерный массив блоков. Каждый блок – 1D, 2D или 3D массив нитей. Все блоки обладают одинаковой размерностью и размером.

Встроенные переменные CUDA 7 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 threadIdx – Идентификатор нити в блоке blockIdx – Идентификатор блока в сетке gridDim – размер сетки (количество блоков в сетке) blockDim – размер блока (количество нитей в блоке) Указанные переменные – трехмерный целочисленный вектор. Доступны только для функций, выполняемых на GPU. Разбиение нитей на группы происходит отдельно для каждого блока. Нити могут взаимодействовать между собой только в пределах блока. Пример (для одномерного случая) int globalIdx = threadIdx.x+blockIdx.x*blockDim.x;

Пример ядра: сложение 1D массивов 8 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 __global__ void sumOfVectors(float *a, float *b, float *c) { int index = blockIdx.x*blockDim.x + threadIdx.x; c[index] = a[index] + b[index]; } Каждая нить находит вычисляет только один элемент массива с. Index определяет номер нити и индекс вычисляемого элемента массива.

Взаимодействие нитей внутри блока 9 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Разделяемая память (быстрая память 16Кбайт, которую нити блока могут совместно использовать). Барьерная синхронизация (синхронизация нитей блока). Вызов функции __synchthreads(). Блокировка вызывающих нитей блока до тех пор, пока все нити не войдут в эту функцию.

Спецификаторы функций 10 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 __global__ определяет ядро и эту функция возвращает результат void Невозможно взять адрес __device__ функции Функции, выполняемые на устройстве, не допускают рекурсию, объявление статических переменных Пример __global__ void KernelFunction(float a) Квалификатор Выполняется на Вызывается на __device__device __host__host __global__devicehost

Структура CUDA-программы 11 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 1. Выделение памяти на GPU 2. Копирование данных из RAM в выделенную память GPU 3. Запуск ядра 4. Копирование результатов из памяти GPU в RAM 5. Освобождение памяти GPU

Выделение и освобождение памяти GPU 12 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 Выделение памяти на GPU cudaMalloc(void ** devPtr, size_t size) devPtr – Указатель на память size – Размер выделяемой памяти в байтах Освобождение памяти GPU cudaFree(void *devPtr) devPtr – Указатель на память для освобождения

Копирование данных из RAM в память GPU и обратно 13 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 cudaMemcpy(void *dest, const void *src, size_t size, enum cudaMemcpyKind kind) dest – Адрес памяти приёма данных src – Адрес памяти отправления данных size – Размер копируемых данных в байтах kind – Направление копирования {cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice}. Примеры cudaMemcpy(aDev, aRAM, 1024, cudaMemcpyHostToDevice); cudaMemcpy(bRAM, cDev, 2048, cudaMemcpyDeviceToHost);

Вызов ядра 14 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 kernelName >>(args) kernelName – Имя __global__ функции. Dg – размерность и размер сетки в блоках Db – размерность и размер блока в нитях Ns – размер дополнительно выделяемой разделяемой памяти (опционально) S- поток СUDA, в котором должен произойти вызов, по умолчанию используется поток 0 (опционально). args – список аргументов функции kernelName

Пример программы GPU для сложения двух векторов: начало 15 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 __global__ void sumOfVectors(float *a, float *b, float *c) { int index = blockIdx.x*blockDim.x + threadIdx.x; c[index] = a[index] + b[index]; } void main(float *a, float *b, float *c) { int numBytes = n*sizeof(float); float *aDev, *bDev, *cDev; cudaMalloc((void**)&aDev, numBytes); cudaMalloc((void**)&bDev, numBytes); cudaMalloc((void**)&cDev, numBytes);

Пример программы GPU для сложения двух векторов: окончание 16 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 //Определение конфигурации сетки dim3 threads = dim3(512, 1); dim3 blocks = dim3(n/threads.x, 1); cudaMemcpy(aDev, a, numBytes, cudaMemcpyHostToDevice); cudaMemcpy(bDev, b, numBytes, cudaMemcpyHostToDevice); sumOfVectors >>(aDev, bDev, cDev); cudaMemcpy(c, cDev, numBytes, cudaMemcpyDeviceToHost); cudaFree(aDev); cudaFree(bDev); cudaFree(cDev); }

Получение информации о GPU 17 Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1 #include int main ( int argc, char * argv [] ) { int deviceCount; cudaDeviceProp devProp; cudaGetDeviceCount ( &deviceCount ); // Определение числа устройств GPU printf ( "Found %d devices\n", deviceCount ); for ( int device = 0; device < deviceCount; device++ ) { cudaGetDeviceProperties ( &devProp, device ); // Получение характеристик GPU printf ( "Device %d\n", device ); printf ( "Compute capability : %d.%d\n", devProp.major, devProp.minor ); printf ( "Name : %s\n", devProp.name ); printf ( "Total Global Memory : %d\n", devProp.totalGlobalMem ); printf ( "Shared memory per block: %d\n", devProp.sharedMemPerBlock ); printf ( "Registers per block : %d\n", devProp.regsPerBlock ); printf ( "Warp size : %d\n", devProp.warpSize ); printf ( "Max threads per block : %d\n", devProp.maxThreadsPerBlock ); printf ( "Total constant memory : %d\n", devProp.totalConstMem ); } return 0; }