?

Log in

No account? Create an account

December 26th, 2009 - Excelsior — LiveJournal

Dec. 26th, 2009

06:40 pm - Зачем нужны GPU? Зачем нужна CUDA?

Этот пост логически принадлежит гайду по CUDA, но из-за ограничений ЖЖ на максимальный размер поста я вынужден был разбить гайд на несколько частей: 1) Зачем нужны GPU? Зачем нужна CUDA? (сейчас вы читаете эту часть), 2) Программная модель CUDA, 3) Хардварная реализация CUDA, 4) CUDA-совместимые видеокарты.

Disclaimer.
  1. Здесь будут затронуты лишь архитектура и спецификации устройств, имеющих compute capability 1.0-1.3, а именно: 1) десктопных карточек GeForce 8, GeForce 9, GeForce 1xx и GeForce 2xx, 2) их мобильных аналогов, 3) их профессиональных или, как сейчас модно выражаться, HPC аналогов семейств Quadro и Tesla. Более подробно читаем в CUDA Programming Guide 2.3 (Appendix A. Technical Specifications).
  2. Соответственно, инфа в этом гайде неприменима к: 1) видеокартам GeForce 7 и ниже (у них кардинально иная микроархитектура), 2) видеокартам ATI/AMD (несмотря на сходства высокоуровневой модели программирования, на низком уровне эти GPU очень сильно отличаются от решений NVIDIA), 3) видеокартам GeForce 300 и Tesla 2xxx (они построены на несколько отличной архитектуре, известной под кодовым названием Fermi; о ней поговорим как-нибудь в другой раз).
  3. В этом посте будет очень сжатый саммари (минимум воды и ноль картинго), поэтому для более мягкого погружения в архитектуру CUDA рекомендую ознакомиться с тремя презентациями на эту тему: общая инфа про архитектуре GPU, описание CUDA - много топиков сразу, но, на удивление, с хорошей детализацией, оптимизация алгоритмов для CUDA - там все последовательно, с необходимой для фтыкания с нуля избыточностью информации и с пачкой картинго. Не стоит забывать и про референсные документы, но они будут потяжелее для восприятия: CUDA Programming Guide 2.3, Best Practices - CUDA 2.3, CUDA PTX ISA 1.4.
  4. Здесь не будет информации о результатах на реальных задачах, ибо они индивидуальны для каждого алгоритма.
  5. Наконец, этот гайд - work in progress, поэтому что-то в нем может быть неточным, а что-то тупо неправильным. Со временем гайд будет расширяться и дополняться. Замечания и дополнения по делу всяко приветствуются. Вопросы о непонятных местах приветствуются еще больше (разрешены анонимные каменты, такчт задать вопрос очень легко - а мне ответить будет только в радость)!

Зачем нужны GPU? Зачем нужна CUDA?

На сегодняшний день графические процессоры (GPU) представляют собой самое яркое воплощение архитектуры SIMD. Их ALU многочисленны и быстры (см. сравнение пиковой single-precision floating point производительности CPU и GPU NVIDIA (2003-2008й годы)), их память обладает впечатляющей пропускной способностью (см. сравнение пропускной способности DRAM CPU и GPU NVIDIA (2003-2007й годы)), а темпы роста производительности являются самыми значительными в индустрии (сравните уклон двух приведенных графиков для CPU и GPU).

Несмотря на провоцирующую букву G (graphical) в названии, GPU хорошо подходят не только для расчета и отрисовки графики, но и для выполнения многих (но не всех!) general-purpose вычислительных задач (отсюда и название области исследований - GPGPU). Производительность GPU на выполнении подходящих алгоритмов в десятки и сотни раз превосходит производительность современных CPU (примеры ускорения алгоритмов на GPU NVIDIA можно посмотреть на страничке CUDA Zone - там еще можно поиграться с чекбоксами в категории "Filter by Application Type").

Впечатляющая вычислительная мощь GPU достигается за счет ограничений программной модели и неудобства программирования. Первую из этих проблем устранить невозможно, ибо это фундаментальный trade-off в проектировании железа - та цена, которую надо платить за непревзойденную производительность (больше на эту тему можно прочитать вот тут). Вторую проблему призвана устранить CUDA, отходя от канонического подхода к выражению параллелизма по данными при помощи векторных инструкций (как, например, сделано в SSE) и предоставляя более удобную модель программирования массово параллельных (massively-parallel) SIMD-устройств (подход CUDA концептуально не оригинален и во многом похож на шейдерные языки HLSL/GLSL, но зато предоставляет существенное расширение программной модели).

Tags:

06:41 pm - Программная модель CUDA

Этот пост логически принадлежит гайду по CUDA, но из-за ограничений ЖЖ на максимальный размер поста я вынужден был разбить гайд на несколько частей: 1) Зачем нужны GPU? Зачем нужна CUDA?, 2) Программная модель CUDA (сейчас вы читаете эту часть), 3) Хардварная реализация CUDA, 4) CUDA-совместимые видеокарты.

