# Problems Understanding Bank Conflicts

**URL:** <https://forums.developer.nvidia.com/t/problems-understanding-bank-conflicts/12167>\
**Category:** CUDA Programming and Performance\
**Created:** [September 15, 2009, 10:53am UTC](https://forums.developer.nvidia.com/t/problems-understanding-bank-conflicts/12167 "2009-09-15T10:53:09Z")\
**Posts on this page:** 2\
**Page:** 1

<div class="post-metadata">

**Author:** ![\_Brian](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@\_Brian](https://forums.developer.nvidia.com/u/_Brian)\
**Post date:** [September 15, 2009, 10:53am UTC](https://forums.developer.nvidia.com/t/problems-understanding-bank-conflicts/12167/1 "2009-09-15T10:53:09Z")

</div>

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

---

<div class="post-metadata">

**Author:** ![\_Brian](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@\_Brian](https://forums.developer.nvidia.com/u/_Brian)\
**Post date:** [September 16, 2009, 7:00pm UTC](https://forums.developer.nvidia.com/t/problems-understanding-bank-conflicts/12167/2 "2009-09-16T19:00:31Z")

</div>

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…