Hello, NV’s experts I have a problem about the bank conflict of shared_mem. I know that the shared_mem is divided into 32 banks, and 4bytes per bank. GPU will raise bank conflict if different threads(in the same warp) access different addresses in a bank. I do some test to check it. test1, as f…

[image] Shaquille: I think it cannot match above theory, obviously, T0, T8, T16, T24 are in the same bank。other threads are similar with them, why GPU didn’t raise bank conflict? When you store (or load) more than 4 bytes per thread, which is like saying more than 128 bytes per warp, the GPU …

thanks Robert, so kindly I want to figure out the detail about shared_mem bank conflict, and try to find method to avoid it. my application is: I want to transpose the existing registers into shared memory. like this: [image] my warp is divided into 8x4, and each thread hold 4 float4(s), or say…

You could take a look at cub::WarpStore CUB: WarpStore< T, ITEMS_PER_THREAD, ALGORITHM, LOGICAL_WARP_THREADS, PTX_ARCH > Class Template Reference

hello, striker: so cool name ! thanks for you suggestion at first I found it cannot help my application. the cub’s warp_transpose is different from my expectation let me check its code, as following: #pragma unroll for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++) { threa…

thank you, it match my expectation

I have verified that if every thread store 8 floats(32 bytes), bank conflicts occur, in your explanation, it should have 8 transactions, every 4 threads make up a transaction, but it is not true, I test on L20 card, please give some light, thanks!

bank conflicts will depend on the access pattern across threads. 32 bytes cannot be stored by a single thread in a single transaction (or in a single SASS instruction). So at the SASS level there will definitely be multiple transactions (instructions) involved, and depending on those instructions, …

Thanks, the following source code on L20 will cause bank conflict, you can verify use ncu: #define N 32 #define factor 8 __global__ void bankConflictExample(float *input, float *output, int M) { // Declare shared memory (assuming 32 banks on the device) __shared__ float sharedMem[N*factor]; /…

Yes, your access pattern with a stride (greater than 1) will cause bank conflicts. I don’t think I said anywhere that accessing multiple bytes per thread automatically eliminates the possibility of bank conflicts. I don’t really understand what question is being asked, if any, and probably won’t b…

How to understand the bank conflict of shared_mem

Accelerated Computing CUDA CUDA Programming and Performance

Robert_Crovella July 26, 2023, 3:02pm 6

Some additional resources that may be of interest: 1 2 3

Topic		Replies	Views
Shared memory bank conflict CUDA Programming and Performance	4	715	July 30, 2025
dont understand bank conflicts for shared mem CUDA Programming and Performance	7	2851	March 31, 2010
Help understanding bank conflicts in transpose example CUDA Programming and Performance	5	6895	February 8, 2009
Requesting clarification for Shared Memory Bank Conflicts and Shared memory access? CUDA Programming and Performance hw , cuda	10	5637	January 23, 2024
Why there is random bank conflicts? CUDA-MEMCHECK cuda	2	1308	September 19, 2023
Shared Memory "Bank Conflicts" I'am confused... CUDA Programming and Performance	14	3758	November 20, 2025
Bank Conflict when each thread accesses 2 elements CUDA Programming and Performance	8	5757	July 9, 2010
Very strange share memory bank conflicts CUDA-MEMCHECK cuda	1	1021	October 15, 2021
shared memory bank conflicts cc 2.0 CUDA Programming and Performance	3	993	December 29, 2011
do not understand bank conflicts please help CUDA Programming and Performance	7	2847	December 22, 2012

How to understand the bank conflict of shared_mem

Related topics