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

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

39624-chasty1-cuda.ppt

  • Количество слайдов: 35

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

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

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

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

>GPU GPU

>GPU GPU

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

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

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

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

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

>Потоки gridDim.x = 3 gridDim.y = 2 gridDim.z = 1  blockIdx.x = 1 Потоки 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) {  Сложение матриц 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) {  Сложение матриц 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) {  Сложение матриц 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) {  Сложение матриц 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) {  Сложение матриц 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) {  Сложение матриц 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) { Сложение матриц (ядро) __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__ Взаимодействие потоков __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, Редукция 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 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 __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 Атомарные операции Арифметические: atomicAdd atomicSub atomicExch atomicMin atomicMax и др. Побитовые: atomicAnd atomicOr atomicXor

>Compute Capability Compute Capability

>Compute Capability Compute Capability

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

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

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

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

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