Часть 1 Технология CUDA Раба Н.О. аспирант кафедры

Часть 1 Технология CUDA Раба Н.О. аспирант кафедры информатики e-mail: [email protected]

Примеры использования GPGPU (GPGPU – General-Purpose computing on Graphics Processing Units) Научные расчеты Перекодирование, проигрывание, монтаж видео Photoshop, Premiere, AutoCAD, MATLAB, Mathematica и др.

Сравнение CPU и GPU

Сравнение CPU и GPU Пиковая производительность Пропускная способность памяти

GPU

GPU

Потоковый мультипроцессор (SM – Streaming Multiprocessor) Ядра (cores) Модули для вычисления специальных функций (SFU) Регистры Разделяемая память/L1 кэш

Технологии использования видеокарт для вычислений общего назначения CUDA (Compute Unified Device Architecture ) OpenCL (Open Computing Language) DirectCompute ATI Stream

CUDA C/C++ с расширениями и ограничениями задействует и CPU (host) и GPU (device) использует потоки (threads) компилятор nvcc Ядро (kernel)

Схема работы CUDA программы Выделение памяти GPU Копирование данных из основной памяти в память GPU Запуск ядра (kernel) или нескольких ядер Копирование результатов из памяти GPU в основную память Освобождение выделенной памяти

Потоки Сетка (grid): gridDim Блок (block): blockIdx, blockDim Поток (thread): threadIdx

Потоки gridDim.x = 3 gridDim.y = 2 gridDim.z = 1 blockIdx.x = 1 blockIdx.y = 1 blockIdx.z = 0 blockDim.x = 4 blockDim.y = 3 blockDim.z = 1 threadIdx.x = 3 threadIdx.y = 1 threadIdx.z = 0

Сложение матриц void sum(int* a, int* b, int* res, int n) { for (int i = 0; i < n; i++) for (int j = 0; j < n; j++) { int index = i + j * n; res[index] = a[index] + b[index]; } }

Сложение матриц void sum(int* a, int* b, int* res, int n) { int* d_a; int* d_b; int* d_res; int size = n * n * sizeof(int); cudaMalloc((void**)&d_a, size) cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_res, size); cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y); sumKernel<<>>(d_a, d_b, d_res, n); cudaMemcpy(res, d_res, size, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_res); }

Сложение матриц void sum(int* a, int* b, int* res, int n) { int* d_a; int* d_b; int* d_res; int size = n * n * sizeof(int); cudaMalloc((void**)&d_a, size) cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_res, size); cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y); sumKernel<<>>(d_a, d_b, d_res, n); cudaMemcpy(res, d_res, size, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_res); }

Сложение матриц void sum(int* a, int* b, int* res, int n) { int* d_a; int* d_b; int* d_res; int size = n * n * sizeof(int); cudaMalloc((void**)&d_a, size) cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_res, size); cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y); sumKernel<<>>(d_a, d_b, d_res, n); cudaMemcpy(res, d_res, size, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_res); }

Сложение матриц void sum(int* a, int* b, int* res, int n) { int* d_a; int* d_b; int* d_res; int size = n * n * sizeof(int); cudaMalloc((void**)&d_a, size) cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_res, size); cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y); sumKernel<<>>(d_a, d_b, d_res, n); cudaMemcpy(res, d_res, size, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_res); }

Сложение матриц void sum(int* a, int* b, int* res, int n) { int* d_a; int* d_b; int* d_res; int size = n * n * sizeof(int); cudaMalloc((void**)&d_a, size) cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_res, size); cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); dim3 block(16, 16); dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y); sumKernel<<>>(d_a, d_b, d_res, n); cudaMemcpy(res, d_res, size, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_res); }

Сложение матриц (ядро) __global__ void sumKernel(int* a, int* b, int* res, int n) { int i = threadIdx.x + blockIdx.x * blockDim.x; int j = threadIdx.y + blockIdx.y * blockDim.y; if ((i < n) && (j < n)) { int index = i + j * n; res[index] = a[index] + b[index]; } }

Спецификаторы функций

Память

Взаимодействие потоков __global__ void kernel(float* a, float* b) { __shared__ float s[BLOCK_SIZE]; s[threadIdx.x] = a[threadIdx.x + blockIdx.x * blockDim.x]; __syncthreads(); … }

Схема выполнения ядра Блоки распределяются по мультипроцессорам Кол-во одновременно выполняемых блоков определяется ресурсами Регистры выделяются потокам Разделяемая память выделяется блокам Блок потоков разбивается на варпы (warp) Все потоки варпа выполняют одну и ту же операцию ядра

Редукция A = a0 + a1 + … + an-1 int reduction(int* a, int n) { int sum = 0; for (int i = 0; i < n; i++) sum +=a[i]; return sum; }

Параллельная редукция

blockSize = 8 а data data итерация 1 data итерация 2 data итерация 3 shift = 4 shift = 2 shift = 1 b

__global__ void reduce (int* a, int* b) { __shared__ int data[blockSize]; //blockSize = blockDim.x int threadIndex = threadIdx.x; int globalIndex = threadIndex + blockIdx.x * 2 * blockSize; data[threadIndex] = a[globalIndex] + a[globalIndex + blockSize]; __syncthreads(); for (int shift = blockSize / 2; shift > 0; shift /= 2) { if (threadIndex < shift) data[threadIndex] += data[threadIndex + shift]; __syncthreads(); } if (threadIndex == 0) b[blockIdx.x] = data[0]; }

Атомарные операции Арифметические: atomicAdd atomicSub atomicExch atomicMin atomicMax и др. Побитовые: atomicAnd atomicOr atomicXor

Compute Capability

Compute Capability

Примеры видеокарт

Оптимизация Максимизация независимого параллелизма Максимизация интенсивности вычислений Минимизация передачи данных между CPU и GPU Минимизация ветвлений внутри варпа Оптимизация работы с памятью

Оптимизация работы с памятью Использование разделяемой памяти Использование coalescing Избежание конфликтов банков разделяемой памяти

Оптимизация работы с памятью Использование разделяемой памяти Использование coalescing Избежание конфликтов банков разделяемой памяти

Источники NVidia CUDA Zone: http://www.nvidia.com/object/cuda_home_new.html www.gpgpu.ru Боресков А.В., Харламов А.А., 2010: Основы работы с технологией CUDA. Джейсон Сандерс, Эдвард Кэндрот, 2011: Технология CUDA в примерах. Введение в программирование графических процессоров.