Category: компьютеры

Category was added automatically. Read all entries about "компьютеры".

glider

Trace Windows 7 boot/shutdown/hibernate/standby/resume issues

Наткнулся на интересный пост с форума MSFN. Оказывается, в Windows SDK есть Windows Performance Toolkit, который позволяет запрофайлить процессы загрузки, шатдауна и саспенда/хибернейта.

Профайлинг происходит следующим образом:
1) Скачиваем тулкит (до семерки его можно скачать с сайта MS, для семерки он есть только на дисках с SDK).
2) Запускаем xbootmgr -trace <what> -traceFlags <flags> -resultPath <folder>.
3) После того, как профилируемый процесс завершится, в папке, указанной в resultPath, появится бинарный лог.
4) Специальной командой этот лог можно преобразовать в XML, который описывает высокоуровневые шаги процесса и их тайминг.
5) По результирующему XML можно проанализировать слабые места системы. В сабжевом посте достаточно детально объяснены этапы стартапа и шатдауна - может быть полезно для понимания происходящего.

Лично я, даже не вдаваясь в матан, смог разобраться почему так долго вырубается операционка. Дело было в том, что сервис TeamViewer выгружался больше пяти секунд. Винда даже успевала показать окошко "программа тормозит загрузку: выкосить или подождать еще?". Но, так как проблема была в сервисе, окошко было пустое, поэтому без профайлинга я бы ни в жизни не догадался.
glider

убунта, часть 2: установка

часть 1: первый взгляд
часть 2: установка
часть 3: софт
часть 4: интероп с виндой
часть 5: обратно на винду

Поставил убунту не только на десктоп, но и на ноут. Три недели - полет нормальный.

Самым главным удивлением при установке ОС было то, что нигде не понадобились драйвера - даже USB-tethering через телефон заработал из коробки сразу после втыкания шнурка. Это же касается и ноута (U31JG). Он вышел в начале этого года, но в убунте уже были под него дрова. Завелись даже комбинации клавиш Fn+XXX (кто настраивал сабжи для vaio под windows xp, тот поймет, почему я так заостряю внимание на этом малозначительном факте =)). Еще вот вспомнил прикол. До установки убунты я был абсолютно уверен, что у меня в ноуте нет блютуса. Уверенность эта базировалась на том, что ни один из блютус-драйверов с сайта асуса на ноут не поставился. А так как я брал самую мегабюджетную вариацию ноута, то подумал, что так и надо. Оказалось, что блютус есть =)

Впрочем, были и неприятные моменты. Во-первых, убунта не задетектила SSD (хотя это очень просто - сам потом написал скрипт) и TRIM пришлось включать самому. Еще непонятно, почему убунта по умолчанию не использует noatime. По-моему, atime (запись времени последнего обращения к файлам) это фича весьма сомнительной необходимости, но высокой деструктивности для SSD. Я погуглил интернет, и он вроде тоже разделяет мое мнение.

Во-вторых, где-то после 10.04 в убунте отломали vsync в связке x-server + opengl + nvidia drivers. В итоге получается вот так, особенно это вырубает при просмотре видео в моменты плавного горизонтального движения камеры. Применив все советы по ссылке выше, я каким-то чудом починил vsync, но какой из 100500 советов сработал проверять боюсь - вдруг, снова отвалится.

В-третьих, на ноуте из коробки не завелись suspend и hibernate. Первый опыт меня вообще поверг в ужас - по нажатию кнопки suspend экран ноута гаснет, но сам ноут продолжает тихо жужжать веником и не выключается. То же было и с hibernate. Оказалось, что это из-за того, что убунта сама не умеет отключать некоторые EHCI устройства, поэтому и показывает черный экран (ибо видеокарта уже уснула), но выключиться полностью не может (ибо некоторые устройства все еще работают). Фикс нагуглился весьма быстро: по сигналу засыпания нужно руками попросить проблемные устройства уснуть, а по пробуждению - включить их обратно (клик). "Руками" это в смысле скриптом, который запускается автоматически, а не то, что вы, возможно, подумали.
glider