Disclaimer.
  1. Здесь будут затронуты лишь архитектура и спецификации устройств, имеющих compute capability 1.0-1.3, а именно: 1) десктопных карточек GeForce 8, GeForce 9, GeForce 1xx и GeForce 2xx, 2) их мобильных аналогов, 3) их профессиональных или, как сейчас модно выражаться, HPC аналогов семейств Quadro и Tesla. Более подробно читаем в CUDA Programming Guide 2.3 (Appendix A. Technical Specifications).
  2. Соответственно, инфа в этом гайде неприменима к: 1) видеокартам GeForce 7 и ниже (у них кардинально иная микроархитектура), 2) видеокартам ATI/AMD (несмотря на сходства высокоуровневой модели программирования, на низком уровне эти GPU очень сильно отличаются от решений NVIDIA), 3) видеокартам GeForce 300 и Tesla 2xxx (они построены на несколько отличной архитектуре, известной под кодовым названием Fermi; о ней поговорим как-нибудь в другой раз).
  3. В этом посте будет очень сжатый саммари (минимум воды и ноль картинго), поэтому для более мягкого погружения в архитектуру CUDA рекомендую ознакомиться с тремя презентациями на эту тему: общая инфа про архитектуре GPU, описание CUDA - много топиков сразу, но, на удивление, с хорошей детализацией, оптимизация алгоритмов для CUDA - там все последовательно, с необходимой для фтыкания с нуля избыточностью информации и с пачкой картинго. Не стоит забывать и про референсные документы, но они будут потяжелее для восприятия: CUDA Programming Guide 2.3, Best Practices - CUDA 2.3, CUDA PTX ISA 1.4.
  4. Здесь не будет информации о результатах на реальных задачах, ибо они индивидуальны для каждого алгоритма.
  5. Наконец, этот гайд - work in progress, поэтому что-то в нем может быть неточным, а что-то тупо неправильным. Со временем гайд будет расширяться и дополняться. Замечания и дополнения по делу всяко приветствуются. Вопросы о непонятных местах приветствуются еще больше (разрешены анонимные каменты, такчт задать вопрос очень легко - а мне ответить будет только в радость)!

Программная модель CUDA

В отличие от программируемых видеокарт пятилетней давности, видеокарты, реализующие программную модель CUDA, представляют собой полноценные сопроцессоры. Их ISA - PTX 1.4 экспозит богатую функциональность для работы с целыми числами, флоатами одинарной и двойной точности, а также поддерживает в полной мере scatter и gather или, проще говоря, случайные чтение и запись памяти. В зависимости от compute capability (о нем ниже в разделе про хардварную реализацию) конкретный GPU может не поддерживать некоторые сабсеты ISA, но корные инструкции работают везде.

По классификации авторов CUDA принадлежит к категории SIMT (Single Instruction Multiple Threads), хотя по факту это просто прикольная вариация SIMD - очень большой ширины (т.е. обрабатывающая много данных за раз) и выраженная в экстравагантной форме. К слову, как бы кому ни хотелось, CUDA не является ни SPMD, ни MIMD - во-первых, потому что на GPU может одновременно выполняться только одна программа, а, во-вторых, потому что треды кернела выполняются в локстепе.

Вместо экспоуза в программной модели широких инструкций (как это сделано в SSE), CUDA выражает параллелизм по данным в одновременном выполнении десятков тысяч потоков (точные цифры зависят от конкретной модели GPU), которые работают с инструкциями одинарной ширины и выполняются в локстепе пачками по 32 (такие пачки называются варпами). Выполнением в локстепе треды CUDA отличаются от тредов OS, которые, даже будучи порождены одинаковыми, выполняются в различном темпе, определяемом планировщиком OS. Экстравагантность программной модели CUDA дает ей следующие преимущества в сравнении с SSE: 1) удобнее записывать программу, т.к. программист нигде явно не указывает ширину инструкций, 2) программы отлично скейлятся вместе с количеством исполняющих устройств и их реальной шириной.

Вычисления в CUDA организуются в следующем виде. Программист пишет алгоритм в виде кернела (тело цикла), определяет индексное пространство (переменные цикла и границы их итераций), после чего отдает получившийся цикл на исполнение GPU. Координатное пространство потоков является трехмерным. Оно разбито на трехмерные блоки (называемые CTA - cooperative thread arrays), а каждый из блоков собственно содержит единичные треды. Каждый тред знает свой ID внутри блока (threadIdx) и ID блока в индексном пространстве (blockIdx), что позволяет поделить пачку данных между отдельными тредами. Двухуровневость координатного пространства обусловлена хардварной начинкой, которую я опишу ниже.

Хоть массово параллельные архитектуры типа SIMD наилучшим образом работают для полностью независимых потоков, CUDA предоставляет способы для организации совместной работы. Взаимодействие между потоками реализуется при помощи общей памяти (см. еще ниже про иерархию памяти). Синхронизировать совместную работу можно двумя способами: 1) при помощи примитивов синхронизации внутри блока (синхронизация между тредами разных блоков невозможна), 2) при помощи интерлокед операций с памятью (поддерживаются не всеми устройствами).

Модель памяти CUDA

Память у GPU совершенно отдельна от системной RAM (т.е. входные параметры для кернелов необходимо копировать из RAM CPU в RAM GPU, а результаты вычислений необходимо перегонять обратно).

Иерархия памяти GPU организована следующим образом: 1) у каждого треда есть приватная рабочая область в массиве регистров, 2) у каждого блока есть рабочая область, разделенная между тредами блока (т.н. shared память), 3) наконец, на девайсе есть глобальная память.

Как и CPU RAM, неприватная память совсем не обязана быть консистентной в условиях совместного доступа нескольких потоков к одной и той же ячейке памяти. Так же, как и в случае с CPU RAM, для обеспечения гарантированного результата совместного чтения-записи необходимо явным образом использовать синхронизацию (при помощи примитивов синхронизации или при помощи интерлокед операций - см. выше насчет деталей).

Спецификация PTX 1.4 определяет и другие типы памяти - они размещаются в одной из трех плоскостей вышеуказанной иерархии, но имеют особую семантику. Вот они: 1) sreg (специальные регистры, хранящие read-only системные значения, см. раздел "Chapter 8. Special Registers" слинкованного выше документа), 2) param (параметры кернела, могут располагаться либо в регистрах, либо в глобальной памяти - на усмотрение конкретной имплементации), 3) константы (read-only память, оптимизированная для 1d-локальности; маппится на глобальную память), 4) текстуры (read-only память, оптимизированная для 2d-локальности; маппится на глобальную память), 5) поверхности (не реализованы в хардваре), 6) приватные данные тредов, не поместившиеся в регистры, и пролитые (spilled) компилятором в глобальную память (т.н. local память).

Ограничения программной модели CUDA

