OpenCL

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-coeur 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. OpenCL distingue le processeur hôte (processeur central faisant office de chef d'orchestre) des devices (CPU, GPU, ou autre) dont la mission est d'exécuter des noyaux de calcul intensifs. OpenCL distingue donc l'application (écrite en C) tournant sur le processeur hôte et qui va appeler l'API OpenCL, des kernels qui sont programmés en OpenCL-C et dont la vocation est d'être exécuté sur les devices. OpenCL permet d'exprimer du parallélisme de tâche et du parallélisme de donnée (SPMD - Single Program Multiple Data) de manière hiérarchique. Un graphe de tache peut être crée dynamiquement via l'API OpenCL. Chaque tache peut être soit une unique instance d'un kernel (appelée task), soit de multiples instances d'un même kernel (appelé NDRange). 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'interieur 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 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'au différents sous-systèmes multimédia embarqués (GPU, DSP, computing array ou autres).

OpenCL a été initialement conçu par Apple (qui en possède les droits d'auteur), et affiné dans le cadre d'une collaboration avec AMD, Intel et Nvidia. Apple soumet d'abord sa proposition initiale au Khronos Group. Le 16 juin 2008, 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 8 décembre. Une nouvelle version, OpenCL 1.1, est publiées en juin 2010 par le Khronos Group. OpenCL 1.1 clarifie certain aspects de la spécification précédente et apporte de nouvelles fonctionnalités tels que les sous-buffers, les vecteurs à 3 éléments, les événements utilisateur, de nouvelles fonctions builtin, le support en standard d'extensions optionnelle 1.0 (tel que les fonctions atomiques 32 bits) .

OpenCL est intégré dans Mac OS X 10.6 (Snow Leopard). AMD décide de supporter OpenCL et DirectX 11 plutôt que Close to Metal dans son framework Stream SDK. RapidMind 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 9 décembre 2008 le support complet de la spécification 1.0 dans son GPU Computing Toolkit.

Des implémentations d'OpenCL existent pour la majorité des plateformes aujourd'hui. IBM pour ses supercalculateurs Linux utilisant des processeur 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 Armada, etc...). Mesa (Implémentation OpenGL/OpenVG de Linux) contient un state-tracker OpenCL pour Gallium3D en cours de développement, nommé Clover

Le 10 décembre 2008, 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émonstratrion 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.

Le 26 mars 2009, à 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.

Le 20 avril 2009, Nvidia annonce la sortie de son pilote OpenCL et du SDK aux développeurs participant à son OpenCL Early Access Program.

Le 5 août 2009, AMD révèle les premiers outils de développement pour sa plateforme OpenCL comme partie de son programmme ATI Stream SDK v2.0 Beta.

Le 28 août 2009, Apple sort Mac OS X Snow Leopard, qui contient une implémentation complète d'OpenCL.

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.

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

Le 13 octobre 2009, AMD sort la quatrième bêta du ATI Stream SDK 2.0, qui fourni 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 Linux et Windows.

Le 30 octobre 2009, IBM sort la version 0.1 de son SDK OpenCL pour Linux sur l'architecture Power utilisé dans une majorité des plus puissants supercalculateurs au monde.

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

Les implémentations OpenCL d'Apple, Nvidia, RapidMind et Mesa Gallium3D sont toutes basées sur la technologie de compilation LLVM et utilisent le compilateur Clang comme frontend.

Le 10 décembre 2009, VIA sort son premier produit supportant OpenCL 1.0 - Le processeur vidéo ChromotionHD 2.0 inclus dans les puces VN1000.

Le 21 décembre 2009, AMD sort la version de production de l'ATI Stream SDK 2.0, qui fournit un support d'OpenCL 1.0 pour les R800 GPUs et un support bêta pour R700.

Cet exemple calcule une Transformée de Fourier rapide

// create a compute context with GPU device context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL); // list the available devices clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &nb_devices); clGetContextInfo(context, CL_CONTEXT_DEVICES, nb_devices, devices, NULL); // create a command queue on the first GPU queue = clCreateCommandQueue(context, devices[0], 0, NULL); // allocate the buffer memory objects 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); // create the compute program program = clCreateProgramWithSource(context, 1, &fft1D_1024_kernel_src, NULL, NULL); // build the compute program executable clBuildProgram(program, 0, NULL, NULL, NULL, NULL); // create the compute kernel kernel = clCreateKernel(program, "fft1D_1024", NULL); // set the args values 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); // create N-D range object with work-item dimensions and execute kernel 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)

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