What is a bank conflict? (Doing Cuda/OpenCL programming)

by

For nvidia (and amd for that matter) gpus the local memory is divided into memorybanks. Each bank can only address one dataset at a time, so if a halfwarp tries to load/store data from/to the same bank the access has to be serialized (this is a bank conflict). For gt200 gpus there are 16 banks (32banks for fermi), 16 or 32 banks for AMD gpus (57xx or higher: 32, everything below: 16)), which are interleaved with a granuity of 32bit (so byte 0-3 are in bank 1, 4-7 in bank 2, …, 64-69 in bank 1 and so on). For a better visualization it basically looks like this:

Bank    |      1      |      2      |      3      |...
Address |  0  1  2  3 |  4  5  6  7 |  8  9 10 11 |...
Address | 64 65 66 67 | 68 69 70 71 | 72 73 74 75 |...
...

So if each thread in a halfwarp accesses successive 32bit values there are no bank conflicts. An exception from this rule (every thread must access its own bank) are broadcasts:
If all threads access the same address, the value is only read once and broadcasted to all threads (for GT200 it has to be all threads in the halfwarp accessing the same address, iirc fermi and AMD gpus can do this for any number of threads accessing the same value).

More Related Contents:

  1. How is CUDA memory managed?
  2. How to measure the inner kernel time in NVIDIA CUDA?
  3. How do CUDA blocks/warps/threads map onto CUDA cores?
  4. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  5. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  6. nvidia-smi Volatile GPU-Utilization explanation?
  7. How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
  8. How do I select which GPU to run a job on?
  9. How can I make tensorflow run on a GPU with capability 2.x?
  10. CUDA determining threads per block, blocks per grid
  11. How does CUDA assign device IDs to GPUs?
  12. Error Message : Cannot find or open the PDB file
  13. CUDA compute capability requirements
  14. Cuda atomics change flag
  15. Modifying registry to increase GPU timeout, windows 7
  16. How can I implement a custom atomic function involving several variables?
  17. Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA
  18. How to create NVIDIA OpenCL project
  19. How to verify CuDNN installation?
  20. What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?
  21. 128 bit integer on cuda?
  22. Any particular function to initialize GPU other than the first cudaMalloc call?
  23. Cuda kernel returning vectors
  24. CUDA and nvcc: using the preprocessor to choose between float or double
  25. Polymorphism and derived classes in CUDA / CUDA Thrust
  26. device function pointers
  27. cudaMemset() – does it set bytes or integers?
  28. Are cuda kernel calls synchronous or asynchronous
  29. CUDA function pointers
  30. Using an array of device function pointers

Leave a Comment