OpenCL

logiciel informatique
OpenCL
Description de l'image OpenCL logo.svg.
Description de l'image Clinfo screenshot.png.
Informations
Créateur AppleVoir et modifier les données sur Wikidata
Développé par Khronos GroupVoir et modifier les données sur Wikidata
Première version Voir et modifier les données sur Wikidata
Dernière version 3.0.15 ()[1]Voir et modifier les données sur Wikidata
Écrit en C++ et CVoir et modifier les données sur Wikidata
Système d'exploitation Microsoft Windows, macOS, Linux, FreeBSD et AndroidVoir et modifier les données sur Wikidata
Environnement X86-64, IA-32 et architecture ARMVoir et modifier les données sur Wikidata
Type Framework
Langage de programmationVoir et modifier les données sur Wikidata
Licence Licence open-source approuvée par l'OSI (d)Voir et modifier les données sur Wikidata
Site web www.khronos.org/openclVoir et modifier les données sur Wikidata

OpenCL (Open Computing Language) est la combinaison d'une API et d'un langage de programmation dérivé du C, proposé comme un standard ouvert par le Khronos Group. OpenCL est conçu pour programmer des systèmes parallèles hétérogènes comprenant par exemple à la fois un CPU multi-cœur et un GPU. OpenCL propose donc un modèle de programmation se situant à l'intersection naissante entre le monde des CPU et des GPU, les premiers étant de plus en plus parallèles, les seconds étant de plus en plus programmables.

Présentation modifier