1. GPU в CUDA - это не самостоятельное устройство, а сопроцессор. Постановку задач, аллокацию памяти, перегонку данных между CPU RAM и GPU RAM - все это делает хостовый CPU при помощи драйвера. Из этого, например, следует, что GPU не могут общаться друг с другом или в процессе выполнения подгружать данные из CPU RAM.

2. Так как аллокацией должен обязательно управлять хост, в кернелах может использоваться лишь статическая аллокация памяти. Единственный способ реализовать динамическое выделение памяти - заранее предвыделить пул в глобальной памяти и написать свой маллок, который будет выполняться на GPU. Однако такой подход весьма сомнителен из-за высокой латентности доступа к глобальной памяти (об этом ниже).

3. У GPU нет стека, поэтому параметры вызовов и локальные переменные подпрограмм размещаются в регистрах. Кроме того, размещение параметров происходит статически на этапе компиляции - один-единственный раз для каждого коллсайта. Из последнего следует тот факт, что рекурсия не поддерживается.

4. В PTX 1.4 задекларирована возможность динамического управления control flow, т.е. перехода на адрес, указываемый регистром, и вызова подпрограммы, адрес которой указывается регистром, но по факту в хардваре эта возможность не реализована.

5. В CUDA нет понятия об эррор-кодах или иксепшнах, поэтому обнаружение и пропагация исключительных состояний должны реализовываться руками. Впрочем, в PTX 1.4 есть инструкции exit (нормальное безусловное завершение работы треда) и trap (аварийный обрыв исполнения кернела с вызовом прерывания CPU).

6. Так как треды выполняются в локстепе (т.е. каждый такт все треды варпа - атомарной единицы исполнения, состоящей из 32 тредов - выполняют одну и ту же операцию), то бранчинг чаще всего убивает перфоманс (абстракция протекает!). Если бранчинг вызывает дивергенцию путей исполнения внутри варпа, то из-за локстепа каждый тред исполняет обе ветки каждого бранча - и ложную, и правдивую (ессно, на корректности алгоритма это не отражается, так как при выполнении "не своей" ветки бранча тред ничего никуда не пищет). Больше вложенность бранчей - большее падение перфоманса.

7. Три уровня иерархии памяти имеют каждый свои инструкции чтения/записи, поэтому поинтеры ниразу не прозрачные. Кроме того, регистровая память (а ее, в отличие от случая с CPU, очень много) вообще не может быть адресована динамически сформированным поинтером. Плюнуть на все и размещать все данные в глобальной памяти не получится, ибо с точки зрения удара по перфомансу это infeasible. Отсюда следует, что рантайм-полиморфизм не поддерживается.

8. Эпичнее всего абстракция протекает в случае с моделью памяти. Кагбэ модель памяти GPU очень похожа на обычную - тоже организована в иерархию, тоже есть кэши, ну еще и небольшой флейвор в виде текстур, а так вроде бы то же самое. На самом деле это совсем не так - по факту хардварная реализация памяти GPU качественно отличается от хардварной реализации памяти CPU: 1) у GPU размер кэшей очень небольшой (десятки килобайт шаред + константной + текстурной памяти против мегабайтов кэша у современных CPU), поэтому у GPU в общем случае вероятность cache hit очень низкая, а у CPU, наоборот - очень высокая, 2) DRAM (глобальная память) видеокарт отличается от системной DRAM тем, что частота у нее выше, а латентности хуже, 3) даже у регистров GPU есть неслабая латентность в 24 такта в read-after-write сценарии. Впрочем, эти факты в большинстве случаев облегчаются тем, что, благодаря большому количеству выполняемых потоков, высокую латентность можно амортизировать - значительным образом или даже полностью.

Tags:

06:42 pm - Хардварная реализация CUDA

Этот пост логически принадлежит гайду по CUDA, но из-за ограничений ЖЖ на максимальный размер поста я вынужден был разбить гайд на несколько частей: 1) Зачем нужны GPU? Зачем нужна CUDA?, 2) Программная модель CUDA, 3) Хардварная реализация CUDA (сейчас вы читаете эту часть), 4) CUDA-совместимые видеокарты.

Disclaimer.
  1. Здесь будут затронуты лишь архитектура и спецификации устройств, имеющих compute capability 1.0-1.3, а именно: 1) десктопных карточек GeForce 8, GeForce 9, GeForce 1xx и GeForce 2xx, 2) их мобильных аналогов, 3) их профессиональных или, как сейчас модно выражаться, HPC аналогов семейств Quadro и Tesla. Более подробно читаем в CUDA Programming Guide 2.3 (Appendix A. Technical Specifications).
  2. Соответственно, инфа в этом гайде неприменима к: 1) видеокартам GeForce 7 и ниже (у них кардинально иная микроархитектура), 2) видеокартам ATI/AMD (несмотря на сходства высокоуровневой модели программирования, на низком уровне эти GPU очень сильно отличаются от решений NVIDIA), 3) видеокартам GeForce 300 и Tesla 2xxx (они построены на несколько отличной архитектуре, известной под кодовым названием Fermi; о ней поговорим как-нибудь в другой раз).
  3. В этом посте будет очень сжатый саммари (минимум воды и ноль картинго), поэтому для более мягкого погружения в архитектуру CUDA рекомендую ознакомиться с тремя презентациями на эту тему: общая инфа про архитектуре GPU, описание CUDA - много топиков сразу, но, на удивление, с хорошей детализацией, оптимизация алгоритмов для CUDA - там все последовательно, с необходимой для фтыкания с нуля избыточностью информации и с пачкой картинго. Не стоит забывать и про референсные документы, но они будут потяжелее для восприятия: CUDA Programming Guide 2.3, Best Practices - CUDA 2.3, CUDA PTX ISA 1.4.
  4. Здесь не будет информации о результатах на реальных задачах, ибо они индивидуальны для каждого алгоритма.
  5. Наконец, этот гайд - work in progress, поэтому что-то в нем может быть неточным, а что-то тупо неправильным. Со временем гайд будет расширяться и дополняться. Замечания и дополнения по делу всяко приветствуются. Вопросы о непонятных местах приветствуются еще больше (разрешены анонимные каменты, такчт задать вопрос очень легко - а мне ответить будет только в радость)!

