Что ещё необходимо узнать про OpenCL C перед тем, как на нём писать

По языку и API есть много примеров, большая часть из которых это сложение двух векторов. Есть хорошая, хоть и сухая, официальная спецификация, несколько англоязычных книг, советы по оптимизации от производителей устройств. Как только программист понимает, как написать свою задачу — он пишет её на знакомом C99 используя новые функции вроде get_global_id() и всё в плане языка кажется понятным и примитивным. Такой знакомый C99, что можно запросто угодить в ловушку и не заметить её до поры до времени. Да, OpenCL C очень похож на С, но в нём есть как очень полезные отличия, незаслуженно забытые потому что аналогов в C99 нет, так и очень коварные отличия, прячущиеся за похожий синтаксис.

Я просматривал много кода на OpenCL С и люди, которые только начинают на нём писать, делают одни и те же ошибки:

• путают векторный литерал с приведением типа;
• не используют замечательные механизмы преобразования типов;
• забывают о нюансах преобразования векторных типов.

Векторные литералы либо явное приведение типа

Новая конструкция в OpenCL C — векторный литерал, с помощью которого можно задать значение вектора. К сожалению, его синтаксис очень похож на явное приведение типа:

Однако (float2)(1) и другие примеры выше это не приведение типов, а новая конструкция (см. 6.1.6 Vector Literals в спецификации OpenCL 1.2).

Внутри вторых скобок должно быть суммарно столько скаляров либо компонентов вектора, сколько в векторном типе внутри первых скобок. Есть одно исключение — если справа только одно значение скаляра в скобках, то оно само «размножается» до необходимого количества компонентов вектора.

Явного приведения векторных типов в стиле C просто нет в языке. Роковая ошибка может быть допущена, если замыленными глазами увидеть «знакомое» приведение типа вместо векторного литерала. Тогда тип в скобках в начале можно убрать: «Ведь и так компилируется, зачем лишнее приведение типов? Уже неявно привелось».

coords задается не векторным литералом, для векторного литерала необходимо было добавить векторный тип:

У нас же получилось следующее: (get_global_id(0), get_global_id(1)) и это уже конструкция из обычного C — в скобках вызов двух функций через оператор « , » (запятая), который означает, что выполнятся обе функции и выражение вернёт результат второй функции, как если бы мы написали:

Сработает неявное преобразование скаляра в вектор (о нём чуть дальше) и в coords будет вектор [get_global_id(1), get_global_id(1)] , а не [get_global_id(0), get_global_id(1)] , как ожидалось.

К счастью, для простых случаев компилятор может выдать предупреждение вроде « warning: expression result unused », но рассчитывать на это не стоит.

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

Код работает, проект сдан. И вдруг понадобилось небольшое изменение — цвет из серого сделать тёмно-синим.

Надо было использовать векторный литерал:

Преобразование булевых значений в векторы

Какое значение лежит в val ? Какое — в val2 ?

Для скаляров действуют правила ISO C99, при преобразовании значения bool (а тип bool и константы true и false есть в C99 и в OpenCL C) false становится нулём, а true — единицей. Это правила для скаляров. Таким образом, в val будет «1». Не всегда это удобно, но такое поведение заложено в мозг программиста — конструкции типа x+=(a>b) уже не удивляют.

Однако, в OpenCL C при преобразовании к векторному целому типу значения типа bool возвращают либо целые со всеми битами в нуле, либо со всеми битами в единице, что соответствует (int)-1 . Вот что говорит на эту тему спецификация (раздел 6.2.2 Explicit Casts):

Таким образом, в val2 будет вектор [-1, -1] . Это немного неожиданно в контексте преобразования типа когда сначала выражение приводится к типу компонента вектора, а потом размножается — как для остальных типов, но для bool заявлено именно такое поведение. При грамотном использовании оно позволяет заменять условные выражения на побитовые операции.

