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

Как было написано	float4 val = (0, 0, 0, 0);
Что хотел написать автор	float4 val = (float4)(0, 0, 0, 0);
Как нужно было написать	float4 val = 0;

Если Вы сталкивались с OpenCL или планируете столкнуться и не видите разницы между первым и вторым вариантом, а третий вызывает у Вас сомнения — «А скомпилируется ли вообще?» — добро пожаловать под кат, там много нюансов языка и совсем ничего про API и оптимизацию производительности.

Самая дорогая компьютерная память — в голове программистов. Вероятно именно поэтому две самые популярные технологии программирования на GPU — CUDA и OpenCL — основаны не на принципиально новых параллельных языках, не на ассемблерах для конкретных архитектур, а на адаптированных C++ и C. В случае OpenCL самый популярный язык написания ядер это OpenCL C 1.2 — диалект C на основе ISO C99. Убрана стандартная библиотека, добавлены векторные типы и операции над ними, несколько функций для блокировок и определения своего места среди прочих потоков, четыре адресных пространства. Добавлены простые атомарные операции, несовместимые с C11 (атомарные операции и блокировки из C11 добавлены в OpenCL C 2.0, который пока что не получил широкого распространения). Добавлены некоторые удобные функции которых не было в C, что-то вроде стандартизованных интринсиков.

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

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

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

Как видно из списка, всё дело в преобразовании типов. В спецификации OpenCL 1.2 это разделы 6.2.* Conversions and Type Casting. Кроме того, коварен следующий раздел 6.3 Operators, который тоже никто не читает. Как показывает опыт, многое в спецификации написано недостаточно ясно и слишком скучно — попробую восполнить пробел в доступной русскоязычной документации на эти темы данной статьёй.

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

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

(векторный тип)(значения скаляров или векторов)

Например

(int2)(1,2);

или

// задаём вектор из двух целых
int2 a = (int2)(1, 2);
// a=[1,2]

// ещё один вектор из двух целых
int2 b = (int2)(3, 4);
// b=[3,4]

// вектор из четырёх целых получается склеиванием пары двухкомпонентных векторов
int4 c = (int4)(a, b);
// c=[1,2,3,4]

// трёхкомпонентный вектор получается из скаляра и двухкомпонентного вектора
int3 d = (int3)(1, c.xy);
// d=[1,1,2]

// как это может не быть приведением типов?!
float2 e = (float2)(1);
// e=[1.0f,1.0f]

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

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

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

Реальный пример:

int2 coords = (get_global_id(0), get_global_id(1));

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

int2 coords = (int2)(get_global_id(0), get_global_id(1));

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

get_global_id(0);
int2 coords = get_global_id(1);

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

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

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

// серый цвет, всё выглядит правильно
float3 color = (0.5f, 0.5f, 0.5f);
// color=[0.5f, 0.5f, 0.5f]

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

// хотели синий, получили опять серый
float3 color = (0.1f, 0.1f, 0.5f);
// color=[0.5f, 0.5f, 0.5f]

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

// правильный вариант
float3 color = (float3)(0.1f,0.1f, 0.5f);
// color=[0.1f,0.1f, 0.5f]

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

int val = true;
int2 val2 = true;

Какое значение лежит в 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):

When casting a bool to a vector integer data type, the vector components will be set to -1 (i.e. all bits set) if the bool value is true and 0 otherwise.

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

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

__kernel void bool_to_int_vec()
{
  int val = true;
  int2 val2 = true;
  printf("int val = true; // val=%d\n", val);
  printf("int2 val2 = true; // val2=%v2d\n", val2);
  if(val2.x == -1 && val2.y == -1)
  {
    printf("Compiler follows specification for bool->intn conversion, OK\n");
  }
  else
  {
    printf("Compiler does not follow specification for bool->intn conversion, FAILED\n");
  }
}

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

На своей машине я в результате эксперимента с двумя платформами по два устройства в каждой наблюдал:

