�ݺ�ߣ

Параллельное программированиеНа современных видеокартах(GPGPU)Алексей Тутубалинwww.gpgpu.rulexa@lexa.ru

Параллельные вычисления: зачем?Производительности всегда не хватает

Рост производительности CPU сильно замедлилсяThe Landscape of Parallel Computer Research: A View from Berkeley (2006)http://www.eecs.berkeley.edu/Pubs/TechRpts/2006/EECS-2006-183.pdf - старые парадигмы во многом устарели, параллельность – единственный (на сегодня) путьОтвет индустрии оборудования: многоядерностьCPUа до того – многопроцессорные машины с общей памятью (SMP) и векторные инструкцииПрограммистам от параллельности деваться некуда

Параллельные вычисленияКак?Одновременное исполнение независимых вычислений на разных процессорах с последующей синхронизациейПроблемы:Правило Амдала: ускорение=P – доля программы, которую можно распараллелить, N – число потоков исполненияСлишком маленькие независимые блоки много затрат на синхронизациюСложно программировать (много объектов, синхронизация, критические секции)

Параллельные системы

Высокая производительность: объективные проблемыЗадержки (Latency)Часто определяют производительностьПример: Умножение матриц наивным способом - медленный доступ «по столбцам»: особенности памяти (и кэшей, если не повезло)Синхронизация: всегда непроизводительное ожиданиеСкорость чтения/записи в память (Bandwidth)Часто определяет производительностьSSE на Core i7: две векторные (16 байт) операции на такт(96 байт/такт), скорость доступа к памяти порядка 6 байт/такт (или 1-1.5 байта на ядро)Нужна быстрая память, кэши частично спасаютНужно повышение сложности: больше операций над данными в регистрах.Распределенные системы: линейный рост bandwidthАрифметическая интенсивность (операций/единица данных)a = b+c: 1 операция, 2 чтения, 1 записьРаспараллеливание может стать бессмысленным, весь выигрыш будет съеден пересылкамиЧтобы был выигрыш от параллельности, алгоритмы должны быть либо сложнее O(n),либо для O(n) – константапри O – большая.

Вычисления на GPU: предпосылки

«Старые видеокарты» (2003-2006)Специализация под 3D-графику (отрисовка2D-проекции сцены)Фиксированные стадии обработки (вершины-геометрия-пиксели)С 2001 – программируемые шейдеры (GeForce3, OpenGL 1.5)Высокая производительность за счет:Архитектуры, сделанной ровно под одну задачуВысокой скорости линейного чтения памятиОтсутствия синхронизации внутри стадии обработки

2001-2006: «Многообещающие результаты»~3000 научных публикацийС общим смыслом «вот на следующих поколениях оборудования все будет хорошо...»Первые средства разработки для не-графики: Brook, ShРезультаты:Ускорение линейной алгебры в разы (относительно CPU), но single precisionСортировка: ускорение в разы, GPUTeraSortвыиграла Sort Benchmark 2006гПромышленное использование:Обработка сигналовОбработка цифрового кино (поток 200-500Mb/sec)

Проблемы начального этапаНеудобная парадигма «Потоковых вычислений» (Stream Computing):«микропрограммы» (шейдеры) – обработка одного элемента входного потокаНикакого взаимодействия между потокамиМногие задачи очень неудобно программироватьТолько одинарная точность вычисленийНеудобство средств разработки

2006: NVidia CUDA и G80 Compute Unified Device ArchitectureПринципиально новая (для GPU) архитектураВзаимодействие между потоками вычисленийПроизвольный код вместо шейдеровИерархическая структура памяти345 GFLOP/s (single), 87Gb/sec memory bandwidthСредства программирования общего назначения (CUDA):Вычислительный код на языке «С»Компактная архитектура запуска вычислителейМного готовых вычислительных примеровГотовые библиотеки BLAS и FFT в поставкеПрекрасная поддержка разработчиков и университетов2007: NVidia Tesla – «Supercomputer on Desktop» - «видеокарты без видеовыхода», только вычисления

2006: интерес других производителейATI (AMD):Инициатива Close-To-Metal: публикация спецификаций кода, надежда на отказ от OpenGL и появление нормальных компиляторовFireStream – первый ускоритель для вычислений а не видеокартаОстаются в парадигме потоковых вычисленийCell (IBM-Sony-Toshiba)8 ядер с локальной памятью (256кб на ядро)Быстрая внешняя шина (25Gb/sec)230 GFLOP/s single, 15-20GFLOP/s doublePlayStation 3 и плата в серверы

