Semaphor doesn't works

Developer Tools Nsight Visual Studio Edition
1

Update

I tried some things and my experience is:
when I put into the code:
while(occupied > 0)
{
occupied = atom_xchg(sem, 1);
}

I can run the kernel without any problem, but as soon as I try to read back anything I got error -5 (CL_OUT_OF_RESOURCES) and the drivers is killed until automatic recover. Now my question is:
1.) is there any problem in my code?
2.) how can I do a working semaphore or mutex (sync mem access)?

Update end

Hi! When I read the buffer it’s always 3. What I didn’t do well? I have more workers. I expected multiple of 3.

__kernel void findCircles(__global uint *inputNodesArray, __global uint *inputNodesArraySizes, uint lastMapCounter,
					volatile __global uint *resultBuffer, volatile __global uint *resultBufferSize, volatile __global uint *sem)
{
uint x = get_global_id(0) ;
long occupied;
while(occupied > 0)
{
	occupied = atom_xchg(sem, 1);
}

(*resultBufferSize) += 3;

long prevVal;
prevVal = atom_xchg(sem, 0);
2

Hi, User Tamax77

When this kernel run in multiple threads, each thread gets “resultBufferSize” as 0 (suppose you input 0) and returns 3.

You expectation counts on the threads run sequencely and writing to the “resultBufferSize” is synchronized. But that’s not true.

3

Hi, thanks.
Why is it nor synct?
Because one of the threads repleace the sem by 1 the others has to wait:

while(occupied > 0)
{
occupied = atom_xchg(sem, 1);
}

until that one set the sem back to 1.

I think this is the main idea of mutexing.

4

I may have read wrong for your code. But in line 5, you need “long occupied = atom_xchg(semaphor, 1);”