Struct shared between threads

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

I have a lot of thread that must write in the same struct. How can I synchronize the struct-accesses between threads?

here some code lines of the kernel

[codebox]contax = threadIdx.x;

contay = threadIdx.y;

mappascelta[((contay*x_fotogramma)+contax)] = -1.0;

for(contdirez=1; contdirez<=8; contdirez++){

for (contordine=1; contordine <= 4; contordine++){

if(someCondition){

 <u>correlazioneOUTPUT[(contdirez-1)*4+contordine-1].numpixel++;</u>

 <u>correlazioneOUTPUT[32+contordine-1].numpixel++;</u>

 <u>correlazioneOUTPUT[(contdirez-1)*4+contordine-1].lunghezza += sqrt(pow((xlocale-contax)*(DECIMIMICRON_X/(x_fotogramma)), 2.0);</u>

}

}

}[/codebox]

The underlined text is the critical part. I need some mutex or some lock. In CUDA it’s possible?

2

Use atomicAdd() to update the struct members.

If this operation is performed often, it might be worth keeping a separate copy within each block and only atomically updating the global struct at the end of each block.