OpenCL distingue le processeur hôte (processeur central faisant office de chef d'orchestre) des périphériques (CPU, GPU, ou autre) dont la mission est d'exécuter des noyaux de calcul intensifs. OpenCL distingue donc d'un côté l'application tournant sur le processeur hôte (et qui va appeler l'API OpenCL), et de l'autre côté les noyaux qui sont programmés en OpenCL-C (et dont la vocation est d'être exécuté sur les périphériques).

OpenCL permet d'exprimer aisément du parallélisme de tâches mais aussi du parallélisme de données sous deux formes ; SPMD (Single Program Multiple Data) et SIMD (Single Instruction Multiple Data), le tout de manière hiérarchique. Un graphe de tâches peut être créé dynamiquement via l'API OpenCL. Chaque tâche peut être représentée soit sous forme d'une instance unique (appelée tâche), soit sous forme d'une collection d'instances (appelée NDRange) d'un même noyau. Les NDRanges peuvent être de 1, 2 ou 3 dimensions. Chaque instance de kernel appartenant à un NDRange est appelée work-item. Le NDrange peut lui-même être structuré en work-groups, ce qui permet aux work-items à l’intérieur des work-groups de partager des données et de se synchroniser via des barrières.

Si parmi certains de ses objectifs techniques, OpenCL semble se rapprocher de C pour CUDA, modèle de programmation propriétaire de la société Nvidia, OpenCL a des objectifs plus larges car n'étant pas uniquement dédié aux GPU. Dans le monde du calcul haute performance ou du jeu, OpenCL permettra de tirer parti de la puissance des processeurs graphiques, des CPU multi-cœurs ou d'autres systèmes de calcul intensifs tels le CELL d'IBM, qui équipe notamment la PlayStation 3 de Sony. Dans le monde des systèmes embarqués sur puce (SoC), tels qu'on les trouve dans les smartphones, OpenCL permettra l'accès, via une infrastructure de programmation unique, au processeur central, ainsi qu'aux différents sous-systèmes multimédia embarqués (GPU, DSP, computing array ou autres).

Il existe actuellement peu de moyens de déboguer des noyaux OpenCL. Tout d'abord, le débogueur NVIDIA Parallel Nsight, capable de déboguer CUDA thread par thread, ne supporte actuellement pas OpenCL, mais permet seulement de traquer les appels à l'API[2]. Ensuite, AMD propose une extension permettant de mettre des traces directement dans le code OpenCL (cl_amd_printf). Enfin, un programme appelé gDebugger (par Gremedy puis par AMD[3]) permet de suivre le déroulement de l'algorithme thread par thread. gDebugger a évolué, se nomme maintenant CodeXL et est hébergé[4] par l'initiative GPUOPEN.

OpenCL est également utilisé comme langage de description de matériel, pour la programmation des FPGA[5],[6].

Historique modifier

OpenCL a été initialement conçu par Apple[7] (qui en a déposé la marque), et affiné dans le cadre d'une collaboration avec AMD, Intel et Nvidia. Apple soumet d'abord sa proposition initiale au Khronos Group. Le , le Khronos Compute Working Group est formé, comprenant des représentants des fabricants de matériel informatique et de logiciels. Celui-ci travaille durant cinq mois à boucler les détails techniques de la spécification OpenCL 1.0. La spécification est révisée par les membres de Khronos et approuvée pour une version d'essai le . Une nouvelle version, OpenCL 1.1, est publiée en par le Khronos Group. OpenCL 1.1 clarifie certains aspects de la spécification précédente et apporte de nouvelles fonctionnalités telles que les sous-buffers, les vecteurs à 3 éléments, les événements utilisateur, de nouvelles fonctions builtin, le support en standard d'extensions optionnelles 1.0 (telles que les fonctions atomiques 32 bits).

OpenCL est intégré dans Mac OS X 10.6[7] (Snow Leopard). AMD décide de supporter OpenCL et DirectX 11 plutôt que Close to Metal dans son framework Stream SDK. RapidMind (en) annonce l'adoption de OpenCL sous sa plate-forme de développement, afin de supporter les processeurs graphiques de plusieurs fabricants avec une interface unique. Nvidia confirme également le le support complet de la spécification 1.0 dans son GPU Computing Toolkit[8].

Le , le Khronos Group a publié les spécifications d'OpenCL 1.2. On y trouve notamment des fonctionnalités liées à la mobilité et à la portabilité, avec par exemple la possibilité de dissocier compilation et édition de liens des noyaux[9].

WebCL modifier

Le Khronos Group a également développé une intégration d'OpenCL, bibliothèque de calcul parallèle, dans l'ensemble des interfaces de programmation d'HTML5[10]. Actuellement, les navigateurs utilisent des extensions pour gérer OpenCL.

Nokia et Mozilla [11] ont développé des extensions pour Firefox. Samsung pour WebKit et Motorola pour Node.js.

Historique des implémentations modifier

Des implémentations d'OpenCL existent pour la majorité des plateformes aujourd'hui. IBM pour ses supercalculateurs sous GNU/Linux utilisant des processeurs Power, les processeurs X86 d'Intel et AMD et les GPU les accompagnant traditionnellement (ATI, nVidia, VIA), les processeurs ARM Cortex-A9 (parties SSE et fpu 128bits Neon), ainsi que les DSP, GPU et autres computing array les accompagnant dans les nombreuses implémentations des System on chip (SoC) (nVidia Tegra2, Qualcomm Snapdragon, Apple A4, Marvell Armadaetc.). Mesa (Implémentation OpenGL/OpenVG sous GNU/Linux) contient un state-tracker OpenCL pour Gallium3D en cours de développement, nommé Clover  [12],[13]

Le , AMD et Nvidia font la première démonstration publique d'OpenCL, une présentation de 75 minutes à SIGGRAPH Asia 2008. AMD effectue une démonstration d'OpenCL accélérée sur CPU et explique la scalabilité d'OpenCL sur un ou plusieurs cœurs tandis qu'Nvidia fait une démonstration accélérée par GPU[14],[15].

Le , à la GDC 2009, AMD et Havok font une démonstration de la première implémentation accélérée par OpenCL, Havok Cloth sur un GPU de la série Radeon HD 4000 d'AMD[16].

Le , Nvidia annonce la sortie de son pilote OpenCL et du SDK aux développeurs participant à son OpenCL Early Access Program[17].

Le , AMD révèle les premiers outils de développement pour sa plateforme OpenCL comme partie de son programme ATI Stream SDK v2.0 Beta[18].

Le , Apple sort Mac OS X Snow Leopard, qui contient une implémentation complète d'OpenCL[19].

Dans Snow Leopard, OpenCL est initialement supporté sur les puces ATI Radeon HD 4850, ATI Radeon HD 4870 ainsi que les puces Nvidia Geforce 8600M GT, GeForce 8800 GS, GeForce 8800 GT, GeForce 8800 GTS, Geforce 9400M, GeForce 9600M GT, GeForce GT 120, GeForce GT 130, GeForce GTX 285, Quadro FX 4800, et Quadro FX 5600[20].

Le , Nvidia sort ses propres pilotes OpenCL et son implémentation du SDK.

Le , AMD sort la quatrième bêta du ATI Stream SDK 2.0, qui fournit une implémentation OpenCL complète sur tous les GPU des familles R700/R800, utilisant également les unités SSE3 des CPUs. Le SDK est disponible à la fois pour GNU/Linux et Windows[21].

Le , IBM sort la version 0.1 de son SDK OpenCL pour GNU/Linux sur l'architecture Power utilisé dans une majorité des plus puissants supercalculateurs au monde[22].

Le , Nvidia sort des pilotes pour l'implémentation OpenCL 1.0 (rev 48).

Les implémentations OpenCL d'Apple[23], Nvidia[24], RapidMind (en)[25] et Mesa Gallium3D[26] sont toutes basées sur la technologie de compilation LLVM et utilisent le compilateur Clang comme frontend.

Le , VIA sort son premier produit supportant OpenCL 1.0 - Le processeur vidéo ChromotionHD 2.0 inclus dans les puces VN1000[27].

Le , AMD sort la version de production de l’ATI Stream SDK 2.0[28], qui fournit un support d'OpenCL 1.0 pour les GPU R800 et un support bêta pour R700.

Le , Intel sort la version finale de son kit de développement supportant OpenCL version 1.1[29].

Le , AMD annonce son kit de développement ATI Stream SDK 2.5[30], qui améliore, entre autres, la bande passante CPU/GPU pour tirer un meilleur parti de ses récents APU.

Chez Intel, les processeurs graphiques inclus dans sa gamme de processeurs Ivy Bridge, sortis le , Intel HD 2500 ainsi que HD 4000 et supérieurs, sont les premières architectures à supporter matériellement OpenCL, en version 1.1[31].

La bibliothèque Beignet[32] est une bibliothèque OpenCL libre pour les processeurs Intel Ivy Bridge GT2, dans le cadre du projet freedesktop.org, et développé principalement par Intel. Il utilise principalement LLVM, mais est également compatible avec GCC.

Fin 2013, ARM annonce à son tour son « Mali OpenCL SDK », pour ses processeurs graphiques Mali T600 et supérieurs, dont les premiers modèles sont sortis en 2012. Les sources sont disponibles, mais la licence est propriétaire et contraignante quant à la redistribution[33].

Implémentations Open-source modifier

Clover [1] et libCLC [2] pour les fonctions de la bibliothèque standard OpenCL
POCL (Portable OpenCL) [3]

Exemple modifier

Cet exemple calcule une Transformée de Fourier rapide :

/* creation d'un contexte de calcul sur GPU */
context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);

/* récupération de la liste des cartes disponibles */
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &nb_devices);
clGetContextInfo(context, CL_CONTEXT_DEVICES, nb_devices, devices, NULL);

/* creation d'une queue de commande sur le premier GPU */ 
queue = clCreateCommandQueue(context, devices[0], 0, NULL);

/* allocation des tampons mémoire */
memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA, NULL);
memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL, NULL);