$ ./clrun ../kernels/bool_to_int_vec.cl
...
Running "bool_to_int_vec" kernel on AMD Accelerated Parallel Processing / Tonga
int val = true; // val=1
int2 val2 = true; // val2=-1,-1
Compiler follows specification for bool->intn conversion, OK
...
Running "bool_to_int_vec" kernel on AMD Accelerated Parallel Processing / Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
int val = true; // val=1
int2 val2 = true; // val2=1,1
Compiler does not follow specification for bool->intn conversion, FAILED
...
Running "bool_to_int_vec" kernel on Intel(R) OpenCL / Intel(R) HD Graphics
int val = true; // val=1
int2 val2 = true; // val2=1,1
Compiler does not follow specification for bool->intn conversion, FAILED
...
Running "bool_to_int_vec" kernel on Intel(R) OpenCL / Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
int val = true; // val=1
int2 val2 = true; // val2=1,1
Compiler does not follow specification for bool->intn conversion, FAILED

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

Спустя некоторое время я проверил то же ядро на другой машине с тремя реализациями OpenCL — от AMD, Intel и NVidia:

...
Running "bool_to_int_vec" kernel on AMD Accelerated Parallel Processing / Ellesmere
int val = true; // val=1
int2 val2 = true; // val2=-1,-1
Compiler follows specification for bool->intn conversion, OK
...
Running "bool_to_int_vec" kernel on AMD Accelerated Parallel Processing / Intel(R) Core(TM) i5-7400 CPU @ 3.00GHz
int val = true; // val=1
int2 val2 = true; // val2=1,1
Compiler does not follow specification for bool->intn conversion, FAILED
...
Running "bool_to_int_vec" kernel on Intel(R) OpenCL / Intel(R) HD Graphics 630
int val = true; // val=1
int2 val2 = true; // val2=-1,-1
Compiler follows specification for bool->intn conversion, OK
...
Running "bool_to_int_vec" kernel on Intel(R) OpenCL / Intel(R) Core(TM) i5-7400 CPU @ 3.00GHz
int val = true; // val=1
int2 val2 = true; // val2=-1,-1
Compiler follows specification for bool->intn conversion, OK
...
Running "bool_to_int_vec" kernel on NVIDIA CUDA / GeForce GTX 1060 6GB
int val = true; // val=1
1,1
Compiler does not follow specification for bool->intn conversion, FAILED


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

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

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

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

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

int a = 1 > 0; // a=1
int4 b = (int4)(1) > (int4)(0); // b=[-1,-1,-1,-1]

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

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

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

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

int a = 1 ? 1 : 0; // a=1
int4 b = (int4)(1, 0, 1, 0) ? (int4)(1) : 0; // b=[0,0,0,0]
int4 c = (int4)(-1, 0, -1, 0) ? 1 : (int4)(0); // c=[1,0,1,0]
int4 d = (uint4)(0x80000000u, 0, 0, 0) ? (int4)(1) : (int4)(0); // d=[1,0,0,0]

// в C99 это допустимо, но не все компиляторы OpenCL С позволяют так делать
float e = 0.0f ? 1 : 2;

// ошибка компиляции, expr1 должно быть вектором целых
float4 f = (float4)(0) ? (float4)(1) : (float4)(2);

// expr2 и expr3 могут быть вещественными векторами
float4 g = (int4)(-1, 1, -1, 1) ? (float4)(1) : (float4)(0);
// g=[1.0f, 0.0f, 1.0f, 0.0f]

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

int a = 1 ? 1 : 0; // a=1
int4 b = (int4)(1) ? (int4)(1) : (int4)(0); // b=[0, 0, 0, 0]

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

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

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

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

The OpenCL language extends the union to allow the program to access a member of a union object using a member of a different type.

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

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

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

В OpenCL пункты 1 и 2 аналогичны C99, пункты 3 и 4 — новшества для удобства и ясности работы с векторными типами.

Неявные преобразования и явное приведение типа в стиле C

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

float a = 5.1f;
int   b = a;      // b = 5
float c = 1.6f;
int   d = (int)c; // d = 1 

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