Высокоуровневая картина

Железо GPU заточено на выполнение следующих двух императивов любой ценой: 1) максимизацию пиковой производительности вычислительных устройств, 2) максимизацию пропускной способности памяти (1 и 2 могут выполняться одновременно, так как ALU и память не зависят друг от друга). Из этого следует то, что: 1) если можно слабать больше ALU (они же, в терминологии CUDA называются SP), но на меньшей частоте, то лабаем больше ALU, 2) если можно увеличить пропускную способность памяти за счет увеличения латентности, то увеличиваем, 3) мы отказываемся от предсказателя ветвлений и out-of-order выполнения для того, чтобы освободить место на кристалле для еще большего количества ALU. Схематическая разница между кристаллами CPU и GPU NVIDIA представлена на картинке из CUDA Programming Guide 2.3.

По сей день GPU успешно справляются с наращиванием чистой вычислительной мощи (обе картинки взяты из CUDA Programming Guide 2.3): сравнение пиковой single-precision floating point производительности CPU и GPU NVIDIA (2003-2008й годы), сравнение пропускной способности DRAM CPU и GPU NVIDIA (2003-2007й годы). За эту мощь приходится платить низкой малопоточной производительностью и высокими задержками памяти, что делает GPU пригодными лишь для специального (но большого) класса вычислительных задач, но зато с ними GPU справляются просто великолепно.

Важной идиомой хардварной реализации CUDA является скрытие латентности памяти очень большим количеством исполняемых потоков. Дело в том, что: 1) скедулер CUDA GPU достаточно умен, чтобы сменять потоки, ожидающие результатов операции с памятью, на потоки, готовые к исполнению, 2) контроллер памяти GPU достаточно продвинут для того, чтобы работать параллельно арифметическим устройствам. Поэтому, благодаря большому количеству одновременно исполняющихся потоков, скедулер иногда может полностью скрыть латентность обращения к памяти (это работает часто, но не всегда - детали в разделе про хардварную реализацию). Благодаря этому, CUDA GPU могут себе позволить тратить транзисторы не на большие кэши (как это делают CPU), а на большое количество ALU.

CUDA-совместимые графические процессоры NVIDIA

GPU, реализующие модель CUDA, представляют собой пачку single-precision floating point числодробилок (streaming processors, SP), объединенных по восемь штук в так называемые SM (streaming multiprocessors). Кроме того, в моделях GPU на ядре GT200 на каждом SM стоит по одному double-precision floating point устройству.

Каждый SM имеет следующее количество on-chip памяти:
  * N (обычно 8192, для GPU на ядре GT200 - 16384) 32-битных регистров.
  * 16 килобайт шаренной памяти.
  * 8 килобайт кэша для константной памяти.
  * 8 килобайт кэша для текстурной памяти.

Кроме того, все SM шарят доступ к on-device (т.е. off-chip) глобальной памяти размером до 4 Gb (ее еще называют "видеопамять"). В глобальной памяти кроме обычных данных могут находиться до 16 килобайт локальной памяти (per-thread), до 64 килобайт констант (per-kernel), а также некоторое количество текстур (per-kernel). Текстуры могут иметь форму 1d, 2d или 3d. Суммарный размер адресуемых текстур не ограничен, но каждая конкретная текстура ограничена в размерности - от 2^11 по каждому измерению для 3d-текстур, до 2^27 в наилучшем случае для 1d-текстур. Детали по размеру текстур читаем в CUDA Programming Guide 2.3.

Исполнение кернела

Исполнение кернела происходит следующим образом:
  * Хостовый CPU с помощью драйвера аллоцирует области в глобальной памяти видеокарты и копирует входные данные из RAM в GPU RAM.
  * Хостовый CPU передает драйверу: 1) текст кернела - инструкции PTX-ассемблера в текстовом виде, 2) в состав PTX также входят требования кернела к памяти – количество регистров и объем шаренной памяти (они должны быть известны заранее, т.е. динамическая аллокация внутри кернела не поддерживается!), 3) характеристики индексного пространства.
  * Если на видеокарте уже выполняется кернел, то драйвер ждет его завершения. Несколько кернелов одновременно исполняться не могут.
  * Драйвер видеокарты валидирует переданные ему параметры: 1) убеждается в том, что в блоке индексного пространства не более 512 тредов, 2) проверяет то, что каждая размерность индексного пространства не превосходит соответствующую размерность блока более, чем в 65536 раз, 3) убеждается в том, что каждый SM видеокарты может исполнить хотя бы один блок кернела, т.е. проверяет на положительность величину Blocks_per_SM = Min(Registers_per_SM/(Registers_per_kernel*Threads_per_block), Shared_memory_per_SM/Shared_memory_per_block, (Max_warps_per_SM*Threads_per_warp)/Threads_per_block, Max_blocks_per_SM), где значения Threads_per_block, Registers_per_kernel и Shared_memory_per_block являются константами кернела, а Registers_per_SM, Shared_memory_per_SM, Max_warps_per_SM, Threads_per_warp, Max_blocks_per_SM являются константами хардвары и варьируются в зависимости от конкретного GPU (16384/16384/24/32/8 для ядра GT200, 8192/16384/16/32/8 для остальных), 4) проверяет то, что длина кернела не превосходит 2 миллиона инструкций.
  * После валидации драйвер выполняет JIT-компиляцию кернела и передает управление GPU.
  * После этого планировщик GPU делит блоки индексного пространства между SM и запускает кернел на выполнение. Каждый блок выполняется от начала и до конца на единственном SM, т.е. блоки между SM не мигрируют.
  * Каждый SM в локстепе выполняет инструкции с гранулярностью в один варп (32 треда). Если для выполнения очередной инструкции варпу необходимы значения из памяти, то, пока он их ждет, на выполнение запускается следующий варп и так далее. Так как варпов обычно много, то таким образом неплохо скрывается латентность доступа к памяти (см. конкретные цифры ниже). Но, как указано выше, количество параллельно выполняемых тредов ограничено объемом регистровой и шаренной памяти SM, поэтому, чем тяжелее алгоритм кернела, тем больше он будет страдать от латентности обращений к памяти.
  * После завершения работы кернела драйвер оповещает хостовый CPU и тот копирует результаты из GPU RAM в RAM.