Для проведения быстрых тестов вроде «Скомпилируется или нет? Какое значение в переменной?» я написал и выложил на гитхаб проект opencl-sandbox. Все примеры из этой статьи я проверил на своей машине. В том числе и такой:

Как известно, разработчики компиляторов тоже люди и не помнят спецификации наизусть.

Две OpenCL платформы — AMD и Intel. У каждой платформы по два устройства — GPU и CPU. И только компилятор AMD под GPU (самый зрелый) следует спецификации, остальные три записывают в val2 вектор из единиц, а не из -1.

Всего в системе пять устройств. Компиляторы AMD ведут себя так же. Более свежий компилятор от Intel «исправился» и теперь ведёт себя в соответствии со стандартом. Компилятор NVidia не справился не только с преобразованием в векторный тип, но и просто с отображением строки во втором printf() .

Выводов из этого два:

1. без знания спецификации переносимый код не написать;
2. необходимо покрывать OpenCL ядра тестами, потому что каждая платформа понимает спецификации по-своему.

Логические операторы и операторы сравнения для векторов

Как и для приведения bool к int , у соответствующих операторов поведение для скаляров и векторов разное. Значения результата выполнения операторов > , < , >= , <= , == , != , && , || , ! , это int . Для скаляров — 0 или 1. Для векторов — вектор соответствующей длины из int 'ов со значениями 0 или -1 (все биты выставлены в 1).

При проверке на 4-х компиляторах на этот раз все выдали правильный результат.

Тернарный оператор для векторов

Тернарный оператор вида « exp1 ? expr2 : expr3 » тоже ведёт себя аналогично по-разному для скаляров и векторов. Для скаляров — как в C99, результат выражения это expr2 если expr1 не ноль и exp3 если expr1 ноль.

Для векторов во-первых, тип expr1 может быть только целым. Во-вторых, при проверке условия в expr1 проверка идёт не на равенство нулю и даже не по первому биту, а по старшему биту. При этом оператор работает покомпонентно. Если одно из выражений expr2 и expr3 это вектор, а другое — скаляр, то скаляр неявно преобразуется к векторному типу с соответствующими компонентами.

Как видите, тут можно опять угодить в ловушку похожести. Сравните одинаковый с точностью до векторизации код:

Вектор b заполнен нулями, в полном соответствии со спецификацией и к недоумению программистов на C.

Преобразование вещественных и целых типов в OpenCL C

Для скалярных типов преобразования из целых типов в вещественные и из вещественных в целые производятся по тем же правилам, что в C99 — то есть при преобразовании из вещественного числа в целое у него отбрасывается дробная часть, при преобразовании из целого числа в вещественное получается вещественное число с тем же значением, что исходное целое. В случае, если число не влезает в диапазон типа, к которому происходит преобразование — результат зависит от реализации.

Если необходимо интерпретировать данные одного типа как данные другого, то единственный всегда работающий способ сделать это в C99 — использовать функцию memcpy . В OpenCL нет memcpy , зато в отличие от C99 абсолютно законно пользоваться union 'ами для интерпретации данных как данных другого типа:

Ко всему прочему, поддерживаются векторные типы и возможности железа по операциям с насыщением — это диктует особенности преобразования типов в OpenCL.

Поддерживаются следующие типы преобразований типов:

1. неявные преобразования (Implicit Conversions);
2. явное приведение (Explicit Casts);
3. явные преобразования (Explicit Conversions);
4. интерпретация данных как данные другого типа (Reinterpreting Data As Another Type).

Как и в C99, если в выражении встречаются операнды различных типов, то они преобразуются к одному общему типу. Разница в том, как это работает для векторов. Для скалярных типов неявное преобразование типа и явное приведение типа поддерживаются так же, как в C99:

При явном или неявном преобразовании из скалярного типа в векторный сначала скаляр приводится к типу элемента вектора по правилам аналогичным C99, а потом размножается до размера векторного типа:

Неявное преобразование и явное приведение в стиле C одного векторного типа в другой — запрещены. Даже если у них одинаковое количество компонентов.

