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
} |
Problème semaphore openCL
Répondre avec citation
Partager