2007-2008: быстрое взросление>2000 публикаций про использование CUDAТоповые конференции (Supercomputing)Хорошие готовые библиотекиНовое оборудование:Поддержка двойной точности у всехNVidia: G200 – 800 GFLOP/s single, 100 – double, ATI HD4870: 1000/240 GFLOP/s single/double, но все еще потоковое программированиеIBM PowerXCell: ~200/100 GFLOP/s

Ускорение приложений на CUDA2008 г.: NVidia 8800GTX в сравнении с AMD Opteron 248 (2.2GHz, 2 ядра)

Современный этап: оборудованиеВидеокарты:ATI R8002.7TFLOP/s single, 500 GFLOP/s doubleПоддержка OpenCL (не только потоковые вычисления)NVidia GF100 (Fermi):Tesla 1.2 TFLOP/s single, 600 GFLOP/s doubleGeforce GTX 480: 1.3 TFLOP/s single, ~180 GFLOP/s doubleДля сравнения, топовые x86 CPU: 80-110GFLOP/s (double)GPU-кластерыПромышленные решения (Cray, IBM, Т-платформы)Суперкомпьютеры: Nebulae - №2 в Supercomputer Top500 по реальной производительности, №1 по пиковой.

Современный этап: программыСтандарт OpenCL (лето-осень 2009): переносимые GPU-программыДля пользователяускорение multimedia, обработка изображений и видео, 3D-моделирование, инженерные программы...Для программистамного готовых библиотек: линейная алгебра, обработка сигналов, computer visionДля ученого: поддержка в расчетных программахОбщего назначения: Matlab, MathematicaСпециализированные пакеты: проектирование, компьютерная химия, компьютерная биология, молекулярная динамика

Современный этап: производительность

Наглядная агитация4xTesla == Cray XT4 на задачах молекулярной биологии (AMBER11)(слайд с NVidia GPU Technology conference, 22 сентября 2010)

NVidia CUDACompute Unified Device ArchitectureОбщий обзор «сверху вниз»Как устроено оборудованиеИ как это отражается в логической структуре программМинимальные примеры программНекоторые тонкие моментыОбзор средств программиста и дополнительных библиотекНет задачи и возможности заменить чтение документации и самостоятельные упражнения

Общая схема исполнения программВсе стадии независимы, одновременно можно:Передавать в видеопамятьСчитать (другие данные)Передавать из видеопамяти

Архитектура G80До 16 одинаковых мультипроцессоровна чипеОбщая глобальная память (R/W): медленная и много (до 1.5Gb)Общая константная память (RO): быстрая и малоОбщий планировщик потоков (блоками)Общий исполняемый код на вечь чипG80Streaming Multiprocessor (SM) 8 простых процессоров (SP): исполняют 32 потока за 4 такта в SIMD-режиме

2 специальных процессора (SFU) для сложных функций: 32 потока за 16 тактов, SIMD

Один диспетчер команд

до 768 потоков на одном SM в «одновременном» режиме

потоки объединяются в пачки по 32 потока (warp), которые исполняются SIMD

16 kb разделяемой памяти (shared memory)

8192 32-битных регистраНа 2 SM – один блок обработки текстур с текстурным кэшем и блоком аппаратной интерполяции…TPCTPCTPCTPCTPCTPCStreaming MultiprocessorTexture Processor ClusterInstruction L1Data L1Instruction Fetch/DispatchSMShared MemoryTEXSPSPSPSPSMSFUSFUSPSPSPSP

TextureTextureTextureTextureTextureTextureTextureTextureTextureHostInput AssemblerThread Execution ManagerParallel DataCacheParallel DataCacheParallel DataCacheParallel DataCacheParallel DataCacheParallel DataCacheParallel DataCacheParallel DataCacheLoad/storeLoad/storeLoad/storeLoad/storeLoad/storeLoad/storeGlobal MemoryG80: общая схема

Логическая организация программДесятки тысяч – миллионы независимых потоков в программеПоток – обрабатывает часть входных данныхПотоки объединены в блокиОбщая shared memoryСинхронизацияБлоки объединены в gridСинхронизации нетGrid of Thread Blocks

Логическая организация (2)Блок исполняется на одном мультипроцессореРазные блоки – на разныхПорядок не определенМожет быть несколько блоков на SM одновременноНумерация потоков в блоке: 1D, 2D или 3DНумерация блоков: 1D,2DКаждый поток знает свое положение в сетке и блоке

