Solve bank conflict

[yofufufufu]

In my opinion, when loading data from global memory to shared memory(i.e. write shared memory) with vectorized access, because of the transposition, threads within a warp may write the same col in shared memory.
For example, thread 0 reads A[0][0] to A[0][3], thread 1 reads A[0][4] to A[0][7]. So thread 0 writes As[0][0] to As[3][0], thread 1 writes As[4][0] to As[7][0]. For a BM(=128) * BK(=8) size As, it is obvious that As[0][0] and As[4][0] are on the same bank, causing bank conflict.
So I think bank conflict will only occur when writing As not Bs. But in kernel v7 and v8, it seems like you try to optimize wrting to Bs:

SGEMM_CUDA/src/kernels/8_kernel_bank_extra_col.cuh

Lines 56 to 60 in 60cba6f

	tmp = reinterpret_cast<float4 *>(&B[innerRowB * N + innerColB * 4])[0];
	Bs[innerRowB * (BN + extraCols) + innerColB * 4 + 0] = tmp.x;
	Bs[innerRowB * (BN + extraCols) + innerColB * 4 + 1] = tmp.y;
	Bs[innerRowB * (BN + extraCols) + innerColB * 4 + 2] = tmp.z;
	Bs[innerRowB * (BN + extraCols) + innerColB * 4 + 3] = tmp.w;

Did I understand something wrong?

[yofufufufu]

Still looking forward to your reply 😿

[cudawarped]

@yofufufufu Your observation is correct which is most likely why this is not included in the blog, being a work in progress that didn't amount to anything.

In my experiment if I add padding here to avoid bank conflicts

__shared__ float As[(BM+2) * BK];

the increase in the number of registers (173 up from 168) decreases the occupancy to 2 significanlty reducing the performance.

[shixianc]

@cudawarped @yofufufufu I know it's old but post my new findings here anyways:
If I set extraCols to 4 it actually gives me ~0.1ms speed up from 4.3ms on my GPU.
I looked at SASS and found out that if padding cols is not multiple of 4, compiler will fail to do LDS.128 auto vec loading from smem to register.

// padding = 1, or 5 -> breaks down into multiple float32 LDS
LDS R78, [R76+-0x10]
LDS.128 R68, [R2+UR5+-0x10]
LDS.128 R72, [R2+UR5]
...
LDS R80, [R76+-0xc]
LDS R82, [R76+-0x8]
LDS R92, [R76+-0x4]
LDS R77, [R76]

// padding = 4 -> 2 LDS.128 for TM=8 vector
LDS.128 R68, [R3+UR5+-0x10]
LDS.128 R80, [R90]

Other than this, yes it's less possible to conflict on smem_b, but more on smem_a

PS: @cudawarped we may have a different GPU and config, my occupancy is not affected by the slight increase on SMEM, instead my occupancy is mainly bottlenecked by register count.

[cudawarped]

@shixianc I think you have misunderstood what @yofufufufu was asking/confirming. They were pointing out that the bank conflicts are caused by the writes to As not Bs, as a result changing extraCols can only impact performance by creating bank conflicts when writing to Bs.

As you have "may" have observed (the timing difference of 0.1ms could just be noise) using extraCols=4 "should" have better performance because it both removes the bank conflicts when writing to Bs introduced by extraCols=5 and restores the 128 bit writes to shared memory.

Either way if you want to remove the bank conflicts with padding you need to add padding of 2 to As as I mentioned above.

[shixianc]

@cudawarped
to clarify, I am only talking about As conflict as it's much less likely to conflict on Bs.

Here extraCols=4 allows compiler to do LDS.128 loading from As in the inner loop. However I want to point out it's still a 4-way conflict, and padding of 2 you mentioned will be a 2-way conflict meaning every bank still conflict twice in a warp access. But it looks to me from NCU that LDS.128 provides more benefits that offsets the small conflict here.

Ideally padding of 1 should make more sense if we ignore the LDS.128 later on, which offers 0 conlfict. What do you think?

[cudawarped]

@shixianc Apologies its been a while since I looked at this, please ignore my previous comments.

Taking a closer look, in 6_kernel_vectorize there are 2-way bank conflicts when

1. writing to As, and
2. reading from Bs.

Previously when I suggested adding a padding of 2 I was refering to removing the bank conflicts when writing to As in 10_kernel_warp_tiling but I neglected to mention that.

Additionaly in 7_kernel_resolve_bank_conflicts and 8_kernel_bank_extra_col @siboehm is only addressing the conflicts when reading from Bs as there are significanlty more reads from shared memory than writes to it (reason for shared memory in the first place). I completely missed this when replying to @yofufufufu because I was looking at 10_kernel_warp_tiling and focused solely on the less important conflicts arising from writes to As.

Anyway to answer you question correctly.

1. Addressing the bank confilicts when writing to As is simply a case of adding 4 columns of padding to As because the bank conflicts arise from threads accessing different rows of As in shared memory.
2. Unfortunately addressing the bank conflicts when reading from Bs can not be performed with such a simple approach because they arise from threads accessing the same row of Bs in shared memory. One simple solution (swizzling would be another but its not as simple) to this is to use padding of 16 colums combined with adding padding of 4 after each 32 columns of shared memory.

I've added the two fixes to my fork (cudawarped@19b397e) if you want to try this for yourself.

[shhyQAQ]

@cudawarped
In your code,you write const int paddedInnerColB = innerColB + (innerColB / 8),but we know,innerColB in a wrap are 0~31,and they have same innerRowB.So we can think innerRowB * (BN + extraCols) + paddedInnerColB * 4 which in a wrap as a continous row,which means we just need to think of paddedInnerColB * 4.
Your code use paddedInnerColB * 4,regardless of what number paddedInnerColB equals,there are 32 even numbers, there will be bank confit.This is my understanding, though I'm not sure if it's correct,but I think your code can't solve the bank conflict when read Bs