//+------------------------------------------------------------------+
//| Noyau OpenCL |
//+------------------------------------------------------------------+
const string
cl_src=
//--- par défaut certains GPU ne supportent pas les doubles
//--- la directive cl_khr_fp64 est utilisée pour permettre un travail avec des doubles
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable \r\n"
//--- fonction du noyau OpenCL
"__kernel void Test_GPU(__global double *data, \r\n"
" const int N, \r\n"
" const int total_arrays) \r\n"
" { \r\n"
" uint kernel_index=get_global_id(0); \r\n"
" if (kernel_index>total_arrays) return; \r\n"
" uint local_start_offset=kernel_index*N; \r\n"
" for(int i=0; i<N; i++) \r\n"
" { \r\n"
" data[i+local_start_offset] *= 2.0; \r\n"
" } \r\n"
" } \r\n";
//+------------------------------------------------------------------+
//| Test_CPU |
//+------------------------------------------------------------------+
bool Test_CPU(double &data[],const int N,const int id,const int total_arrays)
{
//--- vérifie la taille du tableau
if(ArraySize(data)==0) return(false);
//--- vérifie l'index du tableau
if(id>total_arrays) return(false);
//--- calcule le décalage pour un tableau avec l'index id
int local_start_offset=id*N;
//--- multiplie les éléments par 2
for(int i=0; i<N; i++)
{
data[i+local_start_offset]*=2.0;
}
retourne vrai;
}
//---
#define ARRAY_SIZE 100 // taille du tableau
#define TOTAL_ARRAYS 5 // tableaux totaux
//--- Handles OpenCL
int cl_ctx; // Handle du contexte OpenCL
int cl_prg; // Handle du programme OpenCL
int cl_krn; // Handle du noyay OpenCL
int cl_mem; // Handle du buffer OpenCL
//---
double DataArray1[]; // tableau de données pour le calcul du CPU
double DataArray2[]; // tableau de données pour le calcul du GPU
//+------------------------------------------------------------------+
//| Fonction de démarrage du script |
//+------------------------------------------------------------------+
int OnStart()
{
//--- initialisation des objets OpenCL
//--- création du contexte OpenCL
if((cl_ctx=CLContextCreate())==INVALID_HANDLE)
{
Print("OpenCL not found. Erreur=",GetLastError());
return(1);
}
//--- création du programme OpenCL
if((cl_prg=CLProgramCreate(cl_ctx,cl_src))==INVALID_HANDLE)
{
CLContextFree(cl_ctx);
Print("Le programme OpenCL crée a échoué. Erreur=",GetLastError());
return(1);
}
//--- créer le noyau OpenCL
if((cl_krn=CLKernelCreate(cl_prg,"Test_GPU"))==INVALID_HANDLE)
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("Le noyau OpenCL crée a échoué. Erreur=",GetLastError());
return(1);
}
//--- créer un OpenCL buffer
if((cl_mem=CLBufferCreate(cl_ctx,ARRAY_SIZE*TOTAL_ARRAYS*sizeof(double),CL_MEM_READ_WRITE))==INVALID_HANDLE)
{
CLKernelFree(cl_krn);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("L'OpenCL buffer crée a échoué. Erreur=",GetLastError());
return(1);
}
//--- définir les paramêtres du noyau OpenCL
CLSetKernelArgMem(cl_krn,0,cl_mem);
CLSetKernelArg(cl_krn,1,ARRAY_SIZE);
CLSetKernelArg(cl_krn,2,TOTAL_ARRAYS);
//--- préparer des tableaux de données
ArrayResize(DataArray1,ARRAY_SIZE*TOTAL_ARRAYS);
ArrayResize(DataArray2,ARRAY_SIZE*TOTAL_ARRAYS);
//--- compléter des tableaux de données
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- calcule le décalage pour un tableau journalier
uint local_offset=j*ARRAY_SIZE;
//--- préparer le tableau avec index j
for(int i=0; i<ARRAY_SIZE; i++)
{
//--- compléter les tableaux avec la fonction MathCos(i+j);
DataArray1[i+local_offset]=MathCos(i+j);
DataArray2[i+local_offset]=MathCos(i+j);
}
};
//--- calcul du test CPU
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- calcul du tableau avec l'indice j
Test_CPU(DataArray1,ARRAY_SIZE,j,TOTAL_ARRAYS);
}
//--- préparer les paramètres CLExecute
uint offset[]={0};
//--- taille globale du travail
uint work[]={TOTAL_ARRAYS};
//--- écrit des données sur l'OpenCL buffer
CLBufferWrite(cl_mem,DataArray2);
//---exécuter le noyau OpenCL
CLExecute(cl_krn,1,offset,work);
//--- lit des données sur l'OpenCL buffer
CLBufferRead(cl_mem,DataArray2);
//--- erreur totale
double total_error=0;
//--- comparer les résultats et calculer l'erreur
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- calcule le décalage pour un tableau journalier
uint local_offset=j*ARRAY_SIZE;
//--- comparer les résultats
for(int i=0; i<ARRAY_SIZE; i++)
{
double v1=DataArray1[i+local_offset];
double v2=DataArray2[i+local_offset];
double delta=MathAbs(v2-v1);
total_error+=delta;
//--- montrer les premiers et les derniers tableaux
if((j==0) || (j==TOTAL_ARRAYS-1))
PrintFormat("array %d of %d, element [%d]: %f, %f, [error]=%f",j+1,TOTAL_ARRAYS,i,v1,v2,delta);
}
}
PrintFormat("Total error: %f",total_error);
//--- supprimer des objets OpenCL
//--- OpenCL buffer gratuit
CLBufferFree(cl_mem);
//--- sans noyau OpenCL
CLKernelFree(cl_krn);
//--- sans programme OpenCL
CLProgramFree(cl_prg);
//--- sans contexte OpenCL
CLContextFree(cl_ctx);
//---
retourne(0);
}
|