Discussion :

[ebalexandr]

Bonjour,

J'ai voulu tester l'API OpenCL, mais quelque fois je reçois l'erreur CL_OUT_OF_RESOURCES. Dans la documentation de la fonction clEnqueueNDRangeKernel il est écrit, qu'on obtient cette erreur, quand il y a un débordement global_work_size. J'ai utilisé la fonction clGetDeviceInfo pour calculer quel valeurs de global_work_size et local_work_size.

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:3
CL_DEVICE_ADDRESS_BITS:32
CL_DEVICE_MAX_WORK_GROUP_SIZE:512
CL_DEVICE_MAX_WORK_ITEM_SIZES: (512,512,64)

Je travail dans une seule dimension(work_dim=1), donc si je comprend bien les valeurs maximales que je peux attribuer sont 2^32-1 et 512 c'est à dire global_work_size=0xffffffff et local_work_size=0x200, mais il faut que global_work_size se divise par local_work_size, donc global_work_size=0xfffffe00;

Est-ce que je me trompe ?

merci

[gbdivers]

Bonjour ebalexandr

Le débordement est produit quand global_work_size est supérieur à sizeof(size_t) du GPU (donc à CL_DEVICE_ADDRESS_BITS) ou au nombre maximal de kernel

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

global_work_size se divise par local_work_size, donc global_work_size=0xfffffe00

Presque.
Le travail a effectué, de taille gobal_work_size, est divisé entre différents work-group. Chaque work-group contient plusieurs work-items qui exécutent le kernel (voir chap 3.2 des spécifications OpenCL). La limite pour le nombre d'items par groupe est CL_DEVICE_MAX_WORK_ITEM_SIZES et le nombre de groupe max est CL_DEVICE_MAX_WORK_GROUP_SIZE.
La limite pour le nombre d'item total est donc de CL_DEVICE_MAX_WORK_ITEM_SIZES * CL_DEVICE_MAX_WORK_GROUP_SIZE (de mémoire)

[ebalexandr]

Merci pour ta repense. Maintenant au moins je sais ce que je fait :-)

J'ai mit global_work_size[0]=0x88000 et local_work_size[0]=512, mais j'ai toujours cette erreur CL_OUT_OF_RESOURCES. C'est peut être mon kernel qui ne va pas ?

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
const char kernel_cl[] = {
"__kernel void testopencl(__global float *out)\n"
"{\n"
    "uint i=0xffffffff;\n"
    "float f=0.0;\n"
    "while(i)\n"
    "{\n"
        "i--;\n"
        "f=f+0.000001;\n"
    "}\n"
    "*out=f;\n"
"}"
};

Je n'ai plus d'erreur, que quand je met uint i = 0x20000000; mais le résultat est erroné (out).

[ebalexandr]

Le même code en C(exécuté sur le CPU) donne aussi des résultats erronés. Je pense c'est dû à "float".

Et c'est la fonction clEnqueueReadBuffer qui a généré l'erreur -5 (CL_OUT_OF_RESOURCES=-5 dans le fichier cl.h).
C'est étrange car dans la documentation en ligne http://www.khronos.org/registry/cl/s...eadBuffer.html il n'existe aucune description de CL_OUT_OF_RESOURCES.
C'est pour cette raison que j'ai pensé que l'erreur a été générée par la fonction clEnqueueNDRangeKernel http://www.khronos.org/registry/cl/s...ngeKernel.html .

[gbdivers]

Il y a quelques problèmes avec ton code...

Dans ton kernel :
- tu ne fais qu'une boucle : ce n'est pas du calcul parallèle
- tu as des problème de conflit d'accès mémoire : chaque kernel écrit à la même adresse mémoire
- tu fais une boucle relativement longue : ton kernel doit probablement planter avant la fin de la boucle
- du coup, ton readBuffer lit une portion mémoire qui n'est pas initialisée
- tes retours d'erreur ne doit pas correspondre à l'erreur réelle

En bref, peux tu expliquer ce que tu souhaites faire dans ton kernel et donner ton code coté CPU ?

[ebalexandr]

Re-bonjour,
Oui, je ne fait pas de calcule parallèle, car je n'ai pas appris pour l'instant comment le faire. Je ne sais pas ne plus comment est-ce qu'il faut faire pour éviter, que chaque kernel écrit à la même adresse.
S'il plante c'est peut être à cause de ma variable(f) qui est un float. Malheureusement ma carte ne gère pas l'extension cl_khr_fp64. Mais j'ai testé le même boucle sur le CPU avec un float le résultat est faux et avec un double il est ok.
Je pense que le code que j'ai fait n'est pas adapter pour être exécuter sur le GPU, je l'ai fait juste pour apprendre OpenCL, car toute la documentation est en anglais, donc je voulais tester différents combinaisons pour voir que ce que ça donne, mais comme vous m'avez demandé voici la source de mon programme :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>
 