Параметры SM – не догмаG92:16384регистров, до 1024 потоков на SMGF100 (Fermi)32768 регистров и 32/48 процессоров на SM 1536 потоков (48 warp) на SM64kb shared memory+L1-cache (16/48 или 48/16 shared/cache)Другие параметры другой оптимальный кодДругие улучшенияАтомарные операции (G86, G92)Поддержка двойной точности (G200, GF100)

Пример: cложение двух векторовКод GPU: складываем два элемента:__global__ void VecAdd(float* A,float* B, float* C, int N){ int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < N) C[i] = A[i]+B[i];}threadIdx – предопределенная переменная (3-компонентный вектор) – номер потока в блокеblockDim – размерность одного блока (3D)blockIdx – номер блока в сетке (2D)

код на CPU// аллоцируем память на GPUsize_t size = N * sizeof(float);cudaMalloc((void**)&d_A, size);cudaMalloc((void**)&d_B, size);cudaMalloc((void**)&d_C, size);// копируем данные на GPUcudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);// запускаем kernel на GPU (1D сетка и блок):int Nthreads = 256;int Nblocks = (N+Nthreads-1)/Nthreads;VecAdd<<<Nblocks,Nthreads>>>(d_A, d_B, d_C, N);// получаем результаты с GPUcudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

CUDA:С/C++ с небольшими изменениямиСпецификаторы типа функции:__global__- GPU, может быть вызвана с CPU (с параметрами)__device__ - GPU, вызывается из __global____host__ - CPUРекурсия и указатели на функции: только Fermi (GF100) Спецификаторы местоположения переменной__device__ - глобальная память, R/W__constant__ - константная память, R/O__shared__ - shared memory, R/W, локальна в блокеАвтоматические переменные – могут быть в регистрах, если известен размер и регистров достаточно. Иначе – local memory, т.е. живут в медленной глобальной памяти

Типы данных и переменныеСкалярные:char, short, int, long,longlong, float, double (и беззнаковые для целых)Векторные: <=32 бит/компонент – до 4 элементов (uchar2,float4)64 бита – 2 элемента (double2, longlong2)Векторных операций НЕТ: c.x=a.x+b.x работает, c=a+b - нетdim3 – тип для задания размера сетки блоковВстроенные переменные:gridDim, blockDim– размер сетки и блока (dim3)blockIdx, threadIdx– номер блока и потока (dim3)warpSize – ширина SIMD (пока всегда 32)

Двумерные блоки например, для двумерных моделейНа CPU:dim3 blockSize(16, 16);// некратным размером сетки пренебрегаем... dim3 gridSize(width/blockSize.x, height/blockSize.y);myKernel<<<gridSize,blockSize>>>(…..)На GPUuint x = blockIdx.x*blockDim.x+threadIdx.x;uint y = blockIdx.y*blockDim.y +threadIdx.y;

Синхронизация, атомарные операцииГлобальная синхронизация всех блоков – завершение kernel__syncthreads() – синхронизация потоков блокаАтомарные сложения, вычитания и т.п. (atomicAdd() и т.д.) - G86+Shared или Global memory (G86 – только global)Атомарное исполнениеМедленные

Скалярное произведениеНакапливаем частные произведенияпо потокам:__global__ result;shared__ float psum[NTHREADS];psum[threadIdx.x] = 0;for(j=0;j<K*NTHREADS;j+=NTHREADS)psum[threadIdx.x] += A[base+j]*B[base+j];__syncthreads();// Складываем для блокаif(threadIdx.x == 0){ for(j=1;j<blockDim.x;j++)psum[0]+=psum[j]; // синхронизации не надо!! 0-йthreadatomicAdd(&result,psum[0]);}

ТонкостиОптимальный доступ к глобальной памятиОсобенности доступа к shared memoryУсловные операцииРаспределение блоков по процессорам

Доступ к глобальной памятиСложение векторов, 16 сложений на thread:___________________________________________________int j,base = (blockDim.x*blockIdx.x+threadIdx.x)*16;for(j=0;j<16;j++)C[base+j] = A[base+j] + B[base+j]Thread0: A[0], A[1]Thread1: A[16], A[17]Весь Warp обращается c «размахом» в 512 слов______________________________________________________int j,base = (blockDim.x*blockIdx.x)*16+threadIdx.x;for(j=0;j<blockDim.x*16;j+=blockDim.x)C[base+j] = A[base+j] + B[base+j]Thread0: A[0], A[blockDim.x]...Thread1: A[1], A[blockDim.x+1]....Весь Warp обращается к соседним словамВторой вариант – coalesced access: доступ к соседним адресам

