1 / 35

OpenCL

OpenCL. Programmation de cartes graphiques. OpenCL. API + langage basé sur le C99 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

kylene
Download Presentation

OpenCL

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. OpenCL Programmation de cartes graphiques

  2. OpenCL • API + langage basé sur le C99 • 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)

  3. Un périphérique vu par OpenCL

  4. Comparaison: C standard // Addition de 2 vecteurs de taille N // voidvecadd(int *C, int* A, int *B, int N) { for(int i = 0; i < N; i++) { C[i] = A[i] + B[i]; } }

  5. Comparaison: Programmation multithread voidvecadd(int *C, int* A, int *B, int N, int NP, inttid) { intept = N/NP; // nbre d’éléments par thread for(int i = tid*ept; i < (tid+1)*ept; i++) { C[i] = A[i] + B[i]; } }

  6. Comparaison: OpenCL __kernel voidvecadd(__global int *C, __global int* A, __global int *B) { inttid = get_global_id(0); // fonction OpenCL C[tid] = A[tid] + B[tid]; }

  7. Espace des indices • En OpenCL, 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 (NDRange) • Un indice local unique pour chaque work-item à l’intérieur d’un même workgroup.

  8. NDRange et Workgroups

  9. 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.

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

  11. Plateformes • Une plateforme est une implémentation de OpenCL spécifique à un manufacturier donné • Pour obtenir la liste des plateformes: cl_intclGetPlatformIDs(cl_uintnum_entries, cl_platform_id *platforms, cl_uint *num_platforms) • clGetPlatformInfo pour obtenir de l’information sur une plateforme donné

  12. Exemple 1 cl_intstatus; cl_uintnumPlatforms = 0; cl_platform_id *platforms = NULL; status = clGetPlatformIDs(0, NULL, &numPlatforms); platforms = (cl_platform_id*)malloc( numPlatforms*sizeof(cl_platform_id)); status = clGetPlatformIDs(numPlatforms, platforms, NULL); char Name[1000]; clGetPlatformInfo(platforms[0], CL_PLATFORM_NAME, sizeof(Name), Name, NULL); printf("Name of platform : %s\n", Name);

  13. Périphériques • Pour chacune des plateformes (ex. NVIDIA), on peut obtenir la liste des périphériques associés: cl_intclGetDeviceIDs(cl_platform_idplatform, cl_device_typedevice_type, cl_uintnum_entries, cl_device_id *devices, cl_uint *num_devices)

  14. Exemple 2 cl_uintnumDevices = 0; cl_device_id *devices = NULL; status = clGetDeviceIDs( platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices); devices = (cl_device_id*)malloc( numDevices*sizeof(cl_device_id)); status = clGetDeviceIDs( platforms[0], CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);

  15. Contextes • Un contexte est un type de conteneur permettant de communiquer des données et des instructions à des périphériques • On utilise la fonction: clCreateContext • 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.

  16. Contextes cl_context clCreateContext ( constcl_context_properties *properties, cl_uintnum_devices, constcl_device_id *devices, void (CL_CALLBACK *pfn_notify)( const char *errinfo, constvoid *private_info, size_t cb, void *user_data), void *user_data, cl_int *errcode_ret)

  17. Exemple 3 cl_contextcontext = NULL; context = clCreateContext( NULL, numDevices, devices, NULL, NULL, &status);

  18. File de commandes • On a déjà dit qu’un contexte est 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 une communication hôte-périphérique commencent par clEnqueue.

  19. Création d’une file de commandes cl_command_queue clCreateCommandQueue( cl_contextcontext, cl_device_iddevice, cl_command_queue_propertiesproperties, cl_int* errcode_ret) Note: Le paramètre « properties » sert, entre autres, à indiquer si les éléments seront pris dans l’ordre.

  20. Exemple 4 cl_command_queuecmdQueue; cmdQueue = clCreateCommandQueue( context, devices[0], 0, &status);

  21. 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.

  22. Création d’un tampon cl_mem clCreateBuffer( cl_contextcontext, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode_ret) Note: Le paramètre « flags » sert à indiquer si le tampon est en lecture, en écriture ou les deux.

  23. Écrire dans un tampon cl_int clEnqueueWriteBuffer ( cl_command_queuecommand_queue, cl_mem buffer, cl_boolblocking_write, // CL_TRUE pour appel bloquant size_t offset, size_t cb, constvoid *ptr, cl_uintnum_events_in_wait_list, constcl_event *event_wait_list, cl_event *event)

  24. Lire dans un tampon cl_int clEnqueueReadBuffer( cl_command_queuecommand_queue,   cl_mem buffer,   cl_boolblocking_read,   size_t offset,   size_t size,   void *ptr,   cl_uintnum_events_in_wait_list,   const cl_event *event_wait_list,   cl_event *event)

  25. Exemple 5 cl_membufferA, bufferB, bufferC; bufferA = clCreateBuffer( context, CL_MEM_READ_ONLY, datasize, NULL, &status); bufferB = clCreateBuffer( context, CL_MEM_READ_ONLY, datasize, NULL, &status); bufferC = clCreateBuffer( context, CL_MEM_WRITE_ONLY, datasize, NULL, &status); status = clEnqueueWriteBuffer( cmdQueue, bufferA, CL_FALSE, 0, datasize, A, 0, NULL, NULL); status = clEnqueueWriteBuffer( cmdQueue, bufferB, CL_FALSE, 0, datasize, B, 0, NULL, NULL);

  26. Kernels • Il faut distinguer le programme C de l’ordinateur hôte du programme C OpenCL qui sera exécuté sur un périphérique. • Un programme C OpenCL 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.

  27. Création et exécution d’un kernel • 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 clCreateProgramWithSource() • Un objet programme est compilé à l’aide de la commande clBuildProgram() • Un kernel est obtenu à l’aide de la fonction • Les paramètres du kernel sont affectés à l’aide de la commande clSetKernelArg() • Le kernel est finalement exécuté à l’aide de la commande clEnqueueNDRangeKernel()

  28. Création d’un objet programme cl_program clCreateProgramWithSource ( cl_contextcontext,   cl_uint count,   const char **strings,   const size_t *lengths,   cl_int *errcode_ret)

  29. Compiler un objet programme cl_int clBuildProgram ( cl_program program, cl_uintnum_devices, constcl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data), void *user_data)

  30. Obtenir un kernel cl_kernel clCreateKernel( cl_program program,   constchar *kernel_name,   cl_int *errcode_ret) Note: *kernel_namecontient le nom d’une fonction dans le programme source .

  31. Affecter les paramètres cl_int clSetKernelArg ( cl_kernelkernel,   cl_uintarg_index,   size_targ_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 clSetKernelArg()

  32. Exécuter un kernel cl_int clEnqueueNDRangeKernel( cl_command_queuecommand_queue,   cl_kernelkernel,   cl_uintwork_dim,   const size_t *global_work_offset,   const size_t *global_work_size,   const size_t *local_work_size,   cl_uintnum_events_in_wait_list,   const cl_event *event_wait_list,   cl_event *event)

  33. Organisation de la mémoire

  34. Espaces d’adressage Les qualificat • __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.

  35. 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.

More Related