Special cases of Shared Memory Bank Conflict

Accelerated Computing CUDA CUDA Programming and Performance
1

Does it count, and is reported by Nsight Compute, as a Shared Memory Bank Conflict when:

  1. the same thread makes memory requests to two different locations within the same bank?
  2. two (or more) threads, each from a different warp, make memory requests to different locations within the same bank?
2

It’s not possible, with the definition of “single location” being the 4 adjacent bytes that consist of one bank location. An instruction can only request up to 16 adjacent bytes, per thread. This is documented in the programming guide. Those 16 bytes can at most span up to 4 bytes in a “single” location, in any given bank.

Bank conflicts only have “meaning” when the consideration is within a single warp, and for a specific instruction issued.

There is no concept of a bank conflict considering threads of different warps.

There is no concept of bank conflict considering two different instructions.

3

Consider

__shared__ float cache[1024];
...
//threadIdx.x 0 makes the call
cache[31] += cache[63];

this increment is accessing two different addresses on bank #31.
Does it work as:
instruction: load from 63
instruction+1: load from 31
instruction+2: add
instruction+3: store to 31

4

yes, exactly. When I use the word “instruction” I meant SASS instruction. It doesn’t necessarily correspond to “a line of C++ code” or anything else.

For each of the 4 lines you have written that begin with instruction:, those would each be a SASS instruction. And bank conflicts only relate to one of those instructions at a time, and even at that, only for a single warp at a time.

Therefore, no bank conflicts there!