Доступ к глобальной памяти - 2Прямые рекомендации из Programming Guide:Оптимально, когда все потоки одного (полу)-warp*обращаются к одному 128-байтному сегментуДля 16-байтных элементов (вектор) – два подряд 128-байтных сегмента.В правильном порядке – k-йthread к k-му элементу в сегменте.Идеальная ситуация: Array[base + treadIdx.x]Array[base] – выровнен на 128 байт

Разделяемая памятьРазделена на банки 16 банков на G80-G200, 32 на GF100Interleave 4 байтаДва потока в 1 банк =>конфликт, каждый конфликт – 4 такта ожиданияОптимальный доступ array[threadIdx]

Нет конфликта банковBank 0Thread 0Bank 0Thread 0Bank 1Thread 1Bank 1Thread 1Bank 2Thread 2Bank 2Thread 2Bank 3Thread 3Bank 3Thread 3Bank 4Thread 4Bank 4Thread 4Bank 5Thread 5Bank 5Thread 5Bank 6Thread 6Bank 6Thread 6Bank 7Thread 7Bank 7Thread 7Bank 15Thread 15Bank 15Thread 15

Конфликт банковBank 0Bank 1Thread 0Bank 2Thread 1Bank 3Thread 2Bank 4Thread 3Bank 5Thread 4Bank 6Bank 7Thread 8Thread 9Bank 15Thread 10Thread 11GF100: несколько чтений разных потоков по одному адресу (не банку) – не приводят к конфликту (broadcast данных)

Условные операцииSIMD поэтому если в одном warp разные ветки if, их исполнение будет последовательным:if (threadIdx.x %2) { четные ветки стоят, нечетные работают}else { четные работают, нечетные стоят }

Распределение блоков по SMНе меньше одного блока т.е. блок должен пролезать по ресурсам:Общее число потоковв блоке < лимита оборудованияТребуемое количество shared memory< лимитаОбщее количество регистров (регистров на поток * число потоков) < лимита Несколько блоков на SM – если достаточно ресурсов (регистров, shared memory, потоков)G80-G200 – передача параметров в shared mem, 2 X 8192 – не влезет в 16 килобайт

Загрузка SM (occupancy)Отношение числа работающих warps к из возможному количествуВ общем случае, к 100% надо стремиться:Прячется латентность доступа к памятиНо мало регистров (G92 – 16 регистров/thread)Альтернативный подход:Мало потоков (128-192-256 на SM)Зато много регистров на потокПолезно для блочных алгоритмовПодробнее: Volkov, V. 2010. Use registers and multiple outputs per thread on GPUhttp://www.cs.berkeley.edu/~volkov/volkov10-PMAA.pdf

РазноеЦелая арифметика:G80-G200: 24-бита – быстрая, 32 бита – медленнаяGF100 – 32 бита – быстрая, 24 бита - медленнаяТекстуры: 1-2-3D-массивыКэшируются, кэш локален по координатам (эффективное кэширование по столбцам, при последовательном доступе – медленнее global memory)Аппаратная интерполяция

Продолжение.....Документация NvidiaОчень хорошая!http://developer.nvidia.com/object/cuda_download.htmlCUDA Programming GuideCUDA Reference GuideCUDA Best Practices GuideПримеры SDK и документация к нимhttp://developer.download.nvidia.com/compute/cuda/sdk/website/samples.htmlФорумы Nvidiahttp://forums.nvidia.com/index.php?showforum=62

Программные средстваCUDA Toolkit – компилятор + библиотекиИнтегрируется с Visual C++ (добавление правил компиляции для .cu)Отладчик – gdb (командная строка)Visual Profiler – не на уровне отдельных строк кодаParallel NsightПошаговая отладкаГлобальный профайлингStandard (бесплатная) и Professional (бесплатная для ВУЗов) версииCUDA-x86Был эмулятор, но в Cuda 3.x удаленPGI обещает компилятор

Подробности компиляцииfloat4 me = gx[gtid];me.x += me.y * me.z;C/C++ CUDAApplicationCPU-code компилируется родным компиляторомPTX – неоптимизированный ассемблер, дальнейшие варианты: Компиляция в один executable

