Note that cpu benchmarks are not measured via opencl что это
Note that cpu benchmarks are not measured via opencl что это
Вычисляем на видеокартах. Технология OpenCL. Часть 1a. Как работает OpenCL
22 июня автор курса «Разработчик C++» в Яндекс.Практикуме Георгий Осипов провёл вебинар «Вычисляем на видеокартах. Технология OpenCL».
Мы подготовили для вас его текстовую версию, для удобства разбив её на смысловые блоки.
0. Зачем мы здесь собрались. Краткая история GPGPU.
1a. Как работает OpenCL.
1b. Пишем для OpenCL.
2. Алгоритмы в условиях массового параллелизма.
3. Сравнение технологий.
Мы обещали, что разберём написание полноценной программы уже в этой части, но материала оказалось слишком много, и мы разбили эту часть надвое. В первой половине расскажем про основные принципы, которые должен знать каждый OpenCL-разработчик, а во второй напишем программу.
Есть мнение, что для написания эффективного кода для GPU программист обязан понимать архитектуру видеокарты. И это мнение не чьё-нибудь там, а NVIDIA (см. Лекции NVIDIA по GPGPU). Не будем спорить и разберём базовые принципы работы видеокарты.
Kernel и Host
В предыдущей серии шла речь о том, почему не стоит избавляться от CPU. Вот ещё одна причина: центральный процессор необходим для работы видеокарты, ведь именно он отдаёт команды.
В контексте GPGPU центральный процессор называют английским словом host — хозяин. Видеокарта у него как бы в гостях. CPU говорит видеокарте, что именно надо посчитать — готовит задачу и ставит её в очередь. Видеокарта берёт и считает. В комментариях к предыдущей статье отметили, что уже есть технологии, позволяющие обходиться и без CPU, но мы будем рассматривать только те, которые используются широко.
Нужно понимать, что на хосте работает одна программа, а на GPU — другая. Последняя называется kernel. Итак, host — это CPU, kernel — видеокарта. Host даёт команды, kernel считает. Эти слова я буду использовать дальше, и важно их запомнить.
С GPU одновременно могут работать сколько угодно обычных программ, исполняемых CPU. При возникновении задачи она добавляется в очередь и запускается на общих основаниях. Видеокарта может комбинировать задачи и выполнять их в произвольном порядке, чтобы обеспечить наибольшую эффективность, но при этом не нарушить логику работы каждого приложения. Ей неважно, какая программа добавила задачу в очередь.
Теперь о том, как устроены программы, использующие OpenCL. Вы пишете host-код — обычную программу под свою ОС. Кстати, host-код, работающий с OpenCL, можно писать практически на любом языке программирования. Внутри host-программы есть kernel-код в виде обычной строки-ресурса. Он написан на языке C, немножко модифицированном. Программа присылается пользователю. При запуске она ищет динамическую библиотеку OpenCL. И когда программа её находит, то обращается к библиотеке с просьбой скомпилировать кусок kernel-кода. При этом указывается устройство, под которое нужно скомпилировать код. Библиотека OpenCL для этого идёт к драйверу видеокарты, имеющему компилятор. Получается, что библиотека OpenCL — эдакий посредник между вашей программой и драйверами всех устройств, установленных на компьютере, которые могут исполнять OpenCL.
Компиляция происходит непосредственно перед запуском: нельзя заранее скомпилировать kernel-программу и отправить пользователю бинарный код. Потому что вы не знаете, на какой архитектуре будет выполняться ваш kernel. Host-код вы компилируете как обычно, а kernel поставляется исходным кодом и компилируется уже у пользователя.
Геометрия задачи и рабочие группы
Напомню, что OpenCL предназначен для задач массового параллелизма. Это выражается в многократном и синхронном выполнении одного и того же kernel-кода. Одно выполнение kernel называется work item. Для каждой задачи задаётся размерность — она определяет, в виде какой геометрической фигуры удобно представлять себе work item’ы. Иногда удобно располагать их в линию, тогда размерность равна одному. Например, суммирование массива чисел — одномерная задача в work item. Если нужно что-то делать с изображением, тогда один work item — это обработка одного пикселя изображения. Их удобнее расположить в таблице — размерность равна двум. Или задача может быть трёхмерной, если нужно обработать воксели — трёхмерные пиксели.
Work item’ы объединяются в рабочие группы, workgroup. Это важное понятие, смысл которого станет понятен после описания понятия локальной памяти. Все рабочие группы имеют одинаковые размеры. В случае двумерной задачи они будут прямоугольниками, а в случае трёхмерной — параллелепипедами.
При запуске kernel-кода мы задаём следующие параметры.
Рабочие группы делятся на другие группы, которые называются ворпами (warp) или вейвфронтами (wavefront), в зависимости от производителя видеокарты. Я буду использовать термин «ворп». Особенность этого деления в том, что группировка в ворпы происходит автоматически, за кулисами API. Вы не указываете размер ворпа — он зависит только от архитектуры видеокарты.
Допустим, мы майним биткоины. Эта задача одномерная, для неё нет смысла вводить структуру таблицы. Каждый work item — одна элементарная задача, которую нужно вычислить. В данном случае — вычисление одного хеша. Если нужно преобразовать изображение, то work item лучше расположить в таблице. Это двумерная задача, и один work item в ней — вычисление одного пикселя.
Для лучшего понимания работы work item можно представить, что рабочая группа — это рой маленьких роботов. Их много, и каждый робот, work item, исполняет kernel. У него есть данные, одинаковые для всех роботов:
У роботов есть доступ к глобальной памяти, которая аналогична обычной RAM, и к локальной, отдельной для каждой рабочей группы. К локальной памяти могут обращаться все work item’ы внутри рабочей группы, но роботы из другой рабочей группы видят другую локальную память.
GPU выполняет рабочие группы как попало, в произвольном порядке. Поэтому нет никаких средств синхронизации между рабочими группами. Смысл рабочих групп — в локальной памяти, которая к тому же очень быстрая. Внутри одной рабочей группы роботы могут синхронизироваться между собой. Например, если они уже выполнили подзадачу, то могут подождать, пока это сделают другие, чтобы вместе приступить к следующей.
Можно представить, что каждая рабочая группа — это отдельный рой, и в нём все роботы связаны между собой через локальную память. Но задачу решают несколько рабочих групп, несколько роёв. И делают это независимо. Как правило, задача kernel сводится к чтению входных данных из глобальной памяти и записи ответа в глобальную же память.
«Синхронное плавание» и какие проблемы оно влечёт
Другая аллегория, которая подходит для описания работы ворпа, — синхронное плавание. В нём спортсменки одновременно выполняют одно и то же действие. Так же устроен и ворп. Он объединяет 32 или 64 work item’а, которые работают синхронно, и производят в один момент одно и то же вычисление, но над разными числами. Нужно понимать: выполняемая инструкция одна и та же, но данные регистров у каждого ворпа свои. Иначе смысла в таком вычислении не было бы.
В простых случаях дивергенция не страшна, и не стоит её бояться. Но в более сложных программах она может серьёзно ухудшить производительность. Рассмотрим пример — бинарный поиск.
Из-за дивергенции код бинарного поиска, изображённого на слайде, перестаёт быть логарифмическим и становится линейным.
Разберёмся, как бороться с такими нежелательными явлениями. Совет первый: не использовать рекурсию. С одной стороны, совет бесполезный: в OpenCL рекурсия запрещена. С другой стороны, она может эмулироваться. Тот же бинарный поиск можно реализовать циклом без указанной проблемы.
Совет второй: оптимизируйте код с учётом дивергенции. Если подумать, код можно переписать очень просто: нужно предвычислить интервал, а потом один раз рекурсивно вызвать binsearch. Тогда дивергенция утратит могущество и перестанет быть значимой. Пример показан на слайде.
Дивергенция относится не только к if-else, но также к циклам и switch-case, return, в общем, ко всему, что вызывает ветвление кода. Если в цикле все потоки ворпа, кроме одного, уже завершили свои итерации, то они вместе с отстающим будут продолжать ходить по кругу.
Как выбирать локальные и глобальные размеры
Если размер рабочей группы не делится на размер ворпа, ничего страшного не произойдёт. Но код будет немного неэффективным: видеокарта дополнит рабочую группу до кратного размера «холостыми» потоками.
Чтобы выбрать размер группы, помимо делимости на размер warp нужно учитывать объём локальной памяти. Чем больше рабочая группа, тем больше понадобится памяти, а её размер ограничен. Также у устройств есть жёсткие ограничения на размер рабочей группы, который можно запросить у драйвера.
Часто подходит жёстко зафиксированное значение — 256. Его поддерживают все устройства. Но скажем честно: размер группы не так уж влияет на производительность. Видеокарты умные, они перераспределят нагрузку, даже если вы выберете не самый оптимальный размер. Конечно, совсем маленькие группы делать не стоит: большое их количество влечёт дополнительные издержки.
Следующий вопрос: как распределить выбранный размер рабочей группы? Рассмотрим на примере двумерной задачи. Тут есть минимум три варианта:
Как правило, оптимален третий вариант, горизонтальные полоски. Если задача обращается к глобальной памяти, то при использовании горизонтальных полос будут браться последовательные данные, и получится по максимуму использовать кеш. При обращении к текстурной памяти, которая организована в виде z-кривой, или для задач с непоследовательным доступом к памяти этот принцип не работает. При выборе геометрии рабочей группы нужно исходить из специфики задачи. Использование неоптимальной геометрии может замедлить работу программы в десятки раз.
Как мы увидим из следующих частей, в некоторых задачах желательно сократить количество рабочих групп. Следовательно, размер каждой группы должен быть как можно больше. Если жёстко заданный размер вас не устраивает, то нужно действовать следующим алгоритмом:
На этом подготовка почти закончена. В следующей части статьи приступим к написанию программы.
OpenCL. Подробности технологии
Здравствуй, уважаемое хабрасообщество.
В предыдущей статье про OpenCL был сделан обзор этой технологии, возможностей, которые она может предложить пользователю и ее состояния на настоящий момент.
Теперь рассмотрим технологию более пристально. Постараемся понять, как OpenCL представляет гетерогенную систему, какие предоставляет возможности по взаимодействию с устройством и какой предлагает подход к созданию программ.
OpenCL задумывался как технология для создания приложений, которые могли бы исполняться в гетерогенной среде. Более того, он разработан так, чтобы обеспечивать комфортную работу с такими устройствами, которые сейчас находятся только в планах и даже с теми, которые еще никто не придумал. Для координации работы всех этих устройств гетерогенной системе всегда есть одно «главное» устройство, который взаимодействует со всеми остальным посредствами OpenCL API. Такое устройство называется «хост», он определяется вне OpenCL.
Поэтому OpenCL исходит из наиболее общих предпосылок, дающих представление об устройстве с поддержкой OpenCL: так как это устройство предполагается использовать для вычислений – в нем есть некий «процессор» в общем смысле этого слова. Нечто, что может исполнять команды. Так как OpenCL создан для параллельных вычислений, то такой процессор может, иметь средства параллелизма внутри себя (например, несколько ядер одного CPU, несколько SPE процессоров в Cell). Также элементарным способом наращивания производительности параллельных вычислений является установка нескольких таких процессоров на устройстве (к примеру, многопроцессорные материнские платы PC итд.). И естественно в гетерогенной системе может быть несколько таких OpenCL-устройств (вообще говоря, с различной архитектурой).
Кроме вычислительных ресурсов устройство имеет какой-то объем памяти. Причем никаких требований к этой памяти не предъявляется, она может быть как на устройстве, так и вообще быть размечена на ОЗУ хоста (как например, это сделано у встроенных видеокарт).
Собственно все. Больше об устройстве никаких предположений не делается.
Такое широкое понятие об устройстве позволяет не накладывать каких-либо ограничений на программы, разработанные для OpenCL. Эта технология позволит Вам разрабатывать как приложения, сильно оптимизированные под конкретную архитектуру специфического устройства, поддерживающего OpenCL, так и те, которые будут демонстрировать стабильную производительность на всех типах устройств (при условии эквивалентной производительности этих устройств).
OpenCL предоставляет программисту низкоуровневый API, через который он взаимодействует с ресурсами устройства. OpenCL API может либо напрямую поддерживаться устройством, либо работать через промежуточный API (как в случае NVidia: OpenCL работает поверх CUDA Driver API, поддерживаемый устройствами), это зависит от конкретной реализации не описывается стандартом.
Рассмотрим как же OpenCL обеспечивает такую универсальность, сохраняя при этом низкоуровневую природу.
Далее я приведу вольный перевод части спецификации OpenCL 1.0 с некоторыми комментариями и дополнениями.
Модель платформы (Platform Model).
Платформа OpenCL состоит из хоста соединенного с устройствами, поддерживающими OpenCL. Каждое OpenCL-устройство состоит из вычислительных блоков (Compute Unit), которые далее разделяются на один или более элементы-обработчики (Processing Elements, далее PE).
OpenCL-приложение исполняется на хосте в соответствии с нативными моделями его платформы. OpenCL-приложение отправляет с хоста команды устройствам на выполнение вычислений на PE. PE в рамках вычислительного блока выполняют один поток команд как SIMD блоки (одна инструкция выполняется всеми одновременно, обработка следующей инструкции не начнется, пока все PE не завершат исполнение текущей инструкции), либо как SPMD блоки (у каждого PE собственный счетчик инструкций (program counter)).
То есть OpenCL обрабатывает некие команды, поступающие от хоста. Таким образом приложение не связано жестко с OpenCL, а значит всегда можно подменить реализацию OpenCL, не нарушив работоспособность программы. Даже если будет создано такое устройство, которое не укладывается в модель «OpenCL-устройства», для него можно будет создать реализацию OpenCL, транслирующую команды хоста в более удобный для устройства вид.
Модель исполнения (Execution Model).
Выполение OpenCL-программы состоит из двух частей: хостовая часть программы и kernels (ядра; с Вашего позволения я далее буду употреблять английский термин, как более привычный большинству из нас) исполняющиеся на OpenCL-устройстве. Хостовая часть программы определяет контекст, в котором исполняются kernel’ы, и управляет их исполнением.
Основная часть модели исполнения OpenCL описывает исполнение kernel’ов. Когда kernel ставится в очередь на исполнение, определяется пространство индексов (NDRange, определение будет дано ниже). Копия (instanse) kernel’а выполнятся для каждого индекса из этого пространства. Копия kernel’а выполняющаяся для конкретного индекса называется «Work-Item» (рабочей единицей) и определяется точкой в пространстве индексов, то есть каждой «единице» предоставляется глобальный ID. Каждый Work-Item выполняет один и тот же код, но конкретный путь исполнения (ветвления итп.) и данные, с которыми он работает, могут быть различными.
Work-Item’ы организуются в группы (Work-Groups). Группы предоставляют более крупное разбиение в пространстве индексов. Каждой группе приписывается групповой ID с такой же размерностью, которая использовалась для адресации отдельных элементов. Каждому элементу сопоставляется уникальный, в рамках группы, локальный ID. Таким образом, Work-Item’ы могут быть адресованы как по глобальному ID, так и по комбинации группового и локального ID.
Work-Item’ы в группе исполняются конкурентно (параллельно) на PE одного вычислительного блока.
Выбор размерности NDRange определяется удобством для конкретного алгоритма: в случае работы с трехмерными моделями удобно индексировать по трехмерным координатам, в случае работы с изображениями или двумерными сетками – удобнее, когда размерность индексов – 2. 4х-мерные объекты в нашем мире большая редкость, поэтому размерность ограничена 3. Кроме того, как бы там ни было, но в данный момент основная цель OpenCL – это GPU. GPU Nvidia сейчас нативно поддерживают размерность индексов до 3, соответственно, чтобы реализовать большую размерность, пришлось бы прибегать к хитростям и усложнению либо CUDA Driver API, либо реализации OpenCL.
Контекст исполнения и очереди команд в модели исполнения OpenCL.
Использование очереди команд, позволяет добиться большой универсальности и гибкости при использовании OpenCL. Современные GPU имеют собственный планировщик, который решает, что и когда и на каких вычислительных блоках исполнять. Использование очереди не стесняет работу планировщика, который имеет собственную очередь команд.
Модель исполнения: категории kernel.
Модель памяти (Memory Model).
Спецификация определяет 4 типа памяти, но снова не накладывает никаких требований на реализацию памяти в железе. Все 4 типа памяти могут находиться в глобальной памяти, и разделение типов может осуществляться на уровне драйвера и напротив, может существовать жесткое разделение типов памяти, продиктованное архитектурой устройства.
Существование именно этих типов памяти достаточно логично: у процессорного ядра есть свой кэш, у процессора есть общий кэш и у всего устройства есть некоторый объем памяти.
Программная модель. (Programming Model)
Модель исполнения OpenCL поддерживает две программные модели: параллелизм данных (Data Parallel) и параллелизм заданий (Task Parallel), так же поддерживаются гибридные модели. Основная модель, определяющая дизайн OpenCL, – параллелизм данных.
Программная модель с параллелизмом данных.
Эта модель определяет вычисления как последовательность инструкций, применяемых к множеству элементов объекта памяти. Пространство индексов, ассоциированное с моделью исполнения OpenCL, определяет Work-Item’ы и как данные распределяются между ними. В строгой модели параллелизма данных существует строгое соответствие один к одному между Work-Item и элементом в объекте памяти, с которым kernel может работать параллельно. OpenCL реализует более мягкую модель параллелизма данных, где строгое соответствие один к одному не требуется.
OpenCL предоставляет иерархическую модель параллелизма данных. Существует два способа определить иерархическое деление. В явной модели программист определяет общее число элементов, которые должны исполняться параллельно и так же каким образом эти элементы будут распределены по группам. В неявной модели программист только определяет общее число элементов, которые должны исполняться параллельно, а разделение по рабочим группам выполняется автоматически.
Программная модель с параллелизмом заданий.
Существование двух моделей программирования – также дань универсальности. Для современных GPU и Cell хорошо подходит первая модель. Но не все алгоритмы можно эффективно реализовать в рамках такой модели, а так же есть вероятность появления устройства, архитектура которого будет неудобна для использования первой модели. В таком случае вторая модель позволяет писать специфичные для другой архитектуры приложения.
Из чего состоит платформа OpenCL
Как это все работает?
В следующей статье я подробно разберу процесс создания приложения OpenCL на примере одного из приложений, распространяемых вместе с Nvidia Computing SDK. Приведу примеры оптимизаций работы приложений для OpenCL, предлагаемые Nvidia в качестве рекомендаций.
Стоит отметить что сборка программы осуществляется во время исполнения, практически JIT-комиляция. В стандарте описано, что это сделано для того, чтобы можно было собрать программу с учетом выбранного контекста. Так же это позволяет каждому поставщику реализации OpenCL оптимизировать компилятор под свое устройство. Впрочем, программу можно также создавать из бинарных кодов. Либо создавать ее один раз при первом запуске, а в дальнейшем переиспользовать, такая возможность тоже описана в стандарте. Тем не менее компилятор интегрирован в платформу OpenCL, хорошо это или плохо, но это так.
Заключение.
В итоге модель OpenCL получилась весьма универсальной, при этом она остается низкоуровневой, позволяя оптимизировать приложения под конкретную архитектуру. Так же она обеспечивает кроссплатформенность при переходе от одного типа OpenCL-устройств к другому. Поставщик реализации OpenCL имеет возможность всячески оптимизировать взаимодействие своего устройства с OpenCL API, добиваясь повышения эффективности распределения ресурсов устройства. Кроме того, правильно написанное OpenCL приложение будет оставаться эффективным при смене поколений устройств.
Сравнение OpenCL с CUDA, GLSL и OpenMP
На хабре уже рассказали о том, что такое OpenCL и для чего он нужен, но этот стандарт сравнительно новый, поэтому интересно как соотносится производительность программ на нём с другими решениями.
В этом топике приведено сравнение OpenCL с CUDA и шейдерами для GPU, а также с OpenMP для CPU.
Тестирование проводилось на задаче N-тел. Она хорошо ложится на параллельную архитектуру, сложность задачи растёт как O(N 2 ), где N — число тел.
Задача
В качестве тестовой была выбрана задача симуляции эволюции системы частиц.
На скриншотах (они кликабельны) видна задача N точечных зарядов в статическом магнитном поле. По вычислительной сложности она ничем не отличается от классической задачи N тел (разве что картинки не такие красивые).
Во время проведения замеров вывод на экран был отключен, а FPS означает число итераций в секунду (каждая итерация — это следующий шаг в эволюции системы).
Результаты
Код на GLSL и CUDA для этой задачи был уже написан сотрудниками ННГУ.
NVidia Quadro FX5600
Версия драйвера 197.45 
CUDA обгоняет OpenCL приблизительно на 13%. При этом, если оценивать теоретически возможную производительность для этой задачи для данной архитектуры, реализация на CUDA достигает её.
(В работе A Performance Comparison of CUDA and OpenCL говорится о том, что производительность ядра OpenCL проигрывает CUDA от 13% до 63% )
Несмотря на то, что тесты проводились на карточке серии Quadro, понятно, что обычный GeForce 8800 GTS или GeForce 250 GTS дадут схожие результаты (все три карточки основаны на чипе G92).
Radeon HD4890
ATI Stream SDK версия 2.01 
OpenCL проигрывает шейдерам на карточках от AMD так как вычислительный блоки на них имеют архитектуру VLIW, на которую (после оптимизации) могут хорошо лечь многие шейдерные программы, но компилятор для кода OpenCL (который является частью драйвера) плохо справляется с оптимизацией.
Также этот весьма скромный результат может быть вызван тем, что карточки от AMD не поддерживают локальную память на физическом уровне, а отображают область локальной памяти на глобальную.
Код с использованием OpenMP был скомпилирован при помощи компиляторов от Intel и Microsoft.
Компания Intel не выпустила своих драйверов для запуска кода OpenCL на центральном процессоре, поэтому был использован ATI Stream SDK.
Intel Core2Duo E8200
ATI Stream SDK версия 2.01 
Код на OpenMP, скомпилированный при помощи MS VC++ имеет практически идентичную производительность с OpenCL.
Это ещё при том, что Intel не выпустил своего драйвера для интерпретации OpenCL, и используется драйвер от AMD.
Компилятор от Intel поступил не совсем «честно» он полностью развернул основной цикл программы, повторив его где-то 8k раз (число частиц было задано константой в коде) и получив семикратный прирост производительности также благодаря использованию SSE инструкций. Но победителей, конечно, не судят.
Что характерно, на моём стареньком AMD Athlon 3800+ код тоже запустился, но таких выдающихся результатов, как на Intel, конечно, ждать не приходится.
Заключение
Введение в OpenCL
Компилятор
Модель памяти устройства
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «work-item». Каждый «work-item» не зависим от другого и может исполнять код параллельно с остальными. Если же процесс из одного work-item хочет получить данные, используемые или уже обработанные любым другим work-item он может это сделать через общую память. Общая память весьма медленная, зато имеет большой объём. Чтобы ускорить вычисления имеется локальная память. Если вы знакомы с CUDA, то там она называется «разделяемая память». Она значительно быстрее общей, но не любой процесс может получить к ней доступ. К локальной памяти могут обращаться только work-item одной группы. Эти группы называются «Compute Unit» или «Workgroup» (первое название относится к физическому разбиению на уровне железа, а второе к логическому на уровне программы). В зависимости от устройства в каждой из этих групп различное количество work-item (например 240 для NVIDIA GT200 или 256 для Radeon 5700 Series). Количество этих юнитов ограниченно достаточно маленьким числом (30 для NVIDIA GT200 или 9-10 для Radeon 5700 Series). Так же существует сверхбыстрая «private memory» к которой work-item может обращаться единолично.
Драйвера OpenCL устройств автоматизируют старт и работу work-item и workgroup. Например если нам нужно выполнить миллион процессов, а у нас в распоряжении всего тысяча work-item, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.
Базовые особенности
Объявление процедур
В первую очередь в глаза бросается загадочный «__kernel «. Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.
Типы памяти
Тип данных «__global » обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть «__local «. К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является «__privat «. Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе «__global «. Во-первых, это «__constant «, который может использоваться только для чтения. Во-вторых, это «__read_only», «__write_only» и «__read_write» — структуры, использование которых разрешено только для изображений.
Идентификаторы процессов
Оптимизация расчётов
Разработчики OpenCL и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Встроенные вектора
Простые функции
Следующей особенностью OpenCl является встроенная библиотека функций. Кроме стандартного набора math.lib в OpenCl имеются так называемые native функции. Это функции, основаны непосредственно на использовании некоторых функций видеокарт и на загрублённой математике. Не советуется применять их при сверхточных расчётах, но в случае фильтрации изображений разницу невозможно заметить. К таким функциям, например, относятся: «native_sin», «native_cos», «native_powr». Я не буду приводить более подробное объяснение этих функций, их очень много, да и принципы разные. Если они вам понадобятся — смотрите документацию.
Часто встречающиеся функции
Кроме «простых функций» разработчики создали целый ряд называемый common function. Это функции, часто встречающиеся при обработке изображений. Например: mad(a,b,c) = a*b + c, mix(a,b,c) = a + (b-a)*c. Эти функции выполняются быстрее, чем соответствующие им математические действия.
Пример
Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.
Разрешения
#pragma OPENCL EXTENSION extension name : behavior
Синхронизация
Барьеры
В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды «__local int x[10];» и «__global int x[10];». Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Единичные операции
void GetSemaphor(__global int * semaphor) <
int occupied = atom_xchg(semaphor, 1);
while (occupied > 0)
<
occupied = atom_xchg(semaphor, 1);
>
>
Работа с изображениями
Последней вещью, которую я хочу включить в этот guide является работа с изображениями через OpenCL. Создатели попробовали сделать так, чтобы работа с изображениями требовала минимума мозга пользователя. Это очень приятно. Загрузка изображений возможна в типы image2d_t и image3d_t. Первые — это обычные изображения, вторые — трёхмерные. Так же загружаемое изображение должно быть одного из форматов: » __read_only», » __write_only», «__read_write». Чтение и запись данных из изображения возможны только специальными процедурами: значение = read_imageui(изображение, сэмплер, положение), write_imageui(изображение, положение, значение).
На мой взгляд здесь всё понятно кроме понятия «сэмплер». Сэмплер — это штука, которая будет оптимизировать вашу работу с изображением. У него есть три параметра: «normalized coords», «address mode», «filter mode». Первый имеет два значения: «CLK_NORMALIZED_COORDS_TRUE, CLK_NORMALIZED_COORDS_FALSE». В соответствии с названием он должен показывать, нормализованы ли входные координаты или нет. Второй показывает, что делать в случае, если вы пробуете прочитать координаты из-за пределов границ изображения. Возможные варианты: зеркально продолжить изображение(CLK_ADDRESS_MIRRORED_REPEAT), взять ближайшее граничное значение (CLK_ADDRESS_CLAMP_TO_EDGE), взять базовый цвет (CLK_ADDRESS_CLAMP), ничего не делать (пользователь гарантирует что такого не произойдёт CLK_ADDRESS_NONE). Третий показывает, что делать, если на входе не целые координаты. Возможные варианты: приблизить ближайшим значением (CLK_FILTER_NEAREST), линейно проинтерполировать (CLK_FILTER_LINEAR).
Краткий пример. Замыливаем изображение по среднему значению в области:
__kernel void ImageDiff(__read_only image2d_t bmp1, __write_only image2d_t bmpOut)
Первые шаги с OpenCL или сказ о том как одинаковый код на GPU и CPU запускать
Думаю, что пересказывать Википедию об OpenCL особого смысла нет, но если в двух словах, то OpenCL — это язык, (фреймворк и платформа), который позволяет запускать один и тот же код на разных устройствах с разными архитектурами, а в особенности на высокопараллельных процессорах, вроде видеокарт и современных центральных процессоров. Основан стандарт на C99 и поддерживается The Khronos Group, на этом ликбез будем считать завершенным.
Начну я с того, что покажу небольшой кусочек кода и буду объяснять, что там происходит, параллельно рассказывая о том, как OpenCL работает.
Сначала я опишу достаточно тривиальный код и те, кому совсем не терпится увидеть магиюOpenCL, могут пропустить первую часть (только прочитайте последний абзац, где я описываю функцию MathCalculations, это важно. А если вы знаете об OpenCL и вам хочется увидеть результаты тестов, то идите сразу в пятый раздел, но все равно загляните в MathCalculations).
Вот так выглядит main моей небольшой программки для тестирования OpenCL, а если точнее, то для расчета некоего абстрактного математического выражения, до которого мы попозже дойдем. Итак, давайте построчно разбираться, что же тут происходит.
Часть первая — Инициализация исходных данных и традиционный способ вычислений
В ней первый цикл
нужен для того, чтобы провести тест несколько раз для получения более точного времени выполнения. Время вычисления каждого теста сохраняются в массиве timeValues из которого потом вычисляется среднее значение и сохраняется в hostPerformanceTimeMS.
последовательно производит некие математические вычисления над элементами входных массивов и сохраняет их в выходном массиве.
Как мы видим, в этом коде нет ничего необычного, он компилируется обычными сишным компилятором и выполняется последовательно на центральном процессоре, как и большая часть кода, который мы все пишем каждый день. А нужен он нам для того, чтобы впоследствии сверить с ним результаты, полученные OpenCL, а также понять, что за прирост производительности мы получаем.
Тут же стоит заглянуть в MathCalculations и увидеть, что там все совсем скучно:
Часть вторая — Инициализация OpenCL
Итак, терпеливые дочитали до этой части и обрадовались, что начинается интересное, а нетерпеливые этого чувства испытать не смогут, они прошлый абзац пропустили:)
Сначала я скажу о том, что OpenCL Runtime API представляет из себя именно API для C, а не для C++. В целом, в этом нет ничего плохого кроме того, что для проверки ошибок надо проверять код, возвращаемый каждой функцией и это не очень удобно. А также надо вручную следить за освобождением выделенных ресурсов.
Но есть также и официальная C++ обертка (ее можно найти на сайте Khronos), которая представляет из себя набор классов, соответствующих объектам OpenCL и поддерживающим подсчеты ссылок (reference counting который) и бросание исключений в случае ошибок (исключения надо включать при помощи #define __CL_ENABLE_EXCEPTIONS). Вот эту самую обертку я и буду использовать в нашем тесте.
Итак первым делом мы получаем список доступных платформ:
Платформа в OpenCL соответствует вендору, т.е. у NVidia будет одна платформа с ее устройствами, у Intel другая итд итп. В моем случае мне доступны как раз две платформы NVidia и Intel.
Сразу еще один маленький трюк, C++ wrapper может пользоваться своими собственными векторами (если ему об этом сказать) или векторами из STD, так что если где-то в примерах попадется что-то вроде cl::vector, не пугайтесь, он знает оба формата.
После того как мы получили список платформ, для каждой платформы мы получаем список доступных устройств:
Собственно устройства — это то, что будет выполнять наши вычисления. Это может быть и GPU, и CPU и какой-то специальный ускоритель, который подключен к хосту, т.е. той системе, на которой запускается OpenCL. Вместо CL_DEVICE_TYPE_ALL можно передать CL_DEVICE_TYPE_GPU, тогда он будет выдавать только видеокарты или CL_DEVICE_TYPE_CPU для центральных процессоров.
Для каждого найденного устройства я запускаю тест, о котором расскажу чуть ниже, и пытаюсь отловить исключения, которые бросит OpenCL в случае проблем, а если все прошло хорошо, то CheckResults сравнивает результаты с теми, которые мы насчитали в первой части на хосте и рассчитывает статистику ошибок.
Часть третья — Создание и запуск ядра
Первым делом мы выводим имя устройства, полученное таким путем:
Таким же образом можно получить информацию о количестве ядер, частоте, версии, итд итп
Затем мы создаем контекст:
С контекстами все не так просто… При создании контекста, мы передаем список устройств, которые мы хотим в него включить, но тут есть ограничение: только устройства на одной платформе могут быть в одном контексте, т.е. сделать контекст с GPU и CPU (в случае Intel/NVidia) не получится. В случае нескольких устройств в одном контексте, все буферы будут синхронизироваться автоматически на разных устройствах. С одной стороны, это упрощает поддержку multi-GPU, а с другой стороны никто не знает как, что и когда драйвер будет синхронизировать, а эффективность передачи данных является критичным для получения высокой производительности ради которой все и затевается. Поэтому я обычно создаю отдельный контекст для каждого устройства и вручную распределяю данные. Таким образом всегда известно, что, где, когда происходит.
Следующий шаг — это создание очереди команд для устройства:
Эта самая очередь привязывается к конкретному устройству и, в теории, может быть Out of Order, но по факту, я такого поведения не замечал. Очередей для одного устройства может быть несколько, причем можно синхронизировать команды из разных очередей, но в пределах одного контекста.
Далее мы создаем буферы для входных и выходного векторов:
При создании буфера указывается контекст (а не конкретное устройство), его объем и, при желании и использовании флага CL_MEM_COPY_HOST_PTR, указатель на данные, которые будут в него скопированы при создании. Как я говорил ранее, C++ wrapper использует подсчет ссылок, поэтому удалять буфер вручную не надо, в отличие от чистого C API.
Далее нам необходимо создать ядро, код которого хранится в файле «OpenCLFile1.cl». Для этого мы читаем текст из файла, создаем OpenCL программу, компилируем ее и получаем из нее ядро с именем «TestKernel», которое вы увидите в следующей части.
При компиляции надо указать на каких устройствах мы ее планируем запускать, в нашем случае это одно выбранное устройство для теста, хотя можно указать все сразу. Также можно передавать флаги компиляции, но в этом примере мы этого не делаем.
Далее нам нужно установить аргументы, которые будут передаваться ядру. В отличие от CUDA, нужно вызывать специальные функции (в случае C++ wrapper’а, методы) для каждого аргумента и при необходимости указывать размер аргумента.
Теперь мы подошли к самому главному — запуску ядра:
Собственно queue.enqueueNDRangeKernel добавляет команду запуска ядра в очередь команд и устанавливает количество элементов, которые будут обработаны, а также размер группы. О группах я расскажу отдельно (в другой статье), но сейчас упомяну лишь тот факт, что все элементы всегда разбиваются на группы и от размера группы может сильно зависеть производительность. В нашем случае количество элементов равно DATA_SIZE, а размер группы 128. Во время выполнения ядра, оно будет запущено DATA_SIZE раз (в неизвестной последовательности и возможно одновременно) и при каждом запуске ему будет передана информация о том, какой именно элемент обрабатывается.
enqueueNDRangeKernel является не блокирующей, поэтому после запуска ядра, мы должны дождаться его завершения, для чего и служит:
Фактически finish выполняет две задачи:
1) Пересылает все команды в устройство (выполнение enqueueNDRangeKernel гарантирует, что драйвер получил команду и поставил ее в очередь, но не гарантируют ее запуск на устройстве, причем довольно часто может проходить достаточно длительное время перед реальным запуском ядра).
2) Ждет завершения всех команд в очереди.
Если нужно выполнить только первую часть, существует команда push (clFlush), которая является не блокирующей, но заставляет драйвер начать выполнение команд из очереди.
После выполнения расчетов, мы подсчитываем затраченное время и загружаем результаты расчетов обратно на хост командой:
В зависимости от второго аргумента, enqueueReadBuffer может быть блокирующей или не блокирующей. В нашем случае, она блокирующая, поэтому нет необходимости вызывать finish отдельно. Синтаксис простой: первый аргумент — откуда читать, четвертый аргумент — сколько читать и последний аргумент — куда читать. Есть еще параметр, который задает смещение от начала входного буфера, которое надо использовать в случае, если нужно считать данные не сначала, так как мы не можем использовать адресную арифметику для буферов OpenCL на хосте.
Часть четвертая — Код OpenCL kernel
А вот тут мы и дошли до того места, где нам надо начинать писать код (хотя это и кодом назвать сложно, так… баловство:)) на OpenCL. Вот так выглядит OpenCLFile1.cl:
Итак по порядку:
Первым делом мы включаем в наш код файл MathCode.cpp, который содержит математическую функцию, ту самую на которую я просил обратить внимание ранее и ту самую, которая используется для традиционных вычислений на хосте. Как вы видите, мы даже не копируем код, мы используем один и тот же файл с математическим кодом.
Дальше мы создаем ядро, которое помечаем ключевым словом __kernel. Некоторые аргументы ядра также помечены ключевым словом __global, которое указывает на то, что это буфер в глобальной памяти устройства, созданный нами в коде хоста.
В коде ядра мы получаем номер элемента, который необходимо обработать:
Параметр get_global_id указывает на измерение, так как обрабатываемые элементы могут представлять из себя 1, 2 или 3мерный массив.
Затем проверяем граничные условия:
Это необходимо делать по той причине, что количество элементов для обработки должно быть всегда кратно размеру группы и таким образом оно может превышать количество, которые нужно обработать.
А после проверки мы делаем главную часть: вычисления, причем точно таким же образом, как и на хосте:
Часть пятая — Тестирование и замеры производительности
Вот и пришло время запустить приложение, оценить производительность и сделать некоторые выводы.
Я запускал тест на двух машинах и получил интересные результаты:
Ноутбук (CPU: Intel® Core™ i7-820QM, GPU: NVidia Quadro FX 2800M):
Итак, приступим к разбору результатов, а результаты, надо сказать, очень даже впечатляющие. GPU на ноутбуке в
110X быстрее хоста, а на десктопе и вовсе в
340X быстрее, впечатляющий результат, однако. Перед тем, как в меня начнут бросать тапки и говорить, что такое сравнение не правильное, я скажу, что в нем действительно есть несколько лукавств, но не более того.
Во-первых, я мы тут не учитываем время копирования данных на устройство и обратно. С одной стороны, это неправильно, так как с учетом копирования все может выглядеть не так радостно. С другой стороны, копирование можно выполнять одновременно с вычислениями, а может его и вовсе не нужно производить, если данные уже находятся на устройстве. В общем все далеко не так однозначно и зависит от конкретной задачи.
Во-вторых, помните как выглядел математический код? Для тех, кто не смотрел на него, скажу, что это много много математических операций над одними и теми же данными, причем получился он путем простого копипаста и замены цифр в коэффициентах, а изначально он был проще и занимал всего одну строку, только вот когда я начал тестировать, результаты были не такие радостные, GPU было всего в 4-5 раз быстрее. Как думаете, почему? (вопрос риторический, можно не думать:)). А все просто, мы уперлись в производительность памяти. Я надеюсь, что попозже у меня дойдут руки и я напишу статью о взаимосвязи производительности памяти и процессора, но это отдельная история, в этой статье нам интересен лишь тот факт, что с данным ядром у нас получился чистый тест арифметической производительности процессора.
Учитывая эти два момента, можно сказать, что GPU действительно в сотни раз быстрее не-параллельного кода на CPU для чистой арифметики, что в целом, соответствует разнице в теоретической производительности. (Еще одна надежда на то, что дойдут руки замерить реальные цифры и их соответствие теории для другой статьи).
Но о том, что GPU быстро считает мы знаем, а в результате нашего теста получилось, что и CPU выполняет OpenCL код довольно быстро, если быть точным, то в 13X и 25Х раз быстрее, чем обычный код скомпилированный MSVC10 с дефолтными настройками. Давайте разбираться, как так получается и откуда взялись эти цифры.
Оба процессора содержат 4 реальных и 8 виртуальных ядер, а OpenCL как раз и сделан для того, чтобы все ядра использовать, но улучшение у нас гораздо больше, чем 4Х. А тут надо сказать спасибо Intel, которая в своей реализации OpenCL, добавила поддержку автоматической векторизации, т.е. без каких-либо изменений в коде, OpenCL использует SSE или AVX, в зависимости от того, что доступно. Учитывая, что SSE у нас 128битное, а AVX работает с 256битами, получается, что производительность должна подняться в 16X и 32X соответственно. Это уже ближе к истине, но все еще не совсем точное совпадение. А дальше нам надо вспомнить о такой радостной штуке, как TurboBoost. Процессоры эти работают на частотах 1,73GHz/3,06GHz (ноутбук) и 3,4GHz/3,8GHz (десктоп), но по факту могу сказать, что частота ноутбучного процессора скачет от 1,73 до 2,8 непрерывно, да и греется он весьма сильно (тут следует бросить большой каметь в Dell за кривую систему охлаждения), поэтому реально во время теста частоты 3,06GHz сколь нибудь значимое время мы не увидим. Плюс не надо забывать, что практический результат всегда меньше теоретически возможного (десктоп по идее должен работать быстрее), но как мы видим, 25Х улучшение производительности можно получить практически бесплатно на одном и том же железе.
Заключение
Задачей этой статьи не была попытка объяснить все детали работы с OpenCL, скорее это была попытка показать, что все не так уж сложно (вот тут я уже писал, что не все так просто) и в идеальных условиях можно получить очень впечатляющую производительность, причем даже на одном и том же железе, да к тому же можно использовать один и тот же код для всех устройств. Но помните, что это почти идеальные условия, которые бывают далеко не всегда.
PS: Для тех, кто хочет побаловаться с кодом и посмотреть тесты на другом железе, проект (и даже собранный экзешник) лежит на гитхабе. Для запуска может понадобиться OpenCL SDK от производителей вашего железа.
PS2: Если у кого-нибудь есть Ivy Bridge, было бы интересно посмотреть на тест встроенного видеоядра. Дело в том, что в последней версии OpenCL SDK, Intel открыла доступ к IGP, но только для последнего поколения процессоров, а таких у меня под рукой нет. Да и на результаты AMD интересно взглянуть.