Особое внимание хочу обратить на величину Blocks_per_SM. Она важна не столько как абстрактный показатель параллельности алгоритма (все равно больше 8 операций за такт SM выполнять не может - см. детали ниже), а как размер буфера для амортизации латентности операций с памятью (про эту идиому CUDA я упоминал выше).

Инструкции: список

Здесь будет только сжатый саммари - полные детали в документе CUDA PTX ISA 1.4. Итак, инструкции PTX 1.4 делятся на 9 групп:
  * Integer arithmetic: add, sub, addc (учитывает CF), subc (учитывает CF), mul, mad (multiply and add, res = a * b + c), mul24 (работает с 24-битными интами), mad24 (работает с 24-битными интами), sad (sum of absolute differences, res = c + |a - b|), div, rem, abs, neg, min, max.
  * Floating point: add, sub, mul, fma (fused multiply and add; то же самое, что и mad, но без округления между * и +), mad (multiply and add, res = a * b + c), div, abs, neg, min, max, rcp (reciprocal, res = 1/x), sqrt, rsqrt (reciprocal of square root, res = 1/sqrt(x)), sin, cos, lg2, ex2.
  * Comparison and selection: set (сравнивает значения и записывает результат в обычный регистр), setp (сравнивает значения и записывает результат в предикатный регистр), selp (тернарный оператор, юзающий значение в обычном регистре как условие), slct (тернарный оператор, юзающий знак значения в обычном регистре как условие (>= 0 это true, все остальное - false)).
  * Logic and shift: and, or, xor, not (полный негатив битов), cnot (логическое отрицание в стиле C: res = (a == 0) ? 1 : 0), shl, shr.
  * Data movement and conversion: mov, ld, st, cvt (режимы адресации для data movement инструкций описываются ниже).
  * Texture: tex (специальная инструкция для выборки из текстурной памяти).
  * Control flow: @ (префикс к остальным инструкциям - говорит о том, что подлежащую инструкцию необходимо исполнить только если предикатный регистр содержит определенное значение), bra (безусловный бранчинг, с помощью @ можно получить безусловный бранчинг), call (см. комментарии ниже), ret, exit (безусловное окончание работы треда).
  * Parallel synchronization and communication: bar (marks the arrival of threads at a barrier and waits for all other threads to arrive), membar (thread execution resumes after a membar when the thread's prior memory writes are visible to other threads at the specified level, and memory reads by this thread can no longer be affected by other thread writes), atom (атомарная бинарная операция над данными в нерегистровой памяти - только для интов), red (неатомарная бинарная операция (редукция - отсюда и название "red") над данными в нерегистровой памяти - только для интов), vote (performs a reduction of the source predicate (.all, .any, .uni - последнее означает uniform) across threads in a warp).
  * Miscellaneous: trap (безусловное окончание работы кернела и вызов ловушки CPU), brkpt (сигнализация о брейкпоинте), pmevent (индикация одного из 16 событий performance monitor).

1. Фундаментальные типы данных, с которыми умеет работать PTX: 1) знаковые и беззнаковые инты 8, 16, 32, 64 бит, 2) флоаты 16 (half-precision), 32 (single-precision), 64 (double-precision) бит, 3) битовые типы (сырые данные без всякой семантики) 8, 16, 32, 64 бит, 4) предикатные типы.

2. Из них составляются сложные типы данных программной модели: 1) векторы длиной 2 и 4 (общим размером не более 128 бит), 2) массивы константной размерности (динамическая аллокация не поддерживается!), 3) структуры и юнионы аля C. Сложные типы данных могут быть вложены друг в друга. Ессно, сложные типы данных не поддерживаются в ISA - работа с ними транслируется в операции с фундаментальными типами данных.

3. Концептуальное отличие от x86 - локальные переменные кернела за редким исключением размещают в регистрах (PTX имеет неограниченное число виртуальных регистров, которым, к тому же, можно давать символьные имена). Так надо делать потому, что: 1) исполнительные устройства SM имеют большое число 32-битных регистров SM (обычно 8192, для GPU на ядре GT200 - 16384), поэтому места хватит всем, 2) размещение данных в глобальной памяти чревато большими латентностями (см. ниже в разделах про память). Единственная причина проливать (spill) данные из регистров в память - это уменьшение register pressure для того, чтобы на каждом SM могло одновременно исполняться большее число тредов (зачем и когда это нужно см. выше в разделе про высокоуровневую модель хардвары).

4. Режимы адресации, поддерживаемые инструкциями: 1) значение в регистре, 2) указатель, лежащий в регистре, 3) указатель, лежащий в регистре +/- константное смещение, 4) указатель, значение которого известно на этапе компиляции. Для режимов 2-4 при помощи значения в регистре может указываться смещение (это реализует поддержку сложных типов данных). Пример: ld.global.u32 s, a[N-1] (в этом примере s, a и N - регистры).

5. Внимание! PTX не поддерживает указатели на регистровую память. Кроме того, каждая из областей памяти (полный список см. в разделе "Модель памяти CUDA") имеет свои собственные указатели, т.е. любая инструкция, использующая указатели (т.е. ld или st) требует указания области памяти, например: ld.const или st.shared.

6. Несмотря на то, что PTX 1.4 допускает косвенную адресацию для операндов инструкций bra и call, по факту она не реализована в хардваре, поэтому бранчи и вызовы подпрограмм возможны только для смещений, известных на этапе компиляции.

7. Наряду с бранчингом для записи условной логики можно использовать гварды - специальные префиксы, которые говорят о том, что подлежащую инструкцию необходимо исполнить только если предикатный регистр содержит определенное значение. Пример:

