Массивно-параллельные вычислительные системы на основе архитектуры CUDA. - презентация

1 Массивно-параллельные вычислительные системы на основе архитектуры CUDA.

2 Особенности многоядерных систем 1)Как правило вычислительный узел – достаточно маломощный процессор 2)Вычислительные узлы имеют свою оперативную память и свой кэш 3)Вычислительные узлы объединяются в более крупные блоки 4)Крупные блоки могут объединяться с целью наращивания вычислительной мощи

3 План zАрхитектура Tesla zПрограммная модель CUDA zСинтаксические особенности CUDA

4 Архитектура Tesla: Мультипроцессор Tesla 8 TEX SM Texture Processing Cluster Streaming Multiprocessor Instruction $Constant $ Instruction Fetch Shared Memory SFU SP SFU SP Register File

5 TEX SM Texture Processing Cluster SM Архитектура Tesla Мультипроцессор Tesla 10 Streaming Multiprocessor Instruction $Constant $ Instruction Fetch Shared Memory SFU SP SFU SP Double Precision Register File

6 Архитектура Tesla 10 TPC Interconnection Network ROP L2 ROP L2 ROP L2 ROP L2 ROP L2 ROP L2 ROP L2 ROP L2 DRAM Work Distribution CPUBridgeHost Memory

7 Архитектура zМаштабируемость: [+][-] SM внутри TPC [+][-] TPC [+][-] DRAM партиции zСхожие архитектуры: Tesla 8: 8800 GTX Tesla 10: GTX 280

8 План zАрхитектура Tesla zСинтаксические особенности CUDA zПрограммная модель CUDA

9 Программная модель CUDA z GPU (device) это вычислительное устройство, которое: yЯвляется сопроцессором к CPU (host) yИмеет собственную память (DRAM) yВыполняет одновременно очень много нитей

10 Программная модель CUDA zПоследовательные части кода выполняются на CPU zМассивно-параллельные части кода выполняются на GPU как ядра zОтличия нитей между CPU и GPU yНити на GPU очень «легкие» yHW планировщик задач yДля полноценной загрузки GPU нужны тысячи нитей xДля покрытия латентностей операций чтения / записи xДля покрытия латентностей sfu инструкций

11 Программная модель CUDA zПараллельная часть кода выполняется как большое количество нитей zНити группируются в блоки фиксированного размера zБлоки объединяются в сеть блоков zЯдро выполняется на сетке из блоков zКаждая нить и блок имеют свой идентификатор

12 Программная модель CUDA zДесятки тысяч потоков for (int ix = 0; ix < nx; ix++) { pData[ix] = f(ix); } for (int ix = 0; ix < nx; ix++) for (int iy = 0; iy < ny; iy++) { pData[ix + iy * nx] = f(ix) * g(iy); } for (int ix = 0; ix < nx; ix++) for (int iy = 0; iy < ny; iy++) for (int iz = 0; iz < nz; iz++) { pData[ix + (iy + iz * ny) * nx] = f(ix) * g(iy) * h(iz); }

13 Программная модель CUDA zПотоки в CUDA объединяются в блоки: yВозможна 1D, 2D, 3D топология блока xОбщее кол-во потоков в блоке ограничено xВ текущем HW это 512 потоков

14 Программная модель CUDA zПотоки в блоке могут разделять ресурсы со своими соседями float g_Data[gN]; for (int ix = 0; ix < nx; ix++) { pData[ix] = f(ix, g_Data[ix / n]); }

15 Программная модель CUDA zБлоки могут использовать shared память y Т.к. блок целиком выполняется на одном SM y Объем shared памяти ограничен и зависит от HW zВнутри Блока потоки могут синхронизоваться yТ.к. блок целиком выполняется на одном SM

16 Программная модель CUDA zБлоки потоков объединяются в сетку (grid) потоков yВозможна 1D, 2D топология сетки блоков потоков

17 План zАрхитектура Tesla zСинтаксические особенности CUDA zПрограммная модель CUDA

18 Синтаксис CUDA z CUDA – это расширение языка C y [+] спецификаторы для функций и переменных y [+] новые встроенные типы y [+] встроенные переменные (внутри ядра) y [+] директива для запуска ядра из C кода z Как скомпилировать CUDA код y [+] nvcc компилятор y [+].cu расширение файла

19 Синтаксис CUDA Спецификаторы СпецификаторВыполняется наМожет вызываться из __device__device __global__devicehost __host__host zСпецификатор функций zСпецификатор переменных СпецификаторНаходитсяДоступнаВид доступа __device__device R __constant__devicedevice / hostR / W __shared__deviceblock RW / __syncthreads()

20 Синтаксис CUDA Встроенные переменные z Сравним CPU код vs CUDA kernel: __global__ void incKernel ( float * pData ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; pData [idx] = pData [idx] + 1.0f; } float * pData; for (int ix = 0; ix < nx; ix++) { pData[ix] = pData[ix] + 1.0f; } Пусть nx = 2048 Пусть в блоке 256 потоков кол-во блоков = 2048 / 256 = 8 [ ][ == 256][ ]

21 Синтаксис CUDA Встроенные переменные zВ любом CUDA kernele доступны: ydim3 gridDim; yuint3 blockIdx; ydim3 blockDim; yuint3 threadIdx; yint warpSize; dim3 – встроенный тип, который используется для задания размеров kernelа По сути – это uint3.

22 Синтаксис CUDA Директивы запуска ядра zКак запустить ядро с общим кол-во тредов равным nx? incKernel ( pData ); dim3 threads(256, 1, 1); dim3 blocks(nx / 256, 1); float * pData; >> угловые скобки, внутри которых задаются параметры запуска ядра: Кол-во блоке в сетке Кол-во потоков в блоке … Неявно предпологаем, что nx кратно 256

23 Как скомпилировать CUDA код zNVCC – компилятор для CUDA yОсновными опциями команды nvcc являются: y-deviceemu - компиляция в режиме эмуляции, весь код будет выполняться в многонитевом режиме на CPU и можно использовать обычный отладчик (хотя не все ошибки могут проявится в таком режиме) y--use_fast_math - заменить все вызовы стандартных математических функций на их быстрые (но менее точные) аналоги y-o - задать имя выходного файла zCUDA файлы обычно носят расширение.cu

24 Google-группа zCUDA.CS.MSU.SUCUDA.CS.MSU.SU