CODEWITHSUNDEEP
ccudaatomicgpu-shared-memory
lina
lina

Reputation: 1707

atomic operations on shared memory

How can I do an atomic operation on a shared memory?
I have something similar to this:

__shared__ int a[10];
//set a
if(tid<5)
  a[2]++;

Therefore 5 threads are incrementing a. How can I do this?
I know that in this way I am serializing the execution of 5 threads, but how does this effect the warp? would all the threads in the warp be serialized or just the first 5?

Upvotes: 2

Views: 751

Answers (1)

William Pursell
William Pursell

Reputation: 212544

Replace a[2]++ with

atomicAdd( a + 2, 1 );

if you can make a unsigned, you might prefer to use atomicInc() instead, but either one is going to kill performance.

Upvotes: 4

Related Questions