const char kernel_cl[] =
{
//"#pragma OPENCL EXTENSION all : enable\n" //n'est pas suporter par ma carte
    "__kernel void testopencl(__global float *out)\n"
    "{\n"
    "uint i=0xffffffff;\n"
//  "double f=0.0;\n" float pressision n'est pas suporter
    "float f=0.0;\n"
    "while(i)\n"
    "{\n"
    "i--;\n"
    "f=f+0.000001;\n"
    "}\n"
    "*out=f;\n"
    "}"
};
 
 
int main()
{
    cl_int err;
    cl_platform_id platform_id;
    cl_device_id device_id;
    cl_uint platforms, devices;
    cl_context_properties properties[]={CL_CONTEXT_PLATFORM,0,0};
    cl_context context;
    cl_command_queue command_queue ;
    cl_program program;
    cl_kernel kernel;
    cl_mem output_buffer;
    cl_float results;
    const char* kernel_src=kernel_cl;
    size_t len;
    size_t maxworkitemsize[3], global_work_size[3], local_work_size[3];//3 pour 3 dimention
    char buffer[4096];
 
    //la meme fonction sur le CPU
    uint i=0xffffffff;
    double f=0.0;//avec double pressision on obtien un resultat correcte
    while (i)
    {
        i--;
        f=f+0.000001;
    }
    printf("Resultat CPU:%f\n",f);
    // recuperation platform_id
    err=clGetPlatformIDs(1, &platform_id, &platforms);
    if (err != CL_SUCCESS)
    {
        printf("Unable to get Platform ID. err Code=%d\n",err);
        exit(1);
    }
    // recuperation device_id
    err=clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, 1, &device_id, &devices);
    if (err != CL_SUCCESS)
    {
        printf("Unable to get Device ID. err Code=%d\n",err);
        exit(1);
    }
 
    // des info sur ma carte
    err=clGetDeviceInfo(device_id,CL_DEVICE_VERSION,sizeof(buffer),buffer, NULL);
    printf("CL_DEVICE_VERSION:%s\n",buffer);
    err=clGetDeviceInfo(device_id,CL_DEVICE_EXTENSIONS,sizeof(buffer),buffer, NULL);
    printf("CL_DEVICE_EXTENSIONS:%s\n",buffer);
    err=clGetDeviceInfo(device_id,CL_DRIVER_VERSION,sizeof(buffer),buffer, NULL);
    printf("CL_DRIVER_VERSION:%s\n",buffer);
    err=clGetDeviceInfo(device_id,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(global_work_size[0]),global_work_size, NULL);
    printf("CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:%d\n",global_work_size[0]);
    err=clGetDeviceInfo(device_id,CL_DEVICE_ADDRESS_BITS,sizeof(global_work_size[0]),global_work_size, NULL);
    printf("CL_DEVICE_ADDRESS_BITS:%d\n",global_work_size[0]);
    err=clGetDeviceInfo(device_id,CL_DEVICE_MAX_WORK_GROUP_SIZE,sizeof(len),&len, NULL);
    printf("CL_DEVICE_MAX_WORK_GROUP_SIZE:%d\n",len);
    err=clGetDeviceInfo(device_id,CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(maxworkitemsize),&maxworkitemsize, NULL);
    printf("CL_DEVICE_MAX_WORK_ITEM_SIZES:(%d,%d,%d)\n",maxworkitemsize[0],maxworkitemsize[1],maxworkitemsize[2]);
 
    properties[1]=(cl_context_properties)platform_id;
    // creation du contenu
    context = clCreateContext(properties, 1, &device_id, NULL, NULL, &err);
    if (err != CL_SUCCESS)
    {
        printf("Unable to create context. Error Code=%d\n",err);
        exit(1);
    }
    // creation command queue
    command_queue = clCreateCommandQueue(context,device_id, 0, &err);
    if (err != CL_SUCCESS)
    {
        printf("Unable to create command queue. Error Code=%d\n",err);
        exit(1);
    }
    // creation program
    program = clCreateProgramWithSource(context, 1 ,&kernel_src, NULL, &err);
    if (err != CL_SUCCESS)
    {
        printf("Unable to create program object. Error Code=%d\n",err);
        exit(1);
    }
    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Build failed. Error Code=%d\n", err);
        // affichage de log de compillation
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        printf("%s\n",buffer);
        exit(1);
    }
    // creation du kernel
    kernel = clCreateKernel(program, "testopencl", &err);
    if (err != CL_SUCCESS)
    {
        printf("Unable to create kernel object. Error Code=%d\n",err);
        exit(1);
    }
    // creation du buffer de sortie
    output_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float), NULL ,NULL);
    // set the kernel arguments
    err=clSetKernelArg(kernel, 0, sizeof(cl_mem), &output_buffer);
    if (err != CL_SUCCESS)
    {
        printf("Unable to set kernel arguments. Error Code=%d\n",err);
        exit(1);
    }
    // global & local work size
    global_work_size[0]=(512+512+64)*512;//(512+512+64)*512CL_DEVICE_MAX_WORK_ITEM_SIZES(x,y,z)*CL_DEVICE_MAX_WORK_GROUP_SIZE
    local_work_size[0]=512;// pour ma carte (512;512;64)
    //execution
    err = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,&global_work_size,&local_work_size, NULL, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Unable to enqueue kernel command. Error Code=%d\n",err);
        exit(1);
    }
    // on attend que la commande se termine
    clFinish(command_queue);
    // lescture du buffer de sortie
    err = clEnqueueReadBuffer(command_queue, output_buffer, CL_TRUE, 0, sizeof(float), &results, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error enqueuing read buffer command. Error Code=%d\n",err);
        //exit(1);
    }
    printf("Resultat GPU:%f %d\n",results,sizeof(float));
    // liberation de memoire
    clReleaseMemObject(output_buffer);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(command_queue);
    clReleaseContext(context);
 
    return 0;
}

