Эта статья посвящена основам программирования на OpenCl. OpenCl -это язык программирования на GPU/CPU, по своей структуре близкий к стандарту c99. Его развитием занимается Khronos Group, где на их сайте доступна полная документация. Во избежание полемики на тему «ну это же всё тривиально, достаточно покопаться в инете» сразу оговорюсь: в рунете информация на эту тематику практически полностью отсутствует, а в западном инете доступна весьма в разрозненном состоянии на десятке сайтов. Здесь будет приведена некоторая компиляция базовых принципов, максимально упрощающая начинающему программисту жизнь, а так же позволяющая с самого первого проекта максимально задействовать вычислительные мощности видеокарты. Людям написавшим 2-3 серьёзных программы на OpenCl это будет уже неинтересно. Статья в некотором смысле является продолжением моей прошлой статьи.
В первую очередь вопрос: где писать сам код. Насколько мне известно под .NET пока что нет никакой свистелки, позволяющей обрабатывать код ядра непосредственно в студии. Поэтому приходиться использовать сторонние редакторы. AMD, nVidia и Intel прилагают их к своим пакетам SDK. Мне почему-то больше нравиться именно Интеловский. Так же, как вариант, есть несколько редакторов, написанных фанатами. Из них мне больше всего нравиться редактор, прилагающийся к OpenCLTemplate. Стоит отметить, что это именно редакторы. Компиляция кода происходит непосредственно перед запуском на GPU/CPU.
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «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, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.
Так как в основе языка лежит практически стандартный с++ я буду рассматривать только те особенности, которые отличают OpenCL от него. Рассмотрим код простейшей программы ядра, который я приводил в прошлой статье. Этот код складывает два вектора, v1 и v2, кладя результат в первый из них.
В первую очередь в глаза бросается загадочный "__kernel ". Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.
Тип данных "__global " обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть "__local ". К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является "__privat ". Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе "__global ". Во-первых, это "__constant ", который может использоваться только для чтения. Во-вторых, это "__read_only", "__write_only" и "__read_write" — структуры, использование которых разрешено только для изображений.
После запуска на видеокарте все процессы равнозначны и исполняют равнозначный код. Но, очевидно, нам не нужно многократное повторение одного и того же действия — каждый процесс должен делать свой кусок задачи. Для осознания своего места в окружающем мире служат идентификаторы процессов. Самый простой идентификатор — " get_global_id(0)". В случае приведённого примера он указывает на i номер вектора, который должен сложить этот процесс. Если же мы обрабатываем не одномерный вектор, а двухмерное изображение — нам нужно знать положение процесса по двум осям. Конечно, это значение можно вычислить. Но это лишние операции. Поэтому для удобства при запуске можно указать, что нам нужно пространство двухмерной размерности. Тогда в процессе можно получить оба идентификатора положения: «get_global_id(0)», «get_global_id(1)». Так же можно сделать и для трёхмерного пространства. Часто может потребоваться и размерности пространства в которой мы работаем. Например для изображения практически при любой его обработке нам нужны его ширина и высота. Для получения размерности пространства используется идентификатор «get_global_size(i)». Кроме этого есть идентификаторы процессов внутри рабочей группы — «get_local_id(i)», «get_local_size(i)» и идентификатор самой группы -" get_group_id(i)", «get_num_groups(i)». Большая часть этих соотношений связанна друг с другом: num_groups * local_size = global_size, local_id + group_id * local_size = global_id, global_size % local_size = 0.
Разработчики OpenCL и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Самой первой особенностью являются вектора и векторная математика. В OpenCl можно объявлять в качестве переменных 2, 4, 8 и 16-мерные вектора. Это делается соответственно: int2, int4, int8, int16. Так же можно объявлять double, byte и все остальные типы. Вектора соответствующей размерности можно складывать/вычитать/делить/умножать, а так же любой вектор можно делить/умножать на число:
Кроме того есть ряд функций, оптимизированных под вектора и позволяющих работать непосредственно с ними. К таким функциям относятся функции вычисления расстояния, функции векторного произведения. Например:
Ещё можно сливать вектора, беря части от одного и другого, а так же склеивать в более большие:
Следующей особенностью 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. Эти функции выполняются быстрее, чем соответствующие им математические действия.
На сайте www.cmsoft.com.br есть замечательный пример, показывающий возможности при оптимизации кода средствами native и common функций:
Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.
Стоит отметить, что в OpenCl есть ряд директив, позволяющих включать различные дополнительные функционалы. Причин к этому две. Первая — исторически не все из этих функционалов поддерживались. Вторая — эти функционалы могут влиять на производительность. Обычно функционал включается следующей командой:
Для примера. Слудющие команды включают: возможность использования типа byte, двойную точность рассчётов и все математические функции
Часто в вычислениях нужна синхронизация. Это достигается несколькими способами. Первый — барьеры. Барьер это такая команда, на которой остановится процесс пока все прочие процессы или процессы его рабочей группы не достигнут. Приведём два примера:
В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды "__local int x[10];" и "__global int x[10];". Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Второй вариант синхронизации между потоками — atomic. Это функции, предотвращающие одновременное обращение к памяти. Перед их использованием нужно включить следующие директивы:
Простейший пример работы этой функции:
Если бы вместо «atom_inc(&num[0]);» было написано num++, то результат исполнения программы был непредсказуем, так как все процессы бы одновременно обратились к памяти и считали там одно и то же значение. Всего есть 11 функций еденичных операций: «add, sub, xchg, inc, dec, cmp_xchg, min, max, and, or, xor».
При помощи этих функций не сложно создать семафор.
Последней вещью, которую я хочу включить в этот 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).
Краткий пример. Замыливаем изображение по среднему значению в области:
Ну, думаю, с кратким описанием я справился. Теперь несколько ссылок для более подробного изучения, ежели кому понадобиться.
Официальный сайт с документацией.
Сайт с примерами и понятными описаниями.
Неплохой пдфничек, там структура OpenCl устройств неплохо прорисованна.
Так же есть 2 презентации про OpenCL на русском. В них довольно мало инфы и нет связующего текста. Правда есть примеры неплохие. Первая. Вторая.
Компилятор
В первую очередь вопрос: где писать сам код. Насколько мне известно под .NET пока что нет никакой свистелки, позволяющей обрабатывать код ядра непосредственно в студии. Поэтому приходиться использовать сторонние редакторы. AMD, nVidia и Intel прилагают их к своим пакетам SDK. Мне почему-то больше нравиться именно Интеловский. Так же, как вариант, есть несколько редакторов, написанных фанатами. Из них мне больше всего нравиться редактор, прилагающийся к OpenCLTemplate. Стоит отметить, что это именно редакторы. Компиляция кода происходит непосредственно перед запуском на GPU/CPU.
Модель памяти устройства
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «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, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.
Базовые особенности
Так как в основе языка лежит практически стандартный с++ я буду рассматривать только те особенности, которые отличают OpenCL от него. Рассмотрим код простейшей программы ядра, который я приводил в прошлой статье. Этот код складывает два вектора, v1 и v2, кладя результат в первый из них.
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
{
int i = get_global_id(0);
v1[i] = v1[i] + v2[i];
}
* This source code was highlighted with Source Code Highlighter.Объявление процедур
В первую очередь в глаза бросается загадочный "__kernel ". Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.
Типы памяти
Тип данных "__global " обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть "__local ". К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является "__privat ". Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе "__global ". Во-первых, это "__constant ", который может использоваться только для чтения. Во-вторых, это "__read_only", "__write_only" и "__read_write" — структуры, использование которых разрешено только для изображений.
Идентификаторы процессов
После запуска на видеокарте все процессы равнозначны и исполняют равнозначный код. Но, очевидно, нам не нужно многократное повторение одного и того же действия — каждый процесс должен делать свой кусок задачи. Для осознания своего места в окружающем мире служат идентификаторы процессов. Самый простой идентификатор — " get_global_id(0)". В случае приведённого примера он указывает на i номер вектора, который должен сложить этот процесс. Если же мы обрабатываем не одномерный вектор, а двухмерное изображение — нам нужно знать положение процесса по двум осям. Конечно, это значение можно вычислить. Но это лишние операции. Поэтому для удобства при запуске можно указать, что нам нужно пространство двухмерной размерности. Тогда в процессе можно получить оба идентификатора положения: «get_global_id(0)», «get_global_id(1)». Так же можно сделать и для трёхмерного пространства. Часто может потребоваться и размерности пространства в которой мы работаем. Например для изображения практически при любой его обработке нам нужны его ширина и высота. Для получения размерности пространства используется идентификатор «get_global_size(i)». Кроме этого есть идентификаторы процессов внутри рабочей группы — «get_local_id(i)», «get_local_size(i)» и идентификатор самой группы -" get_group_id(i)", «get_num_groups(i)». Большая часть этих соотношений связанна друг с другом: num_groups * local_size = global_size, local_id + group_id * local_size = global_id, global_size % local_size = 0.
Оптимизация расчётов
Разработчики OpenCL и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Встроенные вектора
Самой первой особенностью являются вектора и векторная математика. В OpenCl можно объявлять в качестве переменных 2, 4, 8 и 16-мерные вектора. Это делается соответственно: int2, int4, int8, int16. Так же можно объявлять double, byte и все остальные типы. Вектора соответствующей размерности можно складывать/вычитать/делить/умножать, а так же любой вектор можно делить/умножать на число:
uint4 sumall = (uint4)(1,1,1,1);
small += (uint4)(1,1,1,1);
sumall = sumall/2;
* This source code was highlighted with Source Code Highlighter.Кроме того есть ряд функций, оптимизированных под вектора и позволяющих работать непосредственно с ними. К таким функциям относятся функции вычисления расстояния, функции векторного произведения. Например:
float4 dir1 = (float4)(1, 1, 1, 0);
float4 dir2 = (float4)(1, 2, 3, 0);
float4 normal = cross(dir1, dir2);
* This source code was highlighted with Source Code Highlighter.Ещё можно сливать вектора, беря части от одного и другого, а так же склеивать в более большие:
int4 vi0 = (int4) -7 ;
int4 vi1 = (int4) ( 0, 1, 2, 3 ) ;
vi0.lo = vi1.hi; // слияние
int8 v8 = (int8)(vi0.s0123, vi1.s0123); //склейка
* This source code was highlighted with Source Code Highlighter.Простые функции
Следующей особенностью 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. Эти функции выполняются быстрее, чем соответствующие им математические действия.
Пример
На сайте www.cmsoft.com.br есть замечательный пример, показывающий возможности при оптимизации кода средствами native и common функций:
kernel void regularFuncs()
{
for (int i=0; i<5000; i++)
{
float a=1, b=2, c=3, d=4;
float e = a*b+c;
e = a*b+c*d;
e = sin(a);
e = cos(b);
e = a*b+c*d;
e = sin(a);
e = cos(b);
e = a*b+c*d;
e = sin(a);
e = cos(b);
float4 vec1 = (float4)(1, 2, 3, 0);
float4 vec2 = (float4)(-1, 3, 1, 0);
float4 vec = distance(vec1, vec2);
double x=1, y=2, z=3;
double resp = x*y+z;
}
}
kernel void nativeFuncs()
{
for (int i=0; i<5000; i++)
{
float a=1, b=2, c=3, d=4;
float e = mad(a,b,c);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
e = mad(a,b,c*d);
e = native_sin(a);
e = native_cos(b);
float4 vec1 = (float4)(1, 2, 3, 0);
float4 vec2 = (float4)(-1, 3, 1, 0);
float4 vec = fast_distance(vec1, vec2);
double x=1, y=2, z=3;
double resp = mad(x,y,z);
}
}
* This source code was highlighted with Source Code Highlighter.Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.
Разрешения
Стоит отметить, что в OpenCl есть ряд директив, позволяющих включать различные дополнительные функционалы. Причин к этому две. Первая — исторически не все из этих функционалов поддерживались. Вторая — эти функционалы могут влиять на производительность. Обычно функционал включается следующей командой:
#pragma OPENCL EXTENSION extension name : behavior
* This source code was highlighted with Source Code Highlighter.Для примера. Слудющие команды включают: возможность использования типа byte, двойную точность рассчётов и все математические функции
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
* This source code was highlighted with Source Code Highlighter.Синхронизация
Барьеры
Часто в вычислениях нужна синхронизация. Это достигается несколькими способами. Первый — барьеры. Барьер это такая команда, на которой остановится процесс пока все прочие процессы или процессы его рабочей группы не достигнут. Приведём два примера:
kernel void localVarExample()
{
int i = get_global_id(0);
__local int x[10];
x[i] = i;
barrier(CLK_LOCAL_MEM_FENCE);
if (i>0) int y = x[i-1];
}
kernel void globalVarExample()
{
int i = get_global_id(0);
__global int x[10];
x[i] = i;
barrier(CLK_GLOBAL_MEM_FENCE);
if (i>0) int y = x[i-1];
}
* This source code was highlighted with Source Code Highlighter.В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды "__local int x[10];" и "__global int x[10];". Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Единичные операции
Второй вариант синхронизации между потоками — atomic. Это функции, предотвращающие одновременное обращение к памяти. Перед их использованием нужно включить следующие директивы:
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
* This source code was highlighted with Source Code Highlighter.Простейший пример работы этой функции:
__kernel void test(global int * num)
{
atom_inc(&num[0]);
}
* This source code was highlighted with Source Code Highlighter.Если бы вместо «atom_inc(&num[0]);» было написано num++, то результат исполнения программы был непредсказуем, так как все процессы бы одновременно обратились к памяти и считали там одно и то же значение. Всего есть 11 функций еденичных операций: «add, sub, xchg, inc, dec, cmp_xchg, min, max, and, or, xor».
При помощи этих функций не сложно создать семафор.
void GetSemaphor(__global int * semaphor) {
int occupied = atom_xchg(semaphor, 1);
while(occupied > 0)
{
occupied = atom_xchg(semaphor, 1);
}
}
void ReleaseSemaphor(__global int * semaphor)
{
int prevVal = atom_xchg(semaphor, 0);
}
* This source code was highlighted with Source Code Highlighter.Работа с изображениями
Последней вещью, которую я хочу включить в этот 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)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;
int2 coords = (int2)(get_global_id(0), get_global_id(1));
uint4 sumall = (uint4)(0,0,0,0);
int sum = 0;
for(int i=-10;i<11;i++)
for(int j=-10;j<11;j++)
{
int2 newpol = (int2)(i,j)+coords;
sumall+= read_imageui(bmp1, smp, newpol);
sum++;
}
sumall = sumall/sum;
write_imageui(bmpOut, coords, sumall);
}
* This source code was highlighted with Source Code Highlighter.Полезности
Ну, думаю, с кратким описанием я справился. Теперь несколько ссылок для более подробного изучения, ежели кому понадобиться.
Официальный сайт с документацией.
Сайт с примерами и понятными описаниями.
Неплохой пдфничек, там структура OpenCl устройств неплохо прорисованна.
Так же есть 2 презентации про OpenCL на русском. В них довольно мало инфы и нет связующего текста. Правда есть примеры неплохие. Первая. Вторая.
