Problems Understanding Bank Conflicts

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

i have some Problems understanding Bank Conflicts in shared memory.

In particular its the reduction example, which illustrates bank conflicts.

In the reduce1 example from the sdk:

[codebox]// do reduction in shared mem

for(unsigned int s=1; s < blockDim.x; s *= 2) 

{

    int index = 2 * s * tid;

if (index < blockDim.x)

    {

        sdata[index] += sdata[index + s];

    }

    __syncthreads();

}

[/codebox]

At the first Iteration each Thread access 2 successive Elements from shared memory.

In the Programming Guide its stated that: “any memory read or write request made of n addresses that fall in n distinct memory banks”

For my understanding thats the case in the example.

Thread 0 accesses Bank 0 and 1,

Thread 1 accesses Bank 2 and 3,

…

So there should not be any Bank Conflict in my opinion.

Can somebody clarify things a bit for me?

thanks,

Brian

2

I think i found the answer. If somebody else is interested:

Thread 0 till Thread 7 cause no Bank Conflict, but Thread 8 for example accesses Bank 0 like Thread 0 → Bank Conflict.

Its illustrated in the sdk programming guide 2.3 as well :">. Page 95, Figure 5-7, left: linear addressing with a stride of 2.

Thread can be closed…