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. 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 [suite de l’article…].

En construction

#include<stdio.h>
#include<CL/cl.h>
#include <stdlib.h>
#include <string.h>

char *loadfile(const char *filename)
{
long   len ;
FILE * file  ;
char * buffer ;

file = fopen(filename, »r ») ;
if (file==NULL) {
fprintf(stderr, »*** failed to open file ‘%s’\n »,filename) ;
return NULL ;
}

/* Get the file length */
fseek(file,0L,SEEK_END) ;
len = ftell(file) ;

/* 640K should be enough for everyone */
if (len<=0 || len>640*1024) {
fprintf(stderr, »*** unexpected file length %ld for ‘%s’ \n »,len, filename) ;
}

/* allocate and load the file */
buffer = (char*) malloc(len+1);
fseek(file,0L,SEEK_SET) ;
fread(buffer,len,1,file) ;
buffer[len] = ‘\0’ ;

return buffer ;
}

void erreur(char *s)
{
printf(« Erreur %s\n »,s);
exit(0);
}

int main(int argc, char *argv[])
{

unsigned char j;
unsigned int taille_tab;

// Declarations propres a OpenCL

cl_platform_id id_plateforme ;
cl_uint nb_plateformes, nb_devices ;
cl_int bprogram, kernelarg, oclErr ;
cl_device_id id_device[2] ;
cl_context context ;
cl_command_queue queue ;
cl_program program ;
cl_kernel kernel ;
cl_mem cltableau, clresultat ;
char *source ;

unsigned int i ;
int *tableau ;
unsigned long *resultat;
size_t cltaille, cltaille2 ;

taille_tab=atoi(argv[1]);

tableau =(int *)malloc(taille_tab*sizeof(int));
resultat=(unsigned long *)malloc(taille_tab*sizeof(unsigned long));

cltaille=taille_tab*sizeof(cl_int);
cltaille2=taille_tab*sizeof(cl_ulong);

if (argc < 2)
erreur(« Usage: ./a.out taille »);

memset(resultat,0,taille_tab*sizeof(unsigned long));

for (i = 0; i < taille_tab; i++)
tableau[i] = i;

if (clGetPlatformIDs(1,&id_plateforme,&nb_plateformes) != CL_SUCCESS)
erreur(« clGetPlatformIDs »);

clGetDeviceIDs(id_plateforme,CL_DEVICE_TYPE_GPU,0,NULL,&nb_devices);
if (clGetDeviceIDs(id_plateforme,CL_DEVICE_TYPE_GPU,nb_devices,id_device,NULL) != CL_SUCCESS)
erreur(« clGetDeviceIDs »);

if ((context = clCreateContext(NULL,1,&(id_device[0]),NULL,NULL,NULL)) == NULL)
erreur(« clCreateContext »);

if ((queue = clCreateCommandQueue(context,id_device[0],CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,NULL)) == NULL)
erreur(« clCreateCommandQueue »);

source=loadfile(« plusun.cl »);

if ((program = clCreateProgramWithSource(context, 1, (const char **)&source, NULL, NULL)) == NULL)
erreur(« clCreateProgramWithSource »);

if ((bprogram = clBuildProgram(program,0,NULL,NULL,NULL,NULL)) != CL_SUCCESS)
erreur(« clBuildProgram »);

if ((kernel = clCreateKernel(program, »plusun_kernel »,NULL)) == NULL)
erreur(« clCreateKernel »);

if ((cltableau = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,cltaille,tableau,NULL)) == NULL)
erreur(« clCreateBuffer1 »);

if ((clresultat = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,cltaille2,resultat,NULL)) == NULL)
erreur(« clCreateBuffer2 »);

if ((kernelarg = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&cltableau)) != CL_SUCCESS)
erreur(« clSetKernelArg0 »);

if ((kernelarg = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&clresultat)) != CL_SUCCESS)
erreur(« clSetKernelArg1 »);

if ((kernelarg = clSetKernelArg(kernel, 2, sizeof(cl_int), NULL)) != CL_SUCCESS)
erreur(« clSetKernelArg2 »);

const size_t global_item_size = taille_tab ;
const size_t local_item_size = 1 ;

if ((oclErr=clEnqueueNDRangeKernel(queue,kernel,1,NULL,&global_item_size,&local_item_size,0,NULL,NULL)) !=CL_SUCCESS)
erreur(« clEnqueueNDRangeKernel »);

if ((oclErr=clEnqueueReadBuffer(queue,clresultat,CL_TRUE,0,cltaille2,resultat,0,NULL,NULL)) !=CL_SUCCESS)
erreur(« clEnqueueReadBuffer »);

for (i = 0; i < taille_tab ; i++)
printf(« %d %lu\n »,tableau[i],resultat[i]);
printf(« fin\n »);

// Nettoyage

oclErr = clFlush(queue);
oclErr = clFinish(queue);
oclErr = clReleaseKernel(kernel);
oclErr = clReleaseProgram(program);
oclErr = clReleaseMemObject(cltableau);
oclErr = clReleaseMemObject(clresultat);
oclErr = clReleaseCommandQueue(queue);
oclErr = clReleaseContext(context);
free(tableau);
free(resultat);
return 0;
}

Le kernel OpenCL associé :

__kernel void plusun_kernel(__global const int *a, __global unsigned long *res, __local int *tempo)
{
int gid = get_global_id(0);
*tempo = a[gid] ;
*tempo = *tempo + 1 ;
res[gid] = *tempo ;
}