/* création du programme de calcul (le programme qui s'execute sur le GPU) */ 
program = clCreateProgramWithSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL);

/* compilation du programme */
clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

/* création du noyau de calcul */
kernel = clCreateKernel(program, "fft1D_1024", NULL);

/* mise en place des paramètres */
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs[0]);
clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs[1]);
clSetKernelArg(kernel, 2, sizeof(float)*(local_work_size[0]+1)*16, NULL);
clSetKernelArg(kernel, 3, sizeof(float)*(local_work_size[0]+1)*16, NULL);

/* création des objets de travail
   et lancement du calcul */
global_work_size[0] = num_entries;
local_work_size[0] = 64;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);

Le calcul : (basé sur Fitting FFT onto the G80 Architecture)[34]

// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into 
// calls to a radix 16 function, another radix 16 function and then a radix 4 function 

__kernel void fft1D_1024 (__global float2 *in, __global float2 *out, 
                          __local float *sMemx, __local float *sMemy) { 
  int tid = get_local_id(0); 
  int blockIdx = get_group_id(0) * 1024 + tid; 
  float2 data[16]; 
  // starting index of data to/from global memory 
  in = in + blockIdx;  out = out + blockIdx; 
  globalLoads(data, in, 64); // coalesced global reads 
  fftRadix16Pass(data);      // in-place radix-16 pass 
  twiddleFactorMul(data, tid, 1024, 0); 
  // local shuffle using local memory 
  localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4))); 
  fftRadix16Pass(data);               // in-place radix-16 pass 
  twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication 
  localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15))); 
  // four radix-4 function calls 
  fftRadix4Pass(data); fftRadix4Pass(data + 4); 
  fftRadix4Pass(data + 8); fftRadix4Pass(data + 12); 
  // coalesced global writes 
  globalStores(data, out, 64); 
}

