Writing to pointer struct from threads

Hello,

I am passing a raw pointer to a kernel, and using some kind of mutexes (some kind :P), each thread adds a value to the same member of the struct.

The pointer, holds an array of structures.

Each block works on a seperate structure of that array. So, block0 works on structure0, block1 on structure1 etc.

For example:

[codebox]

device static unsigned int kernel_mutexes[N];

struct someStruct {

int counter;

};

glogal void doSomeStuff( struct someStruct* str ) {

const unsigned int structureIdx = blockIdx.x;

// … do some other stuff here …

__syncthreads( );

// Each thread in the current block, will get a chance to

// access and write the member variable as soon as the

// kernel_mutexes[structureIdx] is equal to their id.

// Also, the "static" keyword, guarantees that the thread

// access will start from 0.

while( true ) {

   if( kernel_mutexes[structureIdx] == threadIdx.x ) {

        str[ structureIdx ].counter++;

        atomicAdd( &kernel_mutexes[structureIdx], 1 );

        __threadfence( );

        break;

   }

}

};

[/codebox]

Maybe my thought is waaaayyy off?

Any comments and thoughts are welcomed.

edit: fixed using atomicExch and atomicInc, it`s about 2ms faster than before o_O