setp.lt.s32 p, i, n; // p = (i < n)
@p add.s32 j, j, 1; // if i < n, add 1 to j

8. У GPU нет стека, поэтому параметры вызовов и локальные переменные подпрограмм размещаются в регистрах. Кроме того, размещение параметров происходит статически на этапе компиляции - один-единственный раз для каждого коллсайта. Из последнего следует тот факт, что рекурсия не поддерживается.

Инструкции: производительность

Производительность арифметических операций такова - за один такт каждый SM может выполнить:
  * 8 простых single-precision FP операций (add, mul, mad/fma).
  * Несколько сложных single-precision FP операций в облегченном режиме точности (т.е. не соответствующем IEEE 754 – детали см. в CUDA Programming Guide 2.3), а именно: 2 операции rcp (reciprocal aka 1/x), 2 операции rsqrt (reciprocal square root aka 1/sqrt(x)), 2 операции log, 1 операцию sin/cos, 1 операцию exp. Остальные операции являются производными от вышеуказанных, например: sqrt(x) = x * rsqrt(x) -> 1.6 операций за такт, x/y = x * rcp(y) -> 1.6 операций за такт и так далее.
  * Single-precision FP операции с точностью, совместимой с IEEE 754, стоят на порядок больше тактов + заодно еще некоторые из них дополнительно аллоцируют локальную память = дальнейшее падение перфоманса из-за высокой латентности такой памяти.
  * GPU на ядре GT200 кроме того могут выполнять double-precision FP операции в темпе 1/12 от SP операций (эта цифра не равна 1/8 из-за различия в функциональности SP и DP устройств – сам пока не понял почему). На практике из-за разницы в загруженности SP и DP устройств (так как последних гораздо меньше) относительное падение DP-производительности несколько более терпимо – NVIDIA утверждает, что оно вообще колеблется около 1/4, но, я подозреваю, что здесь кроется подстава. Тут еще есть загон по поводу того, что иногда эмуляция double-precision несколькими single-precision операциями может быть круче нативного double-precision, но с этим я тоже пока что не разобрался. Наконец, double-precision операции всегда выполняются в режиме точности, соответствующем стандарту IEEE 754.
  * 8 операций сложения, 2 операции умножения, 8 битовых операций над 32-битными целыми числами.
  * Для ускорения умножения целых чисел предлагается использовать 24-битное умножение, которое работает с частотой 8 операций за такт.
  * Точный перфоманс деления нацело и остатка от деления для целых чисел в спеке не приводится, только есть просьба по возможности юзать шифты ибо иначе будет зело медленною
  * Кроме того, операции конверсии между типами (16 бит -> 32 бита, double-precision -> single-precision) работают с частотой 8 операций за такт
  * Наконец, операции min, max и compare выполняются по 8 штук за такт независимо от типа операндов.

Производительность инструкций доступа к памяти (чтения или записи) - 8 операций за такт (доступ к регистрам бесплатный). Эти значения не учитывают латентность (см. ниже) и предполагают оптимальные условия использования памяти (отсутствие конфликтов банков, оптимальное коалесцирование и т.д. - об этом читаем ниже).

Получается простое мнемоническое правило - ввиду того, что у каждого SM имеется по 8 SP, то большинство инструкций выполняется по 8 за такт. Исключения: 1) сложные single-precision операции (деление, степени, тригонометрия, логарифмы, экспоненты), 2) double-precision операции, ибо SP не умеют лабать в двойной точности, 3) умножение 32-битных интов, ибо оно эмулируется флоатовым умножением, а у single-precision флоатов на мантиссу отводится только 24 бита, 4) деление интов, ибо оно эмулируется флоатовым делением, а последнее является сложной операцией (см. пункт 1).

Производительность памяти: латентность

Латентность обращения к памяти (в пересчете на такты GPU):
  * К регистру - 0 тактов (24 такта в случае read-after-write; для полного скрытия этой латентности надо не менее 192 одновременно исполняющихся тредов для каждого SM - это следует из того, что варп исполняется как минимум 4 такта (максимальный перфоманс инструкции - 8 операций за такт - см. выше), поэтому, чтобы покрыть 24 такта латентности, для верности надо 6 варпов = 192 треда).
  * К шаренной памяти - 0 тактов.
  * К константной памяти - 0 тактов в случае cache hit, 400-600 тактов в случае cache miss.
  * К кэшу текстур - 400-600 тактов в независимости от того, попали мы в кэш или нет (хит все равно полезен, ибо уменьшает нагрузку на шину видеопамяти).
  * К глобальной памяти - 400-600 тактов.

Для on-chip памяти указана сказочная латентность в 0 тактов, понимать которую надо как амортизированное время доступа. Перфоманс on-chip памяти в оптимальных условиях равен 2 такта на полуварп, что равняется перфомансу самой быстрой инструкции - 8 операций за такт, поэтому латентность легко прячется выполнением операций другими тредами.

Относительная латентность глобальной памяти отмечена неопределенным количеством тактов, так как значительно зависит от частоты GPU и характеристик конкретных модулей памяти. Частота GPU варьируется от 900 Mhz у GeForce 8300 GS до 1.5 Ghz у GeForce GTX 285. Модули памяти, используемые в видеокартах также варьируются - от GDDR2 до GDDR3. Разброс большой, поэтому трудно делать какие-то адекватные обобщения.

Все значения, указанные выше, (в том числе и нули для on-chip) описывают теоретический максимум, предполагая оптимальные условия использования памяти (отсутствие конфликтов банков, оптимальное коалесцирование и т.д. - об этом читаем ниже). В реальности перфоманс может оказаться хуже.

Производительность памяти: пропускная способность

Подсчет ПСП для on-chip памяти имеет мало смысла, ибо ее частота привязана к частоте GPU, поэтому анализ перфоманса для такой памяти можно и, следовательно, нужно проводить в тактах. Соответственно, конкретные цифры для этого типа памяти приводить не буду - их несложно посчитать, но нафига.