Лайфхак по прокачке дропбокса

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

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

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

Пост был бы неполон без моего отчета о применении идеи на практике. Так как на все про все автор слинкованного поста затратил 400 рублей, я забросил на счет adwords 15$. Так же, как и советуется в посте, я создал рекламную кампанию на ключевое слово dropbox продолжительностью три дня. В качестве целевой аудитории выбрал Россию и русский язык (гугл предлагал еще показывать рекламу тем, у кого язык интерфейса английский, но я решил, что у таких челов уже давно есть дропбокс или что покруче).

И вот что случилось. Первые несколько часов вообще не было результатов, но за время вечерней прогулки по моей ссылке зарегались 8 человек. Пока я спал - еще 8. Пока завтракал - еще 2. Пока писал пост - еще 3. Плюс, 5 человек зарегистрировались в дропбоксе, но не поставили себе его клиента (т.е. за них мне дополнительное место еще не дают). На текущий момент, за 16 часов по моей ссылке кликнули 62 раза (из 332 показов). За все это гугл забиллил меня на 2.84$.
glider

Хибернейт

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

Сегодня от нечего делать решил проверить как оно будет работать. Заенаблил хибернейт через консоль (удивительно, но я не нашел как это сделать из GUI настроек винды), замаппил кнопку выключения системника на хибернейт и сделал "тыц" (памяти было занято около двух из восьми гектар). Результат превзошел мои ожидания - система выключилась за 15 секунд (обычный шатдаун занимает около 10 секунд), после чего включилась обратно за 3 (три!!) секунды. Цена удовольствия: 6 гектар зарезервированного места на SSD венике размером в 90 гектар.

Вроде бы на этом можно было бы закончить, но внимательный читатель может спросить: "а почему шесть гектар, если физической памяти восемь гектар?". Меня тоже заинтересовало это несоответствие, и я полез в гугл. Оказалось, что винда по умолчанию устанавливает размер hiberfil.sys в 75% от объема физической памяти, рассчитывая на то, что: 1) почти никогда вся память не занята, 2) сохраняемые страницы памяти сжимаются. Сам я не проверял, но гугл говорит, что в случае нехватки размера hiberfil.sys, винда выкинет синий экран с сообщением: "STOP 0x000000A0 INTERNAL_POWER_ERROR".

Последняя мелочь - в отличие от pagefile.sys, файл хибернейта перенести на другой диск нельзя. Я попробовал обойти это ограничение при помощи симлинка, но ничего не вышло: The system could not create the hibernation file. The specific error code is 0xc0000001.

Чит-шит:
* Включить хибернейт: PowerCfg.exe /HIBERNATE on
* Выключить хибернейт: PowerCfg.exe /HIBERNATE off
* Установить размер файла hiberfil.sys: PowerCfg.exe /HIBERNATE /SIZE <0-100> (дефолт: 75)

upd. По совету wizzard0 включил гибридный хибернейт - саспенд на 30 минут, после чего комп уходит в хибернейт и выключается насовсем. Радости от общения с компом стало гораздо больше. А потом и вовсе перешел на саспенд - лишней энергии потребляется мизер, а включение компа за 2-3 секунды в любое время суток на том же месте, где остановился - это просто праздник какой-то. Конечно, есть риск отключения питания и потери контекста, но за прошедший месяц такого ни разу не случилось => я пока что об этом не беспокоюсь.
glider

Conflux: GPGPU для .NET

К моему дню рождения в августе (серьезно, прямо в ночь перед ДР =)) случилось радостное событие - наконец-то удалось допилить Конфлакс до относительно рабочего состояния. А на недавней конференции Application Developer Days 2010 мне представился шанс выступить с докладом. Конечно же, я рассказывал про Конфлакс. В честь этого радостного события приглашаю вас ознакомиться со слайдами или даже посмотреть видео. Если вы никогда не слышали про CUDA, то перед этим лучше прочитать прекрасную презентацию Саймона Грина (для начального знакомства хватит первых 14 слайдов).

