opencl
Основы ядра

Вступление

Ядро серого

Давайте создадим ядро ​​для генерации изображения в градациях серого. Мы будем использовать данные изображения, которые определяются с помощью uints для каждого компонента и с порядком RGBA.

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

Теперь пройдите через этот код шаг за шагом. Первая строка создает переменную в области памяти __constant типа sampler_t. Этот сэмплер используется для дальнейшего определения доступа к нашим данным изображения. Пожалуйста, обратитесь к документации Khronos для полной документации.

Мы выделили вход как read_only и вывод как write_only, прежде чем мы позвонили в наше ядро, поэтому добавим сюда эти модификаторы.

image2d и image3d всегда выделяются в глобальной памяти, поэтому здесь можно опустить модификатор __global.

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

После того, как мы убедились, что мы являемся законной нитью, мы читаем наш пиксель из нашего входного изображения. Затем мы преобразуем его в float, чтобы избежать потери десятичных знаков, делать некоторые вычисления, преобразовывать их обратно и записывать в вывод.

Ядро Скеллетон

Позволяет пройти через самое простое ядро ​​и некоторые его варианты

__kernel void myKernel() {
}

Ядро, которое можно запустить из основного кода, идентифицируется ключевым словом __kernel. Функция Kernel может иметь только возвращаемый тип void.

__kernel void myKernel(float a, uint b, byte c) {

}

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

Функция может выставлять переменные, как любая другая функция C / C ++. Единственное различие заключается в том, когда вы хотите ссылаться на память. Это относится ко всем указателям, являются ли они аргументами или используются в коде.

float*  ptr;

является указателем на область памяти, к которой имеет доступ только исполняемый поток. На самом деле это то же самое, что

__private float* ptr;

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

• __global: Этот модификатор относится к указателю, который помещается в глобальную память
• __constant: относится к указателю постоянной памяти
• __local: относится к указателю общей памяти
• __private: ссылается на указатель локальной памяти

Кроме того, мы можем определить, как мы хотим получить доступ к памяти

• нет модификатора: чтение и запись
• __read_only
• __write_only

Эти флаги должны соответствовать тому, как мы выделили буфер памяти на нашем хосте.

Идентификатор ядра

Чтобы правильно работать с данными, каждый поток должен знать свою позицию в пуле threadblock / global thread. Это может быть

get_local_id($dim);
get_global_id($dim);

Эти две функции возвращают положение нити относительно поточного блока или всех потоков.

get_working_dim();

Получает общее количество измерений, с которыми было запущено ядро.

get_local_size($dim);
get_global_size($dim);

Получает общее количество потоков в потоковом блоке или в целом для данного измерения.

Предостережение: всегда убедитесь, что ваш поток не превышает ваш размер данных. Это очень вероятно, и он всегда должен быть проверен.

Векторы в OpenCL

Каждый базовый тип opencl имеет векторную версию. Вы можете использовать векторный тип, добавив количество желаемых компонентов после типа. Поддерживаемое количество компонентов - 2,3,4,8 и 16. OpenCL 1.0 не предлагает трех компонентов.

Вы можете инициализировать любой вектор двумя способами:

• Обеспечить единый скаляр
• Удовлетворите все компоненты

float4 a = (float4)(1); //a = (1, 1, 1, 1)

или же

float4 b = (float4)(1, 2, 3, 4);
float4 c = (float4)(1, (float3)(2));

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

a[0] = 2;

или использовать литералы. Преимущество литералов в том, что вы можете комбинировать их так, как вы хотите, хорошо сделайте это в одно мгновение. Вы можете получить доступ ко всем компонентам вектора с помощью

a.s0 = 2; // same as a[0] = 2

вы также можете объединить несколько компонентов в новый вектор

a.s02 = (float2)(0, 0); // same as  a[0] = 0; a[2] = 0; or even a.s20 = (float2)(0, 0)

вы можете изменить порядок компонентов так, как хотите.

a.s1423 = a.s4132; // flip the vector

но вы не можете сделать что-то вроде

a.s11 = ... // twice the same component is not possible

Есть некоторые удобные сокращения для доступа к векторным компонентам. Следующие сокращения применяются только к размерам 2, 4, 8 и 16

a.hi //=a.s23 for vectors of size 4, a.4567 for size 8 and so on.
a.lo //=a.s01
a.even //=a.s02
a.odd //=a.13

Для векторных размеров 2,3 и 4 есть несколько дополнительных сокращений

a.x //=a.s0
a.y //=a.s1
a.z //=a.s2
a.w //=a.s3

Ядро исправления гамма-коррекции

Давайте посмотрим на ядро ​​коррекции гамма-коррекции

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Gamma(__read_only image2d_t input, __write_only image2d_t output, __constant float gamma) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color = pow(color, (float4)(gamma));
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

Теперь пройдите через этот код шаг за шагом. Первая строка создает переменную в области памяти __constant типа sampler_t. Этот сэмплер используется для дальнейшего определения доступа к нашим данным изображения. Пожалуйста, обратитесь к документации Khronos для полной документации.

Мы выделили вход как read_only и вывод как write_only, прежде чем мы позвонили в наше ядро, поэтому добавим сюда эти модификаторы.

image2d и image3d всегда выделяются в глобальной памяти, поэтому здесь можно опустить модификатор __global. Наше гамма-значение находится в __константной памяти, поэтому мы также указываем это.

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

После того, как мы убедились, что мы являемся законной нитью, мы читаем наш пиксель из нашего входного изображения. Затем мы преобразуем его в float, чтобы избежать потери десятичных знаков, делать некоторые вычисления, преобразовывать их обратно и записывать в вывод.