Начать свой блог хотелось бы с рассказа об одном замечательном технологическом новшестве. Это новшество - архитектура CUDA.
Введение
Вобщем-то не для кого не секрет, что видеокарта - это своего рода маленький компьютер - в ней есть своя оперативная память, свой процессор, свой блок адресации и даже свой вентилятор. Процессоры в большинстве видеокарт мощные, но не универсальные - вы не можете запрограммировать этот процессор выполнять произвольные операции - только те программы, которые заложил в видеокарту производитель.
Второй немаловажный момент - память видеокарты подключена к общей шине памяти. Благодаря этому центральный процессор может работать с памятью видеокарты как с обычной ОЗУ. Кроме того, видеокарта может копировать данные из общей озу в свою озу без участия центрального процессора.
CUDA - архитектура видокарты, которая позволяет выполнять на процессоре видеокарты произвольную программу. Написанную разумеется на своем машинном языке, отличном от того который понимает цп
CUDA реализован в некоторых видеокартах NVIDIA. Несмотря на то, что CUDA реализован в видеокарте, эта технология во многом предназанчена не для игр, и даже не обязательно для обработки графики. В принципе, применять эту технологию можно везде, где речь идет о серийной и однотипной обработке большого объема данных. При близкой стоимости видеокарт и цп, видеокарта может давать потрясающие выигрыши в производительности.
Отличия
Вкратце отличие cuda от x86 довольно простое - x86 это однопоточная архитектура, cuda - нет. Когда вы пишете скажем для core 2 quard, вы даже не задумываетесь о том, что у него 4 ядра, а не одно. Ядра в процессорах x86 очень мало зависимы - у них есть свои регистры, они могут выполнять разные инструкции и даже работать в разных виртуальных адресных пространствах. Фактически процессор с двумя ядрами или два разных процессора на одной шине - это для прикладного программиста одно и то же.
CUDA - изначально многопоточная архитектура. В процессоре CUDA много ядер. И написать программу CUDA, не понимая, механизма взаимодействия ядер нельзя.
Аппаратная архитектура
Весь "камень" видеокарты делится на так называемые мультипроцессоры(MP).
Мультипроцессор - это некая независимая единица, которую можно сравнить с ядром процессора. Согласно документации, на данный момент видеокарты nvidia имеют от 1 до 64 мультипроцессоров. Каждый мультипроцессор состоит из некоего числа скалярных процессоров(SP), имеет специальный процессор для шедулинга потоков между SP а также "набортную" память трех типов - rw-память "shared memory" и две ro-памяти texture и constant. Все три типа памяти являются общими(shared) для каждого SP на мультипроцессоре. shared memory может использоваться как место обмена информацией между SP, а texture и constant могут выступать в роли ручного кэша памяти устройства. Важно отметить что память размещенная на мультипроцессоре - это вовсе не ОЗУ видеокарты. Эта память расположена на чипе MP и намного быстрее ОЗУ видеокарты.
Стоит особенно уделить внимание тому факту, что на MP есть отдельное устройство, называемое Instruction Unit(IU), отвечающее за распределение заданий между скалярными процессорами. Таким образом, вопрос планировки задач в мультипроцессоре решен на аппаратном уровне.
У каждого SP есть набор регистров. Однако в отличие от x86 количество регистров не фиксированно - в зависимости от задачи, IU выделяет некоторое количество набортной памяти MP под регистры SP.
У SP нету аналога регистра esp и как следствие, нету стэка. Таким образом на SP не может быть выполнен рекурсивный вызов.
Програмная архитектура
Если вы хоть раз писали многопоточное приложение для windows или nix, то знаете основные принципы:
1) Все потоки выполняются независимо друг от друга. Никак нельзя указать, в каком порядке должны выполнятся потоки.
2) Все потоки одинаковы по своей сути. У них могут быть разные приоритеты и прочие параметры, но никак нельзя сгруппировать потоки. Например никак нельзя заставить два потока выполняться одновременно и синхронно на двух ядрах или процессорах.
3) Синхронизация потоков возможна только через специальные системные вызовы. Поэтому любая синхронизация - это существенный оверхед.
В CUDA не один из этих постулатов не верен.
Каждый поток в CUDA выполняется на маленьком устройстве, называемом Scalar Processor (SP)
Потоки объединяются в блоки некоего фиксированного размера. Блоки в свою очередь объединяются в гриды(grid). Грид - это некое "задание" которая выполняется на видеокарте. Структура грида задается при его создании, и не может изменяться. Грид может быть запущен только один раз - будучи выполненным он отмирает.
Таким образом когда вам нужно вызвать некоторый критичный по времени блок кода вы делаете следующее:
1) [CPU] Подготавливаете в ОЗУ все необходимые входные данные
2) [CPU] Выделяете блок в памяти видеокарты. Это делается одним вызовом: cudaMalloc(...)
3) [CPU] Копируете входные данные из основного ОЗУ в ОЗУ видеокарты. Это также делается простым системным вызовом: cudaMemcpy(...). Более того, так как для копирования большого куска данных особого внимания со стороны CPU не требуется, вы можете запустить копирование в асинхронном режиме, а сами делать что нибудь еще.
4) [CPU] Когда копирование завершится нужно инициализировать и запустить грид. На ответвлении языка C для CUDA это выглядит примерно так:
- Код: Выделить всё
matrixMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
Где grid - переменная(вектор) указывающая на количество блоков в гриде, threads - переменная указывающая на количество потоков в блоке. Все что в скобках - это просто некоторые числовые параметры, которые в неизменном виде передаются в процедуру matrixMul.
Ах да, собственно matrixMul - процедура, которую мы с вами написали, и которая будет выполняться на видеокарте.
Увидев такую строчку компилятор просто вставит на ее место несколько системных вызовов.
Думаю очевидно, что раз matrixMul выполняется на устройстве, то ее также можно также запускать асинхронно.
Если вызов синхронный, то поток на CPU заморозится, вплоть до завершения matrixMul на устройстве.
5) Получив задание выполнить грид, драйвер CUDA перенесет его код на устройство(если нужно) и запустит механизм планировки. Кто именно занимается распределением блоков по мультипроцессорам, GPU или CPU я так и не нашел в документации. Так или иначе, на свободные мультипроцессоры начнут подаватся задания. Очередность, в котором мультипроцессоры будут получать блоки никак не контролирется.
6) [GPU] Получив задание выполнить блок, IU на мультипроцессоре выделит для каждого SP необходимое количество памяти под регистры. Так как IU знает какую процедуру он будет выполнять, и так как у нас нет вызова функции в привычном смысле, то он точно может определить, сколько памяти потребуется под регистры для выполнения указанного блока. Количество SP на мультипроцессоре обычно существенно меньше, чем количество логических трэдов в блоке. Вообще говоря, очередность выполнения потоков внутри блока также ничем не регулируется. Тем не менее, существуют правила, по которым следует писать код для трэда, что бы он выполнялся эффективно на SP.
Большую часть из этих правил мы оставим за рамками статьи, но упомянем одно наиболее интересное
Трэды, которые выполняются на SP одновременно выполняются на нем синхронно. Это пораждает следующую проблему - если в коде есть условный переход, и два трэда волею судеб пошли по разным веткам этого перехода, то сначала будут работать те SP, которые пошли по одной ветке, а остальные будут простаивать, а затем начнут выполняться SP которые пошли по другой ветке, в то время как те что шли по первой буду простаивать.
Таким образом полезно распределять трэды по блокам так, что бы все трэды внутри блока шли по возможности по одним и тем же ветвям условных переходов.
Выполнив весь блок, MP уведомит вызывающую сторону о том, что готов принять следующее задание.
7) [CPU] Когда все блоки будут выполнены, поток на CPU выполнявший matrixMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB); разблокируется.
8 ) [CPU] Если мы выполняли задание не для потехи ради, то имеет смысл скопировать результат работы грида в основную память и освободить занятую память на устройстве.
Кратикй итог
Программирование на CUDA вещь прежде всего занятная. К сожалению, ввиду привязки к аппаратной платформе, написание программ с использованием CUDA в ряде задач неприменимо. Тем не менее у технологии есть свои ниши - везде где конечные пользователи готовы приобрести специальное железо для поднятие производительности.
Adobe в фотошопе последней версии также реализовал некоторые критичные по производительности функции с применением описанной технологии. Поэтому некоторые задачи в фотошопе будут быстрее работать у пользователей видеокарт с чипсетами nvidia
Копирайты
Все права на картинки принадлежат nvidia