Пропускная способность глобальной памяти зависит от частоты памяти и ширины шины и очень сильно варьируется - от 6.4 GB/s для GeForce 8300 GS (800 Mhz DDR2 x (64 bit bus / 8)) до 160 GB/s для GeForce GTX 285 (2.5 Ghz GDDR3 x (512 bit bus / 8)). Впрочем, надо помнить, что эти значения описывают теоретический максимум, предполагая оптимальные условия использования памяти (отсутствие конфликтов банков, оптимальное коалесцирование и т.д. - об этом читаем ниже). В реальности перфоманс может оказаться хуже.

С материнской платой видеокарты соединены интерфейсом PCI-Express, содержащим 16 дорожек, каждая из которых выдает 250 MБ/c для v1.0 и 500 МБ/c для v2.0. На практике получается чуть больше трех гектар в секунду для материнки с PCI-Ex 1.0 и чуть больше пяти гектар в секунду для материнки с PCI-Ex 2.0.

Производительность памяти: ограничения на паттерны доступа

Канули в лету времена, когда программы для GPU могли работать лишь со строго упорядоченными линейными областями памяти - сегодня видеокарты CUDA в полной мере поддерживают рандомные чтение и запись для всех плоскостей иерархии памяти. Однако на паттерны доступа к памяти все еще накладываются серьезные ограничения, невыполнение которых чревато значительным падением перфоманса (в десять и более раз).

Запросы к любой памяти формируются с гранулярностью одного полуварпа (16 тредов). Если запрос от полуварпа не удовлетворяет ограничениям, которые накладываются типом памяти, то он выполняется медленнее (например, по этой причине SoA aka Structure of Arrays почти всегда лучше для перфоманса, чем AoS aka Array of Structures).

1. Все уровни памяти разбиты на банки. Так регистровая память разбита на ? банков непонятно каким образом, шаренная память разбита на 16 банков полосами по 32 бита, а глобальная память разбита на M (для GTX 285 - 8, для 8800 GTX - 6; я так подозреваю, это число равно битности шины, деленной на 64 (битность шины можно посмотреть в википедии)) банков полосами по 256 байт. Как я понял, каждый банк - одноканальное устройство и за один цикл своей работы (2 такта для шаренной памяти, ? тактов для глобальной памяти) может обработать только один запрос к памяти, поэтому одновременные обращения к одному и тому же банку (пусть даже и к разным полоскам памяти) будут сериализованы, т.е. для максимального перфоманса необходимо поровну загружать все банки памяти. Примечания:
  * Спека божится, что на конфликты банков регистровой памяти простые смертные повлиять никак не могут. Цитирую: "The compiler and thread scheduler schedule the instructions as optimally as possible to avoid register memory bank conflicts. They achieve best results when the number of threads per block is a multiple of 64. Other than following this rule, an application has no direct control over these bank conflicts. In particular, there is no need to pack data into float4 or int4 types". Немного больше на эту тему можно почитать на форуме NVIDIA в топике "Low down on avoiding register bank conflicts".
  * On-chip память обладает способностью броадкаста. Если весь полуварп читает из одной и той же полоски (кагбэ 16-way конфликт), то запрос обрабатывается за один цикл (2 такта). В спеке еще написано про броадкаст шаренной памяти для части тредов, но я нифига не понял как это работает.
  * В случае шаренной памяти конфликты банков проявляются по отдельности для каждого полуварпа, в случае же с глобальной памятью этот эффект имеет область действия всех исполняющихся в данный момент полуварпов, т.е. к глобальной памяти нужно согласовывать доступ не только внутри полуварпа, но и между блоками (пример на эту тему я видел в доке Optimizing CUDA в разделе "Partition Camping").

2. Чтобы достигнуть максимальной пропускной способности глобальной памяти, нужно полностью загружать шину памяти - понятно, что, каждый раз спрашивая по байтику у шины шириной в 512 бит, ничего хорошего мы не добьемся. Вот как это работает. В худшем случае для каждого запроса от полуварпа контроллеру памяти выдаются 16 транзакций - по одной от каждого треда, но, если обращения к памяти удовлетворяют определенному паттерну, то они могут быть склеены (coalesced) в меньшее число транзакций, вплоть до одной. О правилах коалесцирования подробно написано в специальном документе и официальных бест-практисах.

3. У константной и текстурной памяти (а точнее, у их on-chip кэшей) также есть особенности, без знания которых невозможно достичь максимального перфоманса. Кэш констант имеет нулевую латентность, но всего лишь 1 банк полосками по 32 бита (он, как и шаренная память, тоже поддерживает броадкаст) со всеми вытекающими последствиями. Кэш текстур имеет [высокую] латентность, равную латентности глобальной памяти, не имеет банков (хотя тоже поддерживает броадкаст), но зато оптимизирован под 2d-локальность (поэтому его использование полезно тогда, когда невозможно скоалесцировать запросы к глобальной памяти или удовлетворить требования 1d-локальности кэша констант).

Tags:

06:43 pm - CUDA-совместимые видеокарты

Этот пост логически принадлежит гайду по CUDA, но из-за ограничений ЖЖ на максимальный размер поста я вынужден был разбить гайд на несколько частей: 1) Зачем нужны GPU? Зачем нужна CUDA?, 2) Программная модель CUDA, 3) Хардварная реализация CUDA, 4) CUDA-совместимые видеокарты (сейчас вы читаете эту часть).

