Архитектура Tesla. Программно-аппаратный стек CUDA. zЛекторы: yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) yХарламов А.А. (NVidia)Харламов А.А. (NVidia)

Примеры многоядерных систем zНа первой лекции мы рассмотрели yIntel Core 2 Duo ySMP yCell yBlueGene/L yG80 / Tesla

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

Tesla vs GeForce zУ кого есть вопросы в чем разница?

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

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

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

Архитектура 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

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

Технические детали zRTM CUDA Programming Guide zRun CUDAHelloWorld yПечатает аппаратно зависимые параметры xРазмер shared памяти xКол-во SM xРазмер warpа xКол-во регистров на SM x т.д.

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

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

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

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

Программная модель 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); }

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

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

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

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

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

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

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

Синтаксис 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][ ]

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

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

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

Ресуры нашего курса zCUDA.CS.MSU.SUCUDA.CS.MSU.SU yМесто для вопросов и дискуссий yМесто для материалов нашего курса yМесто для ваших статей! xЕсли вы нашли какой-то интересный подход! xИли исследовали производительность разных подходов и знаете, какой из них самый быстрый! xИли знаете способы сделать работу с CUDA проще! z Steps3d Steps3d z