opencl
Notions de base sur le noyau

Introduction

Noyau en niveaux de gris

Permet de construire un noyau pour générer une image en niveaux de gris. Nous utiliserons des données d'image qui sont définies à l'aide d'uints pour chaque composant et avec la commande 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);
    }
}

Maintenant, parcourons ce code pas à pas. La première ligne crée une variable dans la région de la mémoire __constant du type sampler_t. Cet échantillonneur est utilisé pour spécifier davantage l'accès à nos données d'image. Veuillez vous reporter à la documentation de Khronos pour une documentation complète.

Nous avons attribué l'entrée à read_only et la sortie à write_only avant d'appeler notre noyau. Nous ajoutons donc ces modificateurs ici.

image2d et image3d sont toujours allouées sur la mémoire globale, nous pouvons donc omettre le modificateur __global ici.

Ensuite, nous obtenons notre identifiant de thread qui définit le pixel que nous allons convertir en niveaux de gris. Nous interrogeons également la taille pour nous assurer que notre thread n’accède pas à la mémoire non allouée. Cela va définitivement bloquer votre noyau si vous l'oubliez.

Après nous être assuré que nous sommes un fil légitime, nous lisons notre pixel de notre image d'entrée. Nous le convertissons ensuite en float pour éviter la perte de décimales, effectuer des calculs, le reconvertir et l’écrire dans la sortie.

Noyau Skelleton

Permet de parcourir le noyau le plus simple qui existe et certaines variantes de celui-ci

__kernel void myKernel() {
}

Un noyau pouvant être démarré à partir de votre code principal est identifié par le mot clé __kernel. Une fonction du noyau ne peut avoir que le type de retour void.

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

}

Bien sûr, vous pouvez créer plus de fonctions qui ne sont pas exposées en tant que noyaux. Dans ce cas, vous pouvez simplement omettre le modificateur __kernel.

Une fonction peut exposer des variables comme n'importe quelle autre fonction C / C ++. La seule différence est lorsque vous souhaitez référencer la mémoire. Cela s'applique à tous les pointeurs, qu'ils soient des arguments ou utilisés dans le code.

float*  ptr;

est un pointeur sur une région mémoire à laquelle seul le thread d'exécution a accès. En fait, c'est la même chose que

__private float* ptr;

Quatre modificateurs de région de mémoire sont disponibles. Dans le noyau, vous n'avez généralement pas à vous en préoccuper, mais pour les arguments, c'est essentiel.

• __global: ce modificateur fait référence à un pointeur placé dans la mémoire globale
• __constant: fait référence à un pointeur de mémoire constante
• __local: fait référence à un pointeur de mémoire partagée
• __private: fait référence à un pointeur de mémoire locale

En outre, nous pouvons définir comment nous voulons accéder à la mémoire

• aucun modificateur: lire et écrire
• __lecture seulement
• __write_only

Ces drapeaux doivent correspondre à la manière dont nous avons affecté la mémoire tampon sur notre hôte.

ID du noyau

Pour travailler correctement avec les données, chaque thread doit connaître sa position dans le pool de threads / global thread. Cela peut être réalisé avec

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

Ces deux fonctions renvoient la position du thread par rapport au threadblock ou à tous les threads.

get_working_dim();

Obtient le nombre total de dimensions avec lesquelles le noyau a été lancé.

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

Obtient le nombre total de threads dans le threadblock ou au total pour une dimension donnée.

Avertissement: assurez-vous toujours que votre thread ne dépasse pas la taille de vos données. Cela risque fort de se produire et doit toujours être vérifié.

Vecteurs dans OpenCL

Chaque type Opencl fondamental a une version vectorielle. Vous pouvez utiliser le type de vecteur en ajoutant le nombre de composants souhaités après le type. Le nombre de composants pris en charge est 2,3,4,8 et 16. OpenCL 1.0 n'offre pas trois composants.

Vous pouvez initialiser n'importe quel vecteur de deux manières:

• Fournir un scalaire unique
• Satisfaire tous les composants

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

ou

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

ou toute autre combinaison de vecteurs satisfaisant au nombre de composants. Pour accéder aux éléments d'un vecteur, vous pouvez utiliser différentes méthodes. Vous pouvez soit utiliser l'indexation:

a[0] = 2;

ou utilisez des littéraux. L'avantage des littéraux est que vous pouvez les combiner comme bon vous semble, faites-le en un instant. Vous pouvez accéder à tous les composants vectoriels avec

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

vous pouvez également combiner plusieurs composants dans un nouveau vecteur

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

vous pouvez modifier l'ordre des composants comme vous le souhaitez.

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

mais vous ne pouvez pas faire quelque chose comme

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

Il existe des raccourcis pratiques pour accéder aux composants vectoriels. Les raccourcis suivants ne s'appliquent qu'aux tailles 2, 4, 8 et 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

Pour les tailles vectorielles 2,3 et 4, il existe des raccourcis supplémentaires

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

Noyau de correction gamma

Regardons un noyau de correction 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);
    }
}

Maintenant, parcourons ce code pas à pas. La première ligne crée une variable dans la région de la mémoire __constant du type sampler_t. Cet échantillonneur est utilisé pour spécifier davantage l'accès à nos données d'image. Veuillez vous reporter à la documentation de Khronos pour une documentation complète.

Nous avons attribué l'entrée à read_only et la sortie à write_only avant d'appeler notre noyau. Nous ajoutons donc ces modificateurs ici.

image2d et image3d sont toujours allouées sur la mémoire globale, nous pouvons donc omettre le modificateur __global ici. Notre valeur de gamma est située dans la mémoire __constant, nous le spécifions aussi.

Ensuite, nous obtenons notre identifiant de thread qui définit le pixel que nous allons corriger correctement. Nous interrogeons également la taille pour nous assurer que notre thread n’accède pas à la mémoire non allouée. Cela va définitivement bloquer votre noyau si vous l'oubliez.

Après nous être assuré que nous sommes un fil légitime, nous lisons notre pixel de notre image d'entrée. Nous le convertissons ensuite en float pour éviter la perte de décimales, effectuer des calculs, le reconvertir et l’écrire dans la sortie.