Problème semaphore openCL
Bonjour,
J'ai un problème en utilisant un semaphor pour bloquer l'accès à un tableau pour chaque work-item, voici le code:
Code:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
|
__kernel void ker(__global float* v, __global int* semaphor)
{
int g = get_global_id(0);
// wait for reource are available
while (atom_xchg (semaphor, 1) != 0) {}
// utilisation de la resource pendant qu'elle est bloquée
v[g] = g;
// deblocag de la resource
atom_xchg (semaphor, 0);
return;
} |
Lorsque je lance la fonction ça plante et je ne comprends pas pourquoi, si quelqu'un a un idée, merci. Y a t-il un autre moyen de bloquer cette ressource ? Merci
Algernon
ca ne marche toujours pas !!!
Bonjour,
j'ai essayé de suivre les instructions à la lettre du lien précédent mais ça crash toujours.
Je poste le code de mon kernel en entier
Code:
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
|
#pragma OPENCL EXTENSION cl_khr_int32_base_atomics : enable
#define LOCK(a) atom_cmpxchg (a, 0, 1)
#define UNLOCK(a) atom_xchg (a, 0)
/* Return the distance between the two given points
* param a: Pointer to a constant address memory of the first point
* param b: Pointer to a constant address memory of the second point
* param dim: Dimension of points
*/
float sqrtDistance (__global float* a, __global float* b, const int dim)
{
float dist = 0.0f;
for (int i = 0; i < dim; ++i)
dist += ((a[i] - b[i]) * (a[i] - b[i]));
return sqrt (dist);
}
/*
* Get the semaphor to lock resources for the current work-item
* param semaphor: The semaphor used to lock resources
*/
void GetSemaphor (__global long* semaphor)
{
// Try to write 1 in the location pointed by the semaphor and reads its current content which is stored in "occupied"
// If the value is 0, the function may proceed because the resource is available
while (LOCK (semaphor));
}
/*
* Release the semaphor in order to allow another work-item to use the resource
* param semaphor: The semaphor used to lock resources
*/
void ReleaseSemaphor (__global long* semaphor)
{
UNLOCK (semaphor);
}
/* Return the nearest neighbor of the query point
* param vertices: Vertices coordinates among which to search the nearest neighbors (__constant to allow all work-items to read only)
* param queries: Array containing the query points coordinates
* param nearest: Array containing the nearest neighbor, (__global to allow all work-items to read/write it)
* param nearestDist: Array containing the nearest distances, (__global to allow all work-items to read/write it)
* param k: Number of nearest neighbors required
* param nbVertices: Number of vertices in the cloud
* param dim: Dimension of points
* param semaphor: Semaphor in order to update the array of distance and nearest neighbors
*/
__kernel void nearestNeighbor (__global float* vertices,
__global float* queries,
__global int* nearest,
__global float* nearestDist,
__global const int* k,
__global const int* nbVertices,
__global const int* dim,
__global long* semaphor)
{
__private int verticesIdx = get_global_id (0);
if (verticesIdx >= *nbVertices)
return;
__global float* ptrCurrentVertex = vertices + (*dim * verticesIdx);
float distanceVertexToQuery = sqrtDistance (ptrCurrentVertex, queries, *dim);
if (distanceVertexToQuery >= nearestDist[*k - 1])
return;
// Need to compare the distance computed with the maximum distance
while (LOCK (semaphor));
// If the distance from the current vertex is lower than the greater distance stored in the array of k distances
if (distanceVertexToQuery < nearestDist[*k - 1])
{
for (int i = *k - 2; i >= -1; --i) // Search the position where to insert the new distance: distanceVertexToQuery
{
if ((i == -1) || (distanceVertexToQuery >= nearestDist[i]))
{
for (int j = *k - 1; j > i + 1; --j)
{
nearestDist[j] = nearestDist[j - 1];
nearest[j] = nearest[j - 1];
}
nearestDist[i + 1] = distanceVertexToQuery;
nearest[i + 1] = verticesIdx;
break;
}
}
}
UNLOCK (semaphor);
//ReleaseSemaphor (semaphor); // release the resources, so another work-item can grab it and start working
} |
je appel pour enqueueNDRange avec
Code:
1
2
3
4
|
size_t localThreads = 1;
size_t globalThreads = 1500;
status = clEnqueueNDRangeKernel (commandQueue, Kernel, 1, NULL, &globalThreads, &localThreads, 0, NULL, NULL); |
Au cas où, je reprécise mon problème et ce que fait mon kernel:
en fait j'ai 1500 points 3D et je cherche les k plus proches voisins du point queries (3 coordonnées) parmi ces 1500 points.
Donc je calcule en parallèle les distance entre mon point queries et chacun des 1500 points et ensuite, j'ai un algo de tri pour pouvoir trier mes points les plus proches en fonction de leur distance au point requete.
Mais je ne sais pas comment faire pour bloquer tous les threads pendant que un seul modifie mon tableau à la fois.
La solution que j'ai mise ne fonctionne pas et la carte crash, j'ai aussi testé avec BARRIER mais même problème. Combien de local thread dois-je allouer et combien de global threads ? Car peut-être c'est la mon problème?
Merci pour votre aide
Algernon
