МОСКОВСКИЙ АВИАЦИОННЫЙ ИНСТИТУТ (национальный исследовательский университет) (национальный исследовательский университет) Факультет прикладной математики и физики Кафедра вычислительной математики и программирования Выполнил: Семенов С.А. Руководитель: Ревизников Д.Л. Лекция 9 « Буфер инструкций SM. Регистровый файл SM »

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 2 Введение Буфер инструкций SM Регистровый файл SM Конвейеры исполнения команд

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 3 CUDA поток вычислений

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 4 Почему переключение контекста такое эффективное? Большой размер регистрового файла (16K registers/block) –Каждому треду назначается «окно» физических регистров –Работает в том случае, если размер регистров блока на превысил максимальный размер (в противном случае ошибка компилятора) –Можно составлять расписание из множества блоков одновременно Подобно, затраты разделяемой памяти на должны превышать максимальный размер для всех блоков одновременно исполняющихся Register File Block 0 Thread 0 Block 0 Thread 1 Block 0 Thread 256 Block 8 Thread 0 Block 8 Thread 1 Block 8 Thread 256

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 5 Пример: Расписание тредов G80 Каждый блок выполняется как вара, состоящий из 32 тредов –Эта реализация не часть программной модели CUDA –Варп – единица запуска на SM В 3 блоках SM по 256 тредов, как много карпов на SM? –Каждый блок состоит из 256/32 = 8 Warps –Тогда 8 * 3 = 24 Warps … t0 t1 t2 … t31 … … … Block 1 WarpsBlock 2 Warps SP SFU SP SFU Instruction Fetch/Dispatch Instruction L1 Streaming Multiprocessor Shared Memory … t0 t1 t2 … t31 … Block 1 Warps

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 6 SM Warp Scheduling SM аппаратная реализация предполагает Варп расписание без накладных расходов –Варп, у которого будет выполняться следующая инструкция, готов к выполнению и может запускаться –Очерендые вараы выбираются для выполнения приоритизованной политикой расписания –Все треда одного вараы выполняют одну инструкцию, когда он запущен 4 такта необходимо для выборки одной инструкции для всех карпов на G80 –Если доступ к глобальной памяти требует выполнения 4 инструкций –То всего 13 карпов приведут к задержке 200 таков в доступе к памяти warp 8 instruction 11 SM multithreaded Warp scheduler warp 1 instruction 42 warp 3 instruction 95 warp 8 instruction time warp 3 instruction 96

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 7 SM буфер инструкций – Warp Scheduling Извлекает одну вара инструкцию/цикл –Из кэша инструкций –Из любого буферного слота инструкций Отправляет одну готов к запуску вара инструкцию/цикл –Из любого вараа – буферного слота инструкций –модуль scoreboarding используется для согласования Отправляет выбранный вара на основе цикличного/временного показателя вараа SM бродкаст определяет одну инструкцию для 32 тредов вараа I$ Multithreaded Instruction Buffer R F C$ L1 Shared Mem Operand Select MADSFU

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 8 Scoreboarding Как определить что тред уже выполняется? scoreboard - аппаратная таблица, которая следит: –За выборкой инструкций, отправленных, исполненных –За ресурсами (функциональные блоки и операнды) –Какая инструкция какой регистр изменила Концепция пришла из CDC 6600 (1960), включала в себя разделение памяти и вычислений

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 9 Scoreboarding Все регистровые операнды всех инструкций в буфере инструкций - scoreboarded –Назначается статус готовности как только необходимое значение загружено –Предотвращается несогласованность –Очищается инструкции готовые к отправке Разбивает Память/Процессор пайплайны –Каждый тред продолжает отправлять инструкции пока scoreboarding не остановит его –Позволяет память/процессорные операции в теневом режиме

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 10 Scoreboarding на примере Рассмотрим 3 отдельных потока: warp1, warp3, warp8 warp 8 instruction 11 warp 1 instruction 42 warp 3 instruction 95 warp 8 instruction warp 3 instruction 96 t=k t=k+1 t=k+2 t=l>k t=l+1 WarpCurrent Instruction State Warp 142Computing Warp 395Computing Warp 811Operands ready to go … Schedule at time k

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 11 Scoreboarding на примере Рассмотрим 3 отдельных потока: warp1, warp3, warp8 warp 8 instruction 11 warp 1 instruction 42 warp 3 instruction 95 warp 8 instruction warp 3 instruction 96 t=k t=k+1 t=k+2 t=l>k t=l+1 WarpCurrent Instruction State Warp 142Ready to write result Warp 395Computing Warp 811Computing … Schedule at time k+1

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 12 Буфер инструкций SM Буфер инструкций выбирает вора и инструкцию, которые будут исполнены в следующий момент времени Критерии выбора: Готовность данных Длительность исполнения вораа (старые имеют) Активный вора будет исполнять свои инструкции последовательно, пока не возникнет ситуация, мешающая Нет промежуточных результатов для новой Не прочитаны операнды из памяти Загрузка конвейера Программно 32/16 тредов на вара Аппаратно 8 SP в каждом SM

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 13 Ветвления внутри вораа Ветвления разрушают SIMD структуру исполнения Инструкции ложной ветви не исполняются для текущего треда Время теряется как если бы все треды вораа прошли всеми возможными путями (последовательно) Компилятор может восстанавливать точки ре- синхронизации, асинхронизировавшихся тредов Программист может помогать используя __syncthreads();

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 14 Регистровый файл SM В каждом SM регистровый файл Размер = 32K Распределён неравномерно между SP За один clock можно прочитать 4 операнда для каждого SP TEX и Load/Store могут читать и писать RF

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 15 Конвейеры исполнения команд Скалярный конвейер MAD (SPs) FMUL,FADD,FMAD Целочисленные операции, приведение типов 1 инструкция за 1 clock Скалярный конвейер SFU RCP,RSQ,LG2,EX2,SIN,COS 1 инструкция за 4 clocks можно также реализовывать FMUL, MOV TEX конвейер (RO lдоступ к табулированным константам) LoaD / STore конвейер Перенос содержимого регистров в/из локальную память Доступ к локальной и глобальной памяти

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 16 Память констант память констант находится в DRAM на плате Каждый из SM имеет L1 кэш в операциях с памятью констант Константы могут адресоваться без индекса Линейный индекс на основе threadID

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 17 Разделяемая память В каждом SM – 16 k RW разделяемой памяти 16 банков 32-битных слов Последовательная адреса ячеек принадлежат последовательным банкам Обращения к разным банкам возможны одновременно Обращения с конфликтами реализуются как несколько последовательных обращений Каждое чтение исполняется от тредов полувораа

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 18 Parallel Thread eXecution Virtual Machine (PTX VM) Parallel Thread eXecution (PTX) Virtual Machine, а также ISA Программная модель ISA – Instruction Set Architecture Variable declarations Instructions and operands Транслятор – оптимизирующий Компилятор Трансляция PTX исполняемый код Драйвер видеокарты реализует VM runtime

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 19 PTX код (пример 1)

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 20 PTX Код (пример 2) CUDA PTX

Московский авиационный институт (национальный исследовательский университет ) Факультет прикладной математики и физики Кафедра Вычислительной математики и программирования 21