Disclaimer.
  1. Здесь будут затронуты лишь архитектура и спецификации устройств, имеющих compute capability 1.0-1.3, а именно: 1) десктопных карточек GeForce 8, GeForce 9, GeForce 1xx и GeForce 2xx, 2) их мобильных аналогов, 3) их профессиональных или, как сейчас модно выражаться, HPC аналогов семейств Quadro и Tesla. Более подробно читаем в CUDA Programming Guide 2.3 (Appendix A. Technical Specifications).
  2. Соответственно, инфа в этом гайде неприменима к: 1) видеокартам GeForce 7 и ниже (у них кардинально иная микроархитектура), 2) видеокартам ATI/AMD (несмотря на сходства высокоуровневой модели программирования, на низком уровне эти GPU очень сильно отличаются от решений NVIDIA), 3) видеокартам GeForce 300 и Tesla 2xxx (они построены на несколько отличной архитектуре, известной под кодовым названием Fermi; о ней поговорим как-нибудь в другой раз).
  3. В этом посте будет очень сжатый саммари (минимум воды и ноль картинго), поэтому для более мягкого погружения в архитектуру CUDA рекомендую ознакомиться с тремя презентациями на эту тему: общая инфа про архитектуре GPU, описание CUDA - много топиков сразу, но, на удивление, с хорошей детализацией, оптимизация алгоритмов для CUDA - там все последовательно, с необходимой для фтыкания с нуля избыточностью информации и с пачкой картинго. Не стоит забывать и про референсные документы, но они будут потяжелее для восприятия: CUDA Programming Guide 2.3, Best Practices - CUDA 2.3, CUDA PTX ISA 1.4.
  4. Здесь не будет информации о результатах на реальных задачах, ибо они индивидуальны для каждого алгоритма.
  5. Наконец, этот гайд - work in progress, поэтому что-то в нем может быть неточным, а что-то тупо неправильным. Со временем гайд будет расширяться и дополняться. Замечания и дополнения по делу всяко приветствуются. Вопросы о непонятных местах приветствуются еще больше (разрешены анонимные каменты, такчт задать вопрос очень легко - а мне ответить будет только в радость)!

Классификация

На сегодняшний момент к видеокартам, поддерживающим CUDA, относятся: 1) десктопные карточки GeForce 8, GeForce 9, GeForce 1xx и GeForce 2xx, 2) их мобильные аналоги, 3) их профессиональные аналоги семейств Quadro и Tesla.

CUDA-совместимые видеокарты во многом похожи друг на друга. Во-первых, программная модель одинакова для всех таких устройств. Кроме того, с момента представления первого CUDA GPU и по сей день не менялись размеры шарной памяти (16 КБ) и кэшей на SM (8 КБ для констант и 8 КБ для текстур), а также константы хардвары (типо SM = 8 SP, варп = 32 треда и так далее).

Различия между видеокартами CUDA делятся на две категории - качественные и количественные. Качественные различия определяются параметром Compute Capability, который может принимать одно из четырех значений: 1.0, 1.1, 1.2 и 1.3 (см. ниже). Количественные различия - это: 1) количество SM, 2) частота, на которой они работают, 3) размер видеопамяти, 4) частота видеопамяти, 5) ширина шины видеопамяти (тут все изменяется в широких пределах: от кастратов с 1 SM до монстров с 4x30 SM, ну и так далее).

Касательно Compute Capability (детально описанного в CUDA Programming Guide 2.3 (Appendix A. Technical Specifications)) все гораздо проще, чем в теории. CUDA-видеокарты делятся на две категории - с GPU на ядре GT200 (CC 1.3) и все остальные (CC 1.0 и 1.1). Отличия между CC 1.3 и CC 1.0 следующие: 1) каждый SM имеет в два раза больше регистров (16384 вместо 8192), 2) каждый SM имеет на треть увеличенный максимум одновременно выполняющихся варпов (32 вместо 24), 3) нативно поддерживается double-precision FP (в девайсах более низкой cc double-precision эмулируется несколькими single-precision операциями), 4) поддерживаются атомарные операции над 32- и 64-битными словами, 5) условия на коалесцирование запросов к памяти сильно упрощаются. CC 1.1 от CC 1.0 отличается незначительно - всего лишь поддержкой атомарных операций над 32-битными словами в глобальной памяти.

State of the art

Гигантское семейство GPU CUDA (которому уже более трех лет) венчается ядром GT200 (внимание, маркетинг! не все видеокарты GeForce 2xx построены на ядре GT200!). На основе этого GPU и его обрезков собраны потребительские видеокарты GeForce GTX 260-295 и high-performance решения Tesla 10xx.

Линейка high-end видеокарт на ядре GT200 открывается десктопной карточкой GTX 285 (видеокарты GTX 260 и GTX 260 Core 216 являются ее обрезками - т.е. содержат несколько бракованных и поэтому отключенных SM) и ее HPC-клоном Tesla S1050. Дальше GPU уже не улучшаются, а размножаются - GTX 295 представляет собой два GTX 285 на одной плате, а Tesla S1070 содержит целых четыре GTX 285.

Полноценное ядро GT200 имеет 30 SM, соответствующих compute capability 1.3 (см. выше), которые работают на частоте 1.3-1.5 Ghz (частоты варьируются в зависимости от конкретной модели). Частоты, на которых работает память в рассматриваемых устройствах, очень сильно отличаются: у C1060/S1070 они на уровне 1.6 Ghz (x512 бит шина = 102 GB/s), у GeForce 285 - 2.3-2.5 Ghz (x512 бит шина = 147.2-160 GB/s), у GeForce 295 - 2.0 Ghz (x512 бит шина = 128 GB/s).

Tags:

06:44 pm - Гайд по CUDA

В каментах к потсу про Конфлакс (кросскомпилятор параллельных алгоритмов из MSIL в CUDA/Stream/..., который мы с коллегами разрабатываем в ОИПИ НАН РБ) зашел разговор о теоретических возможностях CUDA, и я решил собрать воедино все материалы на эту тему.

К сожалению, где-то начиная с десятой версии, гайд превысил максимальный размер поста ЖЖ (что-то около 40000 юникод-символов), поэтому я разбил его на несколько логических частей. Вот ссылки:
   1) Зачем нужны GPU? Зачем нужна CUDA?
   2) Программная модель CUDA
   3) Хардварная реализация CUDA
   4) CUDA-совместимые видеокарты

Tags:
Previous day (Calendar) Next day