Компиляция в внешний .cubin-файл и загрузка на runtimeФинальную компиляцию реальные коды делает драйвер видеокартыNVCCCPU CodePTX CodeVirtualPhysicalPTX to TargetCompiler G80 … GPU Target code

Два APICUDA C API: реализуется внешней библиотекой cudartВысокоуровневое, более удобноеудобная поддержка нескольких контекстов (например для использования в многопоточных программах)CUDA Driver API: Аналогично по возможностямНе требует дополнительных библиотекБольше работы программисту

Альтернативы CUDAOpenCLCAL/IL (Compute Abstraction Layer/Intermediate Language)Только видеокарты ATIПотоковое программированиеПлохо поддерживаетсяDirectX 11 (DirectCompute)Только Windows

OpenCLМногоплатформенный стандартНеобычайно похож на CUDA Driver APIДетали:Язык C для kernels, другие ключевые словаКомпиляция из исходных текстов на ходу («в драйвере»)Уже две версии (1.0 и 1.1)Нестандартная конфигурация поддерживается через «расширения»

Сложение векторовКод GPU:__kernel void VectorAdd(__global const float* a,__global const float* b,__global float* c, int N){ int iGID = get_global_id(0); if (iGID < N) c[iGID] = a[iGID] + b[iGID];}

Код CPU// Создаем контекст GPU (до того был поиск устройств...)Context = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr1);// и очередь командCommandQueue = clCreateCommandQueue(Context, cdDevice, 0, &ciErr1);// Аллоцируем массив(ы) на GPUcmDevSrcA = clCreateBuffer(Context, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr1);…. И еще два аллоцирования ....// Создаем (компилируем) программу из исходных текстовcpProgram = clCreateProgramWithSource(Context, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);// создаем из нее kernelckKernel = clCreateKernel(cpProgram, "VectorAdd", &ciErr1);// задаем ему аргумент(ы)clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrcA);.... И еще три аргумента ......// В очередь команд помещаем копирование данных и запуск kernel, все асинхронноclEnqueueWriteBuffer(CommandQueue, cmDevSrcA, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcA, 0, NULL, NULL);.... И еще одно копирование....clEnqueueNDRangeKernel(CommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);// В очередь команд помещаем синхронное чтение результатаclEnqueueReadBuffer(CommandQueue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);

CUDA  OpenCLAMD: Porting CUDA to OpenCLhttp://developer.amd.com/zones/OpenCLZone/Pages/portingcudatoopencl.aspx

Переносимость и производительностьПереносимыйкод (NVidia, ATI,x86, Cell)Производительность не переносится:Либо используем векторные типы на векторных архитектурах (ATI, Cell) и получаем непереносимый кодЛибо код получается переносимый и в разы медленнееПеренос CUDA  OpenCL – потеря 20-30%. Компилятор??CUDA всегда будет впереди т.к. нет этапа согласований с рабочей группой

OpenCL и видеокарты ATIЧеловеческий интерфейс (в отличие от CAL/IL)HD5870 – производительность примеров SDK - на уровне GTX280 (в разы ниже ожидаемого)HD4xxx – производительность плохая т.к. нет shared memory (эмуляция через глобальную)Плохая документация (в сравнении с...)Но что делать, если попалось такое оборудование и не хочется ассемблера.....

Почему GPU НАСТОЛЬКО быстрее?Линейная алгебра: 5-8раз

Обработка сигналов, обработка изображений: десятки раз

Некоторые задачи (молекулярная динамика, N-body): бывают сотни разВот почему:Другой аппаратный дизайн:CPU – оптимизация под любые задачи (включая очень плохой код) и низкую латентность. GPU – оптимизация под общую производительность и только высокопараллельные задачи.Быстрая память:В разы выше bandwidth видеопамятиБольше очень быстрой памяти(2-3Mbпротив 32-256Kb L1 cache CPU)Аппаратная поддержка ряда функций (exp, sin, cos, pow, log….): разница с CPU в десятки и даже сотни (!) разЧасто сравнивают с неоптимизированным кодом (а другого и нет, писать руками SSE-инструкции скучно)

�ݺ�ߣ

Параллельное программирование на современных видеокартах

Recommended

More Related Content

What's hot (20)

Similar to Параллельное программирование на современных видеокартах (20)

Параллельное программирование на современных видеокартах