Références modifier

  1. « The OpenCL Specification »
  2. (en) « Nvidia Parallel Nsight » (consulté le ).
  3. (en-US) « AOCL Archive », sur AMD (consulté le ).
  4. (en) « CodeXL has moved to GPUOPEN! » (consulté le ).
  5. (en) Kevin Morris, « HLS versus OpenCL », sur Electronic-Engineering Journal,
  6. Roozmeh et Lavagno 2018.
  7. a et b (en) « Apple Previews Mac OS X Snow Leopard to Developers », Apple, .
  8. Communiqué de presse de Nvidia.
  9. (en) « Khronos Releases OpenCL 1.2 Specification », sur The Khronos Group, (consulté le ).
  10. (en) « Khronos Releases OpenCL 1.2 Specification », sur The Khronos Group, (consulté le )
  11. http://webcl.nokiaresearch.com/
  12. (en) « Clover git repo », (consulté le ).
  13. (en) « OpenCL Over Mesa, Gallium3D Discussion », (consulté le ).
  14. (en) « OpenCL Demo, AMD CPU », (consulté le ).
  15. (en) « OpenCL Demo, NVIDIA GPU », (consulté le ).
  16. (en) « AMD and Havok demo OpenCL accelerated physics », PC Perspective, (consulté le ).
  17. (en) « NVIDIA Releases OpenCL Driver To Developers », NVIDIA, (consulté le ).
  18. (en) « AMD does reverse GPGPU, announces OpenCL SDK for x86 », Ars Technica, (consulté le ).
  19. (en) Dan Moren, Jason Snell, « Live Update: WWDC 2009 Keynote », macworld.com, MacWorld, (consulté le ).
  20. (en) « Mac OS X Snow Leopard – Technical specifications and system requirements », Apple Inc, (consulté le ).
  21. (en) « ATI Stream Software Development Kit (SDK) v2.0 Beta Program » (consulté le ).
  22. (en) « OpenCL Development Kit for Linux on Power » (consulté le ).
  23. (en) « Apple entry on LLVM Users page » (consulté le ).
  24. (en) « Nvidia entry on LLVM Users page » (consulté le ).
  25. (en) « Rapidmind entry on LLVM Users page » (consulté le ).
  26. (en) « Zack Rusin's blog post about the Mesa Gallium3D OpenCL implementation » (consulté le ).
  27. (en) « WebCL - Heterogeneous parallel computing in HTML5 web browsers », sur The Khronos Group, (consulté le ).
  28. (en) « ATI Stream SDK v2.0 with OpenCL 1.0 Support » (consulté le ).
  29. (en) « Intel® SDK for OpenCL™ Applications », sur Intel (consulté le ).
  30. (en) « AMD APP SDK 2.5 provides enhanced performance and major new capabilities » (consulté le ).
  31. « Nouveaux pilotes chez Intel : OpenGL 4.0 et Ultra HD pour Ivy Bridge », (consulté le ).
  32. « Beignet » (consulté le ).
  33. (en) « Mali OpenCL SDK » [archive du ] (consulté le ).
  34. (en) « Fitting FFT onto G80 Architecture », Vasily Volkov and Brian Kazian, UC Berkeley CS258 project report, (consulté le )

Bibliographie modifier

  • Mehdi Roozmeh et Luciano Lavagno, « Design space exploration of multi-core RTL via high level synthesis from OpenCL models », Microprocessors and Microsystems, vol. 63,‎ , p. 199-208 (DOI 10.1016/j.micpro.2018.09.009, lire en ligne)

Voir aussi modifier