En tous cas merci pour vos réponses. Je continue à chercher sur le Net plus d'informations. L'API OpenCl est souvent comparée à CUDA, donc j'ai commencé à étudier sa documentation.
http://tcuvelier.developpez.com/gpgp.../introduction/

[gbdivers]

Bonjour

S'il plante c'est peut être à cause de ma variable(f) qui est un float. Malheureusement ma carte ne gère pas l'extension cl_khr_fp64. Mais j'ai testé le même boucle sur le CPU avec un float le résultat est faux et avec un double il est ok.

Non. Une erreur d'exécution d'un kernel ne provoquera pas d'erreur d'exécution du programme (il faut passer par les clEvent pour récupérer les erreurs des kernels)
La seule chose que te garantisse clFinish() (qui n'est d'ailleurs pas utile avec clEnqueueNDRangeKernel() sauf erreur de ma part), c'est que tes kernels ont terminé de travailler... mais pas du pourquoi (en particulier si les kernels ont plantés)

Tes calculs avec float ne provoqueront jamais d'erreur d'exécution (simplement l'envoi d'exception dans certain cas). Ton problème est un classique cas de problème de précision des nombres réels. Fais une recherche dans les forums C ou C++ pour les détails.

Pour résumé :
en math :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

x * y <=> additionne y fois la valeur x

en informatique :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
int y;
float x, r1, r2 = 0;
 
r1 = x * y;
for(int yy=0; yy<y; ++yy) { r2 += x; }
 
bool r = (r1 == r2); // sera toujours faux ! (ou presque toujours, ça peut au hasard tomber juste)

Tu ne peux pas remplacer une multiplication par une boucle faisant n additions.

Oui, je ne fait pas de calcule parallèle, car je n'ai pas appris pour l'instant comment le faire.

En gros, tu essaies d'apprendre OpenCL sans connaitre (ni apprendre) le calcul parallèle...
C'est un peu étrange, non ? C'est comme si tu apprenais à conduire une voiture sans connaître le code de la route : tu risques de faire n'importe quoi.

Donc, fais les choses dans l'ordre : apprend le C ou le C++ (en particulier les problèmes de calculs numériques), apprend l'optimisation du code (le GPGPU est une optimisation donc il faut déjà connaître les méthodes d'optimisation sans GPGPU avant) et ensuite le calcul parallèle (en l'appliquant avec OpenCL si tu veux). Je suis de ceux qui pense que l'on peut apprendre le GPGPU sans avoir besoin d'être un expert en programmation (on gagnera quand même en performances même si c'est pas optimal) mais il ne faut quand même pas penser que le GPGPU est la solution miracle pour tous les problèmes de performances.

Bon courage


PS : j'avais oublié de te préciser un point : avec OpenCL, tu n'es pas obligé de donner local_work_size. OpenCL calcul lui même le paramètre local_work_size optimal si on lui fournit pas.

[ebalexandr]

Re-bonjour,

La documentation de CUDA me suffi merci