Тамбнейл доклада
Просто и ёмко Конфлакс описал antilamer: "в дотнетовской программе (на шарпе, например) можно написать метод, который будет выполняться на CUDA - он в рантайме декомпилируется и перекомпилируется в CUDA PTX". Если вас это описание заинтриговало, то вот как все это можно увидеть своими глазами:
  1) Скачать демку умножения матриц на Конфлаксе.
  2) Пока демка качается, можно посмотреть краткое описание.
  3) Демка состоит из двух частей - для CPU и для GPU. Для запуска демонстрации на CPU не нужно ничего, кроме четвертого дотнета. Для запуска демонстрации на GPU нужны CUDA-видеокарта и CUDA-драйвер.
  4) Обратите внимание на артефакты, которые создает демка. Во-первых, это логи, которые выводятся на консоль. Во-вторых, это динамически генерируемый код (ассемблер NVIDIA пишется в трейс, сгенерированный IL после завершения программы сбрасывается на диск - его можно посмотреть рефлектором). В-третьих, это CFG (графы потока управления) дизассемблированных методов, использованных в кернеле.
  5) Можно скачать исходники Конфлакса и посмотреть, как все работает. Для того, чтобы открыть проект, нужна Visual Studio 2010.

В заключение несколько слов саммари. Конфлакс позволяет программисту в 100% managed коде описать параллельный алгоритм, который потом будет исполняться на GPU. Написанный алгоритм можно прозрачно дебагать, используя стандартные средства Visual Studio, одной строчкой переключившись на CPU бакэнд. Наконец, несмотря на происходящую под капотом магию, Конфлакс не является дополнительным уровнем абстракции, которая может неожиданно протечь - он всего лишь один-к-одному преобразует код на дотнете в код для GPU, в нем можно даже писать ассемблерные вставки и вызывать third-party библиотеки CUDA (например, CUBLAS или CUFFT). Все это является шагом вперед в мире GPGPU для дотнета.
glider

Полиэдральная модель оптимизации циклов

Начав разбираться с GPGPU, я почти сразу узнал, что, несмотря на высокий пиковый перфоманс современных видеокарт (еще прошлое поколение GPU почти достигло single-precision терафлопа), программер должен весьма потрудиться для того, чтобы достигнуть хотя бы 30-50% теоретической производительности. Этот труд заключается в том, что после написания и отладки код, выражающий вычислительный алгоритм, уродуется до неузнаваемости и неподдерживаемости ручными оптимизациями, которые учитывают все штучки-огонёчки подлежащей хардварной платформы. Вот один из самых безобидных примеров таких оптимизаций: loop unrolling.

Естественный ужос толкнул меня на поиск математических моделей, которые позволяют все такие штуки производить автоматически на этапе компиляции, приближаясь хотя бы к 80-90% перфоманса хардкорного хэндкрафтинга. Конечно, в общем случае такая задача неразрешима и state of the art современной параллелизации склоняется к недетерминированным рантайм оптимизациями через спекуляцию и профилирование.

Впрочем, есть и хорошие новости. Для общего случая задачу решать совсем необязательно, ибо есть несколько удачных особенностей проблемной области. Во-первых, control flow программ для GPU в большинстве случаев статический, а, во-вторых, в значительном проценте вычислительных алгоритмов доступ к памяти происходит весьма регулярно (например, нет динамического диспатча). Эти ограничения на анализируемый алгоритм делают возможным использование полиэдральной модели (polyhedral model).

***

Полиэдральная модель предназначена для анализа и трансформации гнезд циклов (loop nests, т.е. набора возможно вложенных циклов), которые удовлетворяют ограничениям: 1) скалярности (данные, над которыми ведется работа, либо скалярны, либо являются массивами скаляров) , 2) линейности (т.е. границы циклов и индексы массивов либо константны, либо представляют собой линейную комбинацию счетчиков внешних циклов и констант) и 3) статичности (т.е. условия выхода из цикла статичны, внутри циклов нет ветвления). Удивительно, но очень много вычислительных алгоритмов (например, все дифуры, решаемые конечными разностями) сюда подходят.

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

