opencl
Fundamentos del Kernel

Introducción

Kernel en escala de grises

Permite construir un kernel para generar una imagen en escala de grises. Utilizaremos los datos de imagen que se definen mediante uints para cada componente y con el pedido 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);
    }
}

Ahora veamos ese código paso a paso. La primera línea crea una variable en la __ región de memoria constante de tipo sampler_t. Esta muestra se utiliza para especificar aún más el acceso a nuestros datos de imagen. Consulte la documentación de Khronos para obtener una documentación completa.

Asignamos la entrada como read_only y la salida como write_only antes de llamar a nuestro kernel, así que agregamos esos modificadores aquí.

image2d y image3d siempre se asignan en la memoria global, por lo tanto, podemos omitir el modificador __global aquí.

Luego obtenemos nuestro ID de hilo que define el píxel que vamos a convertir a escala de grises. También consultamos el tamaño para asegurarnos de que nuestro hilo no está accediendo a la memoria no asignada. Esto definitivamente bloqueará tu kernel si olvidas eso.

Después de asegurarnos de que somos un hilo legítimo, leemos nuestro píxel fuera de nuestra imagen de entrada. Luego lo convertimos en flotante para evitar la pérdida de lugares decimales, hacemos algunos cálculos, lo convertimos y lo escribimos en la salida.

Kernel Skelleton

Caminemos por el núcleo más simple que hay y algunas variaciones de él.

__kernel void myKernel() {
}

Un kernel que puede iniciarse desde su código principal se identifica con la palabra clave __kernel. Una función de Kernel solo puede tener el tipo de retorno void.

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

}

Por supuesto, puedes crear más funciones que no estén expuestas como núcleos. En este caso, simplemente puede omitir el modificador __kernel.

Una función puede exponer variables como cualquier otra función C / C ++. La única diferencia es cuando quieres hacer referencia a la memoria. Esto se aplica a todos los punteros, ya sean argumentos o se utilicen en el código.

float*  ptr;

es un puntero a una región de memoria a la que solo tiene acceso el subproceso en ejecución. De hecho es lo mismo que

__private float* ptr;

Hay cuatro modificadores de región de memoria diferentes disponibles. Dentro del núcleo, normalmente no tiene que preocuparse por eso, pero para los argumentos es esencial.

• __global: este modificador se refiere a un puntero que se coloca en la memoria global
• __constante: se refiere a un puntero de memoria constante
• __local: se refiere a un puntero de memoria compartida
• __privado: se refiere a un puntero de memoria local

Además podemos definir cómo queremos acceder a la memoria.

• sin modificador: leer y escribir
• __solo lectura
• __escribir solamente

Esos indicadores deben coincidir con la forma en que asignamos el búfer de memoria a nuestro host.

ID de Kernel

Para trabajar correctamente con los datos, cada subproceso necesita conocer su posición en el conjunto de subprocesos / bloque de subprocesos. Esto puede ser archivado con

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

Esas dos funciones devuelven la posición del hilo en relación con el bloque de subprocesos o todos los subprocesos.

get_working_dim();

Obtiene el número total de dimensiones con las que se lanzó el kernel.

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

Obtiene el número total de subprocesos en el bloque de subprocesos o en total para una dimensión determinada.

Advertencia: asegúrese siempre de que su hilo no exceda el tamaño de sus datos. Es muy probable que esto ocurra y siempre se debe revisar.

Vectores en OpenCL

Cada tipo fundamental de opencl tiene una versión vectorial. Puede usar el tipo de vector agregando el número de componentes deseados después del tipo. El número de componentes admitidos es 2,3,4,8 y 16. OpenCL 1.0 no ofrece tres componentes.

Puedes inicializar cualquier vector de dos maneras:

• Proporcionar un solo escalar
• Satisfacer todos los componentes

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

o

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

o cualquier otra combinación de vectores que satisfagan el número de componentes. Para acceder a los elementos de un vector puedes utilizar diferentes métodos. Puedes usar indexación:

a[0] = 2;

o utilizar literales. La ventaja de los literales es que puedes combinarlos de la forma que quieras, así que hazlo en un momento. Puedes acceder a todos los componentes vectoriales con

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

También puedes combinar múltiples componentes en un nuevo vector.

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

Puede cambiar el orden de los componentes de la forma que desee.

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

pero no puedes hacer algo como

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

Hay algunos métodos abreviados convenientes para acceder a componentes vectoriales. Las siguientes abreviaturas solo se aplican a los tamaños 2, 4, 8 y 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

Para tamaños de vectores 2,3 y 4 hay algunas abreviaturas adicionales.

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

Kernel de corrección gamma

Veamos un núcleo de corrección gamma

__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);
    }
}

Ahora veamos ese código paso a paso. La primera línea crea una variable en la __ región de memoria constante de tipo sampler_t. Esta muestra se utiliza para especificar aún más el acceso a nuestros datos de imagen. Consulte la documentación de Khronos para obtener una documentación completa.

Asignamos la entrada como read_only y la salida como write_only antes de llamar a nuestro kernel, así que agregamos esos modificadores aquí.

image2d y image3d siempre se asignan en la memoria global, por lo tanto, podemos omitir el modificador __global aquí. Nuestro valor gamma se encuentra en la memoria __ constante, por lo que también lo especificamos.

Luego obtenemos nuestro ID de hilo que define el píxel que vamos a gamma correcto. También consultamos el tamaño para asegurarnos de que nuestro hilo no está accediendo a la memoria no asignada. Esto definitivamente bloqueará tu kernel si olvidas eso.

Después de asegurarnos de que somos un hilo legítimo, leemos nuestro píxel fuera de nuestra imagen de entrada. Luego lo convertimos en flotante para evitar la pérdida de lugares decimales, hacemos algunos cálculos, lo convertimos y lo escribimos en la salida.