Технологии высокопроизводительных вычислений на GPU и гибридных

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

2. Преимущества и недостатки SISD архитектуры Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1

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

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

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

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

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

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

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

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

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

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

13. 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); Копирование данных из RAM впамять GPU и обратно Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1

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

15. __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 для сложения двух векторов: начало Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1

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

17. #include <stdio.h> 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; } Получение информации о GPU Технологии высокопроизводительных вычислений на GPU и гибридных вычислительных системах - Лекция 1