Ниже приведен пример представления несложного гнезда циклов в полиэдральной модели (эта картинка и картинки ниже взяты из работ исследователей Стэнфордского университета - Майкла Вольфа и Моники Лэм: "A Loop Transformation Theory and an Algorithm to Maximize Parallelism" и "A Data Locality Optimizing Algorithm"; сорри за качество картинок, но работам почти 20 лет, отсканированы они хреново, а лучших примеров для введения в полиэдральную модель я не нашел). Несколько пояснений: 1) в примере границы циклов константны, но это не является обязательным условием для представления циклов в полиэдральной модели (см. дальнейшие примеры), 2) в примере рассматривается совершенное гнездо циклов (perfect loop nest), т.е. такое, что все операции производятся только в самом внутреннем цикле гнезда, но и это условие не является обязательным для полиэдральной модели, 3) зависимости между точками пространства итераций обозначаются стрелками (еще раз сорри за качество картинки).

Пример представления цикла в полиэдральной модели

***

Итак, в полиэдральной модели задача оптимизации гнезда циклов заключается в нахождении такого обхода пространства итераций, который сохраняет зависимости по данным (т.е. является "легальным") и максимизирует некоторую целевую функцию (например, степень параллелизма или процент кэш-хита). Это будет задача целочисленного линейного программирования (parametric integer programming, PIP). В общем случае PIP является NP-полной, поэтому надо применять какие-то эвристики по отсечению пространства перебора.

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

А сейчас пример. Вот два способа, которыми можно оптимизировать гнездо циклов на картинке выше. Некоторые пояснения: 1) в ручном труде нет необходимости - ужасные границы циклов и индексы массивов генерятся автоматически из матриц преобразования координат, 2) на первой картинке итерации внутреннего цикла работают с данными, лежащими недалеко друг от друга, что позитивно сказывается на перфомансе памяти, 3) на второй картинке мы, несмотря на неслабые зависимости по данным, все же нашли способ выделить параллелизм, 4) разница между первой и второй картинками всего лишь в матрице трансформации.

Пример вымощения пространства итераций для достижения хорошей локальности данных

Пример вымощения пространства итераций для достижения высокой степени параллелизма

То, что описано выше, было state of the art лет 15 назад, но сейчас разработаны гораздо более интересные штуки, которые позволяют достичь в среднем более хороших результатов оптимизации, например, и параллелизм выщемить, и локальность не упустить.

***

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

Применению фреймворка в контексте GPGPU посвящено несколько работ, из которых больше всего мне нравится "A Compiler Framework for Optimization of Affine Loop Nests for GPGPUs", которая, несмотря на свой полуторагодовалый возраст, покрывает подавляющее большинство современных техник автоматической оптимизации. В этой работе составлены несколько вариантов целевых функций, учитывающих особенности современных GPU, приведены эвристики по их комбинации и отсечению пространства решений.

Впрочем, несмотря на неплохие продвижения в применении полиэдральной модели для GPGPU, остается куча непонятных вопросов: 1) как производить более качественное отсечение пространства решений, не теряя хороших решений, но сохраняя полиномиальное время? 2) как моделировать локальные переменные? 3) как моделировать структуры? Есть надежда, что на эти вопросы удастся найти хорошие ответы, ибо уж очень велик соблазн уменьшить ручной труд, необходимый для оптимизации программ для GPU.
glider

Гайд по CUDA

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

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

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).
glider

Хардварная реализация 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-локальности кэша констант).
glider

Программная модель 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 сценарии. Впрочем, эти факты в большинстве случаев облегчаются тем, что, благодаря большому количеству выполняемых потоков, высокую латентность можно амортизировать - значительным образом или даже полностью.