Явного приведения векторных типов нет, однако скаляр привести к векторному типу можно. Это добавляет дополнительную путаницу к векторным литералам. Сравните три способа задать вектор с одинаковыми компонентами:

Для векторов с разными компонентами такой же код не сработает, нужно использовать только векторный литерал. Что самое плохое, весь приведённый ниже код отлично скомпилируется, просто результаты будут соответствующие:

Помимо приведения типов в стиле C, в OpenCL появился механизм приведения типов, который обрабатывает ситуации переполнения и работает с векторами. Это семейство функций

и более общие функции

которые дополнительно принимают режим работы при переполнении и вид округления. Для скаляров и векторов функции работают одинаково. Количество элементов в векторах исходного и результирующего типов должно совпадать.

При приведении к целым типам поведение при переполнении определяется опциональным модификатором _sat . Без него переполнение целого типа происходит как обычно в C99, с ним — работает насыщение, значения вне допустимого типом диапазона приводятся к максимально близкому значению, представимому в преобразованном типе:

При приведении к вещественным типам использование _sat не допускается. В этом нет необходимости, ведь при переполнении вещественных типов они и так становятся ±INF.

Для контроля над округлением предусмотрены модификаторы _rte (round to nearest even), _rtz (round toward zero), _rtp (round toward positive infinity) и _rtn (round toward negative infinity), которые обозначают округление до ближайшего целого, округление к нулю, округление к плюс бесконечности и округление к минус бесконечности соответственно. При отсутствии модификатора округления используется _rtz для преобразования из вещественных в целые и _rte при преобразовании из целых в вещественные. В _rte используется не привычный математический, а так называемый «банковский» вариант округления к ближайшему целому. Когда дробная часть ровно 0.5 то нет одного ближайшего целого числа, из двух ближайших выбирается чётное.

Преобразование float в int с разными режимами округления (проверено тут):

0.5 -0.5 1.1 -1.1 1.5 -1.5 1.7 -1.7 Округление к ближайшему целому (round to nearest even, rte) 0 0 1 -1 2 -2 2 -2 Округление к нулю (round toward zero, rtz) 0 0 1 -1 1 -1 1 -1 Округление к плюс бесконечности (round toward positive infinity, rtp) 1 0 2 -1 2 -1 2 -1 Округление к минус бесконечности (round toward negative infinity, rtn) 0 -1 1 -2 1 -2 1 -2 В англоязычной статье про округление на википедии есть замечательная иллюстрация. Режиму rte на ней соответствует «even», rtz — «round→zero», rtp — «round up», rtn — «round down».

Интерпретация данных как данных другого типа

Для интерпретации данных одного типа как данных другого типа в OpenCL существует, помимо union 'ов, семейство функций as_тип() для скаляров и векторов:

Если размер в байтах исходного и нового типов не совпадают, то as_тип должен вызвать ошибку компиляции:

Если количество элементов в исходном и новом типе не совпадает (но размеры типов одинаковы), то результат зависит от реализации OpenCL (implementation-defined), кроме случая когда операнд это 4-х компонентный вектор, а результат — 3-х компонентный вектор. Так, бывает удобно получить байты 32-х битного слова как элементы вектора:

Но результат при этом может быть как [4, 3, 2, 1] , так и [1, 2, 3, 4] , так и все что угодно, на усмотрение конкретной реализации OpenCL. Впрочем, при оптимизации и работе на какой-либо одной версии OpenCL подобное использование as_тип вполне допустимо.

Если операнд это 4-х компонентный вектор, а результат — 3-х компонентный вектор, то as_тип обязан возвратить биты исходного типа без изменений — по стандарту размеры векторов из трех компонент равны размеру векторов из четырех компонент, если размеры их элементов одинаковы.

Заключение

OpenCL C коварен в своей похожести на обычный C99. Надеюсь, после прочтения этой статьи Вы