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

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

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

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

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

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

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

7 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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

28