writing to an array of 64 ints

Accelerated Computing CUDA CUDA Programming and Performance
1

I still don’t understand fully the problem of bankconflicts.

I have an array of 64 integers and 64 threads. At some moment they write to the array at the same time in the following fashion:

sdata[0].array[ threadId.x ] = aValue;

My question is: How can I optimise this so that there are no or just a few memory conflicts?

2

why do you think there are bank conflicts in this code?

3

I think there are bank conflicts, because the warp size is 32, so there must be at least one conflict

4

Check the guide again, bank conflicts can only occur within each half warp. Your example has each thread writing to shared memory bank threadIdx.x (assuming the size of the value is 32 bits), so each will access a different bank. Bank conflicts are the last thing to optimize for anyways: global memory bandwidth matters much more.

5

As said in prev posts, there are no bank conflicts. Succesive integers of an array are placed in successive banks. And, Bank conflicts is the last thing to worry about when it comes to performance.