Open CL Programmation de cartes graphiques Open CL

Open. CL Programmation de cartes graphiques

Open. CL �API + langage basé sur le C 99 �Conçu pour programmer des systèmes parallèles hétérogènes: CPU multi-coeur, GPU, etc. �On distingue le processeur hôte des périphériques �Un conteneur spécial (context) sert d’interface entre le processeur hôte et les périphériques �Les données et les tâches (kernels) sont transférés à l’aide d’une file d’attente (command queue)

Un périphérique vu par Open. CL

Comparaison: C standard // Addition de 2 vecteurs de taille n // La fonction s’occupe de toutes les cases void vecadd(int *C, int* A, int *B, int n) { for(int i = 0; i < n; i++) { C[i] = A[i] + B[i]; } }

Comparaison: Programmation multithread // La fonction s’occupe d’un groupe de cases void vecadd(int *C, int* A, int *B, int n) { int p = omp_get_num_threads(); int id = omp_get_thread_num(); int t = n/p; // nbre d’éléments par thread for(int i = id*t; i < (id+1)*t; i++) C[i] = A[i] + B[i]; }

Comparaison: Open. CL // La fonction s’occupe d’une seule case __kernel void vecadd(__global int *C, __global int* A, __global int *B) { int tid = get_global_id(0); // fonction Open. CL C[tid] = A[tid] + B[tid]; }

Espace des indices �En Open. CL, l’espace des indices des processeurs peut avoir 1, 2 ou 3 dimensions. �Il y a deux niveaux d’indices: �Un indice global unique pour chaque work-item du périphérique (ce domaine est appelé NDRange) �Un indice local unique pour chaque work-item à l’intérieur d’un même workgroup.

NDRange et Workgroups

Connaître son indice �get_global_id(dim): indice globale du work-item appelant selon la dimension dim=0, 1 ou 2 �get_local_id(dim): indice local du work-item appelant �get_group_id(dim): indice du workgroup auquel appartient le work-item appelant �get_local_size(dim): taille de la dimension dim dans le workgroup du work-item appelant.

Relation entre global et local get_global_id(dim) = get_local_size(dim)*get_group_id(dim) + get_local_id(dim)

Plateformes �Une plateforme est une implémentation de Open. CL spécifique à un manufacturier donné �Pour obtenir la liste des plateformes: cl_int cl. Get. Platform. IDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms) �num_entries: Taille du tableau plateformes �Platforms: Liste des plateformes retournés Si la valeur est NULL alors le paramètre est ignoré. �num_platforms: Nombre de plaetformes retournées. Ignoré si NULL.

Plateformes Pour obtenir de l’information sur une plateforme donné: cl_int cl. Get. Platform. Info( cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret ) Platform: The platform ID returned by cl. Get. Platform. IDs or can be NULL. If platform is NULL, the behavior is implementation-defined. param_name: An enumeration constant that identifies the platform information being queried. It can be one of the values specified in the table below. param_value: A pointer to memory location where appropriate values for a given param_name will be returned. Possible param_values returned are listed in the table below. If param_value is NULL, it is ignored. param_value_size: Specifies the size in bytes of memory pointed to by param_value. This size in bytes must be ≥ size of return type specified in the table below. param_value_size_ret: Returns the actual size in bytes of data being queried by param_value. If param_value_size_ret is NULL, it is ignored

Exemple 1 cl_int status; cl_uint num. Platforms = 0; cl_platform_id *platforms = NULL; status = cl. Get. Platform. IDs(0, NULL, &num. Platforms); platforms = (cl_platform_id*)malloc( num. Platforms*sizeof(cl_platform_id)); status = cl. Get. Platform. IDs(num. Platforms, platforms, NULL); char Name[1000]; cl. Get. Platform. Info(platforms[0], CL_PLATFORM_NAME, sizeof(Name), Name, NULL); printf("Name of platform : %sn", Name);

Périphériques �Pour chacune des plateformes (ex. NVIDIA), on peut obtenir la liste des périphériques associés: cl_int cl. Get. Device. IDs(cl_platform_id platform, cl_device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices)

Types de périphériques CL_DEVICE_TYPE_CPU: An Open. CL device that is the host processor. The host processor runs the Open. CL implementations and is a single or multi-core CPU. CL_DEVICE_TYPE_GPU: An Open. CL device that is a GPU. By this we mean that the device can also be used to accelerate a 3 D API such as Open. GL or Direct. X. CL_DEVICE_TYPE_ACCELERATOR: Dedicated Open. CL accelerators (for example the IBM CELL Blade). These devices communicate with the host processor using a peripheral interconnect such as PCIe. CL_DEVICE_TYPE_CUSTOM: Dedicated accelerators that do not support programs written in Open. CL C. CL_DEVICE_TYPE_DEFAULT: The default Open. CL device in the system. The default device cannot be a CL_DEVICE_TYPE_CUSTOM device. CL_DEVICE_TYPE_ALL: All Open. CL devices available in the system except CL_DEVICE_TYPE_CUSTOM devices.