float  a = 4.7f;
float4 b = 5;       // int 5 → float 5.0f → (float4)(5.0f, 5.0f, 5.0f, 5.0f)
int4   c = 4.7f;    // float 4.7f → int 4 → (int4)(4, 4, 4, 4)
int4   d = 1;       // int 1 → (int4)(1, 1, 1, 1)
int4   e = (int4) a;// float 4.7 → int 4 → (int4)(4, 4, 4, 4) явное приведение типа для скаляра
float4 f = a;       // float 4.7f → (float4)(4.7f, 4.7f, 4.7f, 4.7f)

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

float4 a = (float4)(5.0f, 5.0f, 5.0f, 5.0f); // это векторный литерал, а не приведение типа
int4   b = a; // ошибка, неявного преобразования векторных типов нет в языке
float4 c = 0;
int4   d = (int4)c; // ошибка, явного привидения типа вектора нет в языке
int4   e = (int4)(c); // ошибка, такой векторный литерал составить нельзя — в языке нет неявного преобразования float4 в int4
int4   f = (int4)(c.xy, c.zw); // ошибка, такой векторный литерал составить нельзя — в языке нет неявного преобразования float2 в int2
int4   g = (int4)(c.x, c.y, c.z, c.w); // а такой векторный литерал составить можно, потому что скаляры неявно преобразовались из float к int

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

float2 a = (float2)(1); // векторный литерал
float2 b = (float2)1;   // явное приведение типа скаляра к вектору
float2 c = 1;           // неявное преобразование типа скаляра к вектору

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

float2 a, b, c, d;

// это векторный литерал
a = (float2)(1, 2);
// a=[1, 2]

// 1 явно приводится к вектору из двух единиц, 2 — игнорируется
b = (float2)1, 2;
// b=[1, 1]

// 1 неявно приводится к вектору из двух единиц, 2 — игнорируется
c = 1, 2;
// c=[1, 1]

// 1 игнорируется, 2 неявно приводится к вектору из двух двоек
d = (1, 2);
// d=[2, 2]

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

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

convert_результирующийТип(исходныйТип)

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

convert_результирующийТип<_sat><_режимОкругления>(исходныйТип)

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

float  a = 5.5f;
int    b = convert_int(a);    // b = 5
float4 c = a; // c=[5.5, 5.5, 5.5, 5.5]
float2 d = convert_float2(c); // ошибка, вектор из четырёх компонентов нельзя преобразовать к вектору из двух

// а вектора с разными типами но одинаковым количеством
// компонентов преобразовывать как раз можно и нужно
int4 e = convert_int4(c);
// e=[5,5,5,5]

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

int a   = 257;
uchar b = convert_uchar(a);     // b = 1, сработало переполнение
b       = convert_uchar_sat(a); // b = 255, сработало насыщение

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

int a = convert_int_rtp(4.2f); // a = 5
a     = convert_int(4.2f);     // a = 4
int4 b = convert_int4_rte((float4)M_PI_F); // b = [3, 3, 3, 3]

Преобразование 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_тип() для скаляров и векторов:

float a = 25.0f;
int b = as_int(a);
// b=0x41C80000, что соответствует двоичному представлению 25.0f

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

int a    = 0;
char b   = as_char(a);  //ошибка, sizeof(int)!=sizeof(char)
float2 c = 0;
float8 d = as_float8(c); //ошибка, sizeof(float2)!=sizeof(float8)

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

uint word = 0x01020304;
uchar4 bytes = as_uchar4(word);

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

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

float4 a = 1.0f;
int3 b = as_int3(a); // работает, так как sizeof(int3)==sizeof(float4)
// b=[0x3f800000, 0x3f800000, 0x3f800000]
char3 c = as_char3(a); // ошибка, sizeof(char3)!=sizeof(float4) 

Заключение

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

• никогда не спутаете векторный литерал с явным приведением типов;
• не попадётесь на логических операторах в векторных типах;
• добавите в свой арсенал функции convert_* и as_*;