Exemple 2 cl_uint num. Devices = 0; cl_device_id *devices = NULL; status = cl. Get. Device. IDs( platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &num. Devices); devices = (cl_device_id*)malloc( num. Devices*sizeof(cl_device_id)); status = cl. Get. Device. IDs( platforms[0], CL_DEVICE_TYPE_ALL, num. Devices, devices, NULL);

Contextes �Un contexte est un canal de communication entre l'hôte et des périphérique d'une même plateforme. On peut le voir comme un type de conteneur permettant de communiquer des données et des instructions à des périphériques �On utilise la fonction: cl. Create. Context �Le premier paramètre sert à limité la porté du contexte, par exemple en spécifiant une plateforme particulière. �Le « callback » sert a fournir une fonction qui pourra fournir d’avantage d’information sur les erreurs se produisant tout au long de l’utilisation du contexte.

Contextes cl_context cl. Create. Context ( const cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (CL_CALLBACK *pfn_notify)( const char *errinfo, const void *private_info, size_t cb, void *user_data), void *user_data, // correspond au dernier argument du callback cl_int *errcode_ret)

Exemple 3 cl_context = NULL; context = cl. Create. Context( NULL, num. Devices, devices, NULL, &status);

File de commandes �On a déjà dit qu’un contexte est comme un conteneur permettant de communiquer avec un périphérique. �On ajoute et retire l’information d’un contexte à l’aide d’une file d’attente appelée « file de commande » (command queue). �Toutes les commandes indiquant un transfert d'information hôte-périphérique commencent par cl. Enqueue.

Création d’une file de commandes cl_command_queue cl. Create. Command. Queue. With. Properties( cl_context, cl_device_id device, cl_command_queue_properties, cl_int* errcode_ret) Note: Le paramètre « properties » sert, entre autres, à indiquer si les éléments seront pris dans l’ordre.

Exemple 4 cl_command_queue cmd. Queue; cmd. Queue = cl. Create. Command. Queue. With Properties( context, devices[0], 0, &status);

Objets mémoire �Avant d’être transférés, les données doivent être converties dans un format particulier. �Il y a deux formats: � Les tampons (buffers) � Les images �Les tampons sont l’équivalent des tableaux en C et sont stockés de façon contiguë en mémoire. �Les images sont des objets opaques et sont placées en mémoire de façon à optimiser les performances.

Création d’un tampon cl_mem cl. Create. Buffer( cl_context, cl_mem_flags, size_t size, //taille du tampon void *host_ptr, // adresse du tampon cl_int *errcode_ret) Note: Le paramètre « flags » sert à indiquer si le tampon est en lecture, en écriture ou les deux.

Écrire dans un tampon cl_int cl. Enqueue. Write. Buffer ( cl_command_queue, cl_mem buffer, cl_bool blocking_write, // CL_TRUE pour appel bloquant size_t offset, //décalage dans le tampon size_t cb, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)

Lire dans un tampon cl_int cl. Enqueue. Read. Buffer( cl_command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, //décalage dans le tampon size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)

Exemple 5 cl_mem buffer. A, buffer. B, buffer. C; buffer. A = cl. Create. Buffer( context, CL_MEM_READ_ONLY, datasize, NULL, &status); buffer. B = cl. Create. Buffer( context, CL_MEM_READ_ONLY, datasize, NULL, &status); buffer. C = cl. Create. Buffer( context, CL_MEM_WRITE_ONLY, datasize, NULL, &status); status = cl. Enqueue. Write. Buffer( cmd. Queue, buffer. A, CL_FALSE, 0, datasize, A, 0, NULL); status = cl. Enqueue. Write. Buffer( cmd. Queue, buffer. B, CL_FALSE, 0, datasize, B, 0, NULL);

Kernels � Il faut distinguer le programme C de l’ordinateur hôte du programme C Open. CL qui sera exécuté sur un périphérique. � Un programme C Open. CL est une collection de fonctions appelées « kernels » � Le type de retour de ces fonctions doit être void � Les kernels représentent les tâches que l’on peut ordonnancer de façon dynamique sur les périphériques.

Création et exécution d’un kernel 1. 2. 3. 4. 5. 6. Le code source est stocké sous la forme d’une chaîne de caractères La chaine est convertie en objet programmes à l’aide de la commande cl. Create. Program. With. Source() Un objet programme est compilé à l’aide de la commande cl. Build. Program() Un kernel est extrait à l’aide de la fonction cl. Create. Kernel Les paramètres du kernel sont affectés à l’aide de la commande cl. Set. Kernel. Arg() Le kernel est finalement exécuté à l’aide de la commande cl. Enqueue. NDRange. Kernel()

Création d’un objet programme cl_program cl. Create. Program. With. Source ( cl_context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret)

Compiler un objet programme cl_int cl. Build. Program ( cl_program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program, void *user_data)

Obtenir un kernel cl_kernel cl. Create. Kernel ( cl_program, const char *kernel_name, cl_int *errcode_ret) Note 1: *kernel_name contient le nom d’une fonction dans le programme source. Note 2: Voir cl. Get. Program. Build. Info pour les erreurs de compilation

Affecter les paramètres cl_int cl. Set. Kernel. Arg ( cl_kernel, cl_uint arg_index, size_t arg_size, const void *arg_value) Note: Les paramètres demeurent affectés au kernel tant qu’il n’y a pas de modification explicite à l’aide de cl. Set. Kernel. Arg()

Exécuter un kernel cl_int cl. Enqueue. NDRange. Kernel( cl_command_queue, cl_kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)

Organisation de la mémoire

Espaces d’adressage Les qualificatifs: �__global: Mémoire partagée par tous les work-items du kernel. �__constant: Mémoire en lecture seule partagée par tous les work-items du kernel. �__local: Mémoire partagée par tous les work-items d’un même work-group. �__private: Mémoire privée pour chaque work-item.

Espace d’adressage par défaut �Pour améliorer la performance, on utilise autant que possible la mémoire __local ou __private. �S’il n’y a pas de qualificatif, les paramètres des fonctions et les variables locales sont __private � Le qualificatif des paramètres qui sont des pointeurs doit être spécifié et ne peut pas être __private. �Les mémoires __private et __local ne sont pas préservées à la fin de l’exécution d’un work-item. �La seule façon de retourner une valeur à l’hôte est via la mémoire __global.

Synchronisation 1) Opération mémoire bloquante Ex. cl_int cl. Enqueue. Read. Buffer( cl_command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event )

Synchronisation 2) Attente de la terminaison d'un événement Chaque commande cl. Enqueue produit un événement qui peut être dans l'un des états suivants: �Queued: La commande est placée dans une file d'attente. �Submitted: La commande est retiré de la file d'attente et soumise au périphérique pour être exécutée. �Ready: La commande est prête à être exécuté sur le périphérique. �Running: En cours d'exécution sur le périphérique. �Ended: L'exécution sur le périphérique est terminée. �Complete: La commande et tous ses enfants ont terminé.

Synchronisation On peut déterminer l'état d'un événement à l'aide de la commande cl. Get. Event. Info() en lui pâssant le paramètre CL_EVENT_COMMAND_EXECUTION_STATUS cl_int cl. Get. Event. Info( cl_event, cl_event_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret )

Synchronisation On peut attendre la fin d'un événement à l'aide de la commande: cl_int cl. Wait. For. Events ( cl_uint num_events, const cl_event *event_list )

Synchronisation 3) Une autre méthode de synchronisation consiste à mettre une barrière sur la file de commande ce qui ne bloque pas le programme hôte: cl_int cl. Enqueue. Barrier. With. Wait. List( cl_command_queue , cl_uint num_events_in_wait_list , const cl_event *event_wait_list , cl_event *event )

Synchronisation 4) On peut aussi vouloir que le périphérique attende après un événement usager: cl_event cl. Create. User. Event( cl_context, cl_int *errcode_ret )

Synchronisation L'état initial d'un événement usager est initialement CL_SUBMITTED. On peut le changer pour CL_COMPLETED. cl_int cl. Set. User. Event. Status( cl_event, cl_int execution_status )

Synchronisation 5) Barrière de synchronisation pour les work item d'un même workgroup. void work_group_barrier(cl_mem_fence_flags) Flag une combinaison de valeurs tels que: �CLK_LOCAL_MEM_FENCE �CLK_GLOBAL_MEM_FENCE

Tiré du livre "Heterogeneous Computing with Open. CL"

File de commandes non ordonnée Les files de commandes que nous avons vues jusqu'à présent étaient ordonnées: les commandes sont exécuté et complétées dans l'ordre où elle ont été mises dans la file. Cependant, cette façon de faire n'exploite pas toujours pleinement le parallélisme potentiel de du périphérique. Voici un exemple de création et d'utilisation d'une file de commande non ordonnée. http: //www. uqac. ca/flemieux/INF 856/Open. CL-Out-of-Order-Queue. txt