How to Calculate the Bank Conflicts

Accelerated Computing CUDA CUDA Programming and Performance
,
1

Hi everyone,

My graphics card is an RTX 4070 Super.

I ran into a problem while studying CUDA NCU analysis. When analyzing blocked matrix multiplication, I can’t understand how NCU calculates the Shared Load Bank Conflicts, mainly regarding the code below. This piece of code is taken from the blocked matrix multiplication and simplified to make it easier to understand.

#define TM 8
#define TK 8
#define BM 128
#define BK 128
#define BN 8 

// The blocksize is 16*16
const int tx = threadIdx.x; // 0~15
const int ty = threadIdx.y; // 0~15
const int bx = blockIdx.x;
const int by = blockIdx.y;
const int tid = ty * blockDim.x + tx;
 
int tCx = tid % 16; // 0~15

int c_col = tCx * TK;

__shared__ float Bs[BN][BK];

float c_val[TM][TK] = { {0.0f} };

for (int n = 0; n < BN; n++) {
	for (int k = 0; k < TK; k++) {
		c_val[n][k] += Bs[n][c_col + k];
	}
}

For a warp with 32 threads, the c_col indices are 0, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, 0, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, the Bank indices are

(n*128 + k)%32 = k
(n*128 + k+8)%32 = k + 8
(n*128 + k+16)%32 = k + 16
(n*128 + k+24)%32 = k + 24

(n*128 + k+32)%32 = k

…

(n*128 + k+120)%32 = k + 24

So each read only touches four banks, with 8 threads reading from each bank. I understand this as an 8-way conflict, so Bank Conflicts = 8-1 = 7, and since the loop runs 64 times, the total Bank Conflicts should be 448. But the NCU result shows only 64, I feel like I have a misunderstanding about bank conflict, but I can’t pinpoint exactly what’s wrong.

I would really appreciate it if I could get any answers or any hints about information related to this topic.

Best,

LiTian

2

I don’t know how close the simplified code is to your actual code, but

  • Using 32-bit loads this is a 4-way bank conflict, since each value gets broadcast to two threads (e.g. 0 and 16).
  • However: the compiler can infer that Bs[n][c_col] is 16-byte aligned, and the k-loop reads consecutive elements, so the compiler can actually unroll the loop and use 128-bit loads from SMEM, dividing the number of required load ops by 4. If I drop your example code into Compiler Explorer (and make use of c_val so it doesn’t get removed as unused), I see the SASS output containing LDS.128 ops to read from SMEM.
    You still need 4 wavefronts per load op, since you’re accessing each of the 16 banks in use 4 times (e.g. bank 0 serves data from indices 0, 32, 64, 96). 4 wavefronts is what you’d expect from a 128-bit smem load without broadcast, but in this case the broadcast and the use of only half the banks per load op cancel each other out.
3

Hi，Nanodeoclus，

Thanks a lot for your reply, but I’m a bit confused. I’ve provided a full demo below. This demo is only for testing Shared Load Instructions, Wavefronts, and Bank Conflicts, and doesn’t really have any practical significance.

#include <cuda_runtime.h>
#include <iostream>
#include <stdio.h>
#include <cmath>

#define TILE_SIZE 16   // 线程块的基本tile大小

// 错误处理宏
#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

static void HandleError(cudaError_t err, const char* file, int line) {
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}


__global__ void bank_test(const float* B, float* C) {
    int tx = threadIdx.x; // 0~15
    int ty = threadIdx.y; // 0~15
    int tid = ty * blockDim.x + tx;

    int tCx = tid % 16; // 0~15
    int tCy = tid / 16; // 0~15
    int c_row = tCy * 8;
    int c_col = tCx * 8;

    int tBx = tid & (32 - 1); // tid % 32 0~31
    int tBy = tid / 32; // 0~7

    __shared__ float Bs[8][128];

    float c_val[8][8] = { {0.0f} };

    for (int k = 0; k < 128; k += 32) {
        Bs[tBy][k + tBx] = B[tBy * 128 + k + tBx];
    }

    __syncthreads();

    for (int n = 0; n < 8; n++) {
        for (int k = 0; k < 8; k++) {
            c_val[n][k] += Bs[n][c_col + k];
        }
    }

    __syncthreads();

    for (int i = 0; i < 8; i++) {
        int global_row = c_row + i;
        for (int j = 0; j < 8; j++) {
            int global_col = c_col + j;
            C[global_row * 128 + global_col] = c_val[i][j];
        }
    }

}

// A, B, C are device pointers (i.e. pointers to memory on the GPU)
void main() {
    int M = 128, N = 8, K = 128;
    float* B = new float[N * K];
    float* C = new float[M * K];
    float* dev_A, * dev_B, * dev_C;

    for (int j = 0; j < N * K; j++) {
        B[j] = rand() / (float)RAND_MAX;
    }

    HANDLE_ERROR(cudaMalloc((void**)&dev_B, N * K * sizeof(float)));
    HANDLE_ERROR(cudaMalloc((void**)&dev_C, M * K * sizeof(float)));

    HANDLE_ERROR(cudaMemcpy(dev_B, B, N * K * sizeof(float), cudaMemcpyHostToDevice));

    dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);
    bank_test << <1, threadsPerBlock >> > (dev_B, dev_C);

    cudaDeviceSynchronize();
    cudaError_t error = cudaGetLastError();
    if (error != cudaSuccess) {
        printf("CUDA Error: %s\n", cudaGetErrorString(error));
    }

    HANDLE_ERROR(cudaMemcpy(C, dev_C, M * K * sizeof(float), cudaMemcpyDeviceToHost));

    cudaFree(dev_B);
    cudaFree(dev_C);

    delete[] B;
    delete[] C;
}

I used Nsight Compute to get some metrics for shared memory, but they don’t match what I calculated myself. Here’s my understanding:

  1. The block size is 16*16, so there are 8 warps. The Bs[n][c_col + k] reads are because c_col + k are 8 consecutive floats in memory. The compiler will expand this into two LDS.128 instructions, which also matches the SASS source code. n ranges from 0~7, so one warp will issue 8*2=16 LDS instructions, and 8 warps in total send 8*16=128 instructions. This calculation matches the NCU results.

  2. For a warp with 32 threads, the c_col indices are 0, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, 0, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, the Bank indices are

    (n*128 + k)%32 = k
    (n    *128 + k+8)%32 = k + 8
    (n*128 + k+16)%32 = k + 16
    (n    *128 + k+24)%32 = k + 24

    (n*128 + k+32)%32 = k

    …

    (n*128 + k+120)%32 = k + 24

  3. The first 16 threads access the same addresses as the last 16 threads, so each bank only has four different addresses being accessed by different threads. So it’s considered a 4-way bank conflict, meaning the number of conflicts should be 4-1=3. That means each load operation will generate 4 wavefronts and 3 bank conflicts. So in total, it should be 128*4=512 wavefronts and 128*3=384 bank conflicts, but this doesn’t match the NCU results.

image
image1859×295 11.6 KB

4

Lds128 loads 512 bytes which requires a minimum of 4 128 byte wavefronts. 8 consecutive threads are processed per wavefront (8x 16 byte=128 byte).

With your access pattern threads 0-3 and 4-7 access the same banks leading to a two-way bank conflict.

512 wavefronts x 2way conflict = 1024 wavefronts, from which 512 are the result of bank conflicts.

5

HI, striker159

Thanks a lot for your reply, I’m a bit confused about your answer. The unit for LDS.128 should be bits, not bytes. Threads 0~31 access banks 0, 8, 16, 24 when k=0, so why do you mention a 2-way conflict here?

6

A wavefront is (up to) 128 bytes, i.e. the amount of data smem can serve in a single cycle (4 bytes per bank). Typically a LDS.128 executes as 4 wavefronts, with each wavefront mapping to 8 consecutive threads (0-7, 8-15, 16-23, 24-31), for the common case of lane-consecutive addresses.

However, in the presence of bank conflicts and broadcasts things get a bit complicated. LDS.128 can actually go down to 2 wavefronts (tested on H100 and RTX5090) for certain patterns of broadcast, and worst case it’s 32 wavefronts (32-way conflict).

Take this example code, executing one LDS.128 per kernel launch (templated kernel so ncu counts them separately):

#include <cuda_runtime.h>

template <int EXPERIMENT>
__global__ void bank_test(int *lane_offsets) {
    __shared__ int4 smem[1024];
    int lane = threadIdx.x % 32;
    int idx = lane_offsets[EXPERIMENT * 32 + lane];
    for (int i = 0; i < 1; ++i) {
        smem[i * 32 + lane] = int4{idx, idx + 1, idx + 2, idx + 3};
    }
    __syncthreads();
    int sum = 0;
    for (int i = 0; i < 1; ++i) {
        int4 x = smem[i * 128 + idx];
        sum += x.x + x.y + x.z + x.w;
    }
    // use result somehow
    lane_offsets[EXPERIMENT * 32 + lane] = sum;
}

int main() {
    constexpr int lane_offsets[][32] = {
            { // 2 wavefronts
                    0, 0, 0, 0, 0, 0, 0, 0,
                    1, 1, 1, 1, 1, 1, 1, 1,
                    2, 2, 2, 2, 2, 2, 2, 2,
                    3, 3, 3, 3, 3, 3, 3, 3,
            },
            { // 2 wavefronts
                    0, 1, 0, 1, 2, 3, 2, 3,
                    4, 5, 4, 5, 6, 7, 6, 7,
                    8, 9, 8, 9, 10, 11, 10, 11,
                    12, 13, 12, 13, 14, 15, 14, 15,
            },
            { // 4 wavefronts
                    0, 1, 0, 1, 2, 3, 2, 2,  // << note pattern break
                    4, 5, 4, 5, 6, 7, 6, 7,
                    0, 1, 0, 1, 2, 3, 2, 3,
                    4, 5, 4, 5, 6, 7, 6, 7,
            },
            { // 4 wavefronts
                    0, 1, 2, 3, 0, 1, 2, 3,
                    0, 1, 2, 3, 0, 1, 2, 3,
                    0, 1, 2, 3, 0, 1, 2, 3,
                    0, 1, 2, 3, 0, 1, 2, 3,
            },
            { // 4 wavefronts
                    0, 1, 2, 3, 4, 5, 6, 7,
                    0, 1, 2, 3, 4, 5, 6, 7,
                    0, 1, 2, 3, 4, 5, 6, 7,
                    0, 1, 2, 3, 4, 5, 6, 7,
            },
            { // 4 wavefronts
                    0, 2, 4, 6, 1, 3, 5, 7,
                    8, 10, 12, 14, 9, 11, 13, 15,
                    0, 2, 4, 6, 1, 3, 5, 7,
                    8, 10, 12, 14, 9, 11, 13, 15,
            },
            { // 8 wavefronts
                    0, 2, 4, 6, 8, 10, 12, 14,
                    1, 3, 5, 7, 9, 11, 13, 15,
                    0, 2, 4, 6, 8, 10, 12, 14,
                    1, 3, 5, 7, 9, 11, 13, 15,
            },
    };
    int *dev_lane_offsets;
    cudaMalloc((void **)&dev_lane_offsets, sizeof(lane_offsets));
    cudaMemcpy(dev_lane_offsets, lane_offsets, sizeof(lane_offsets), cudaMemcpyHostToDevice);
    bank_test<0><<<1, 32>>>(dev_lane_offsets);
    bank_test<1><<<1, 32>>>(dev_lane_offsets);
    bank_test<2><<<1, 32>>>(dev_lane_offsets);
    bank_test<3><<<1, 32>>>(dev_lane_offsets);
    bank_test<4><<<1, 32>>>(dev_lane_offsets);
    bank_test<5><<<1, 32>>>(dev_lane_offsets);
    bank_test<6><<<1, 32>>>(dev_lane_offsets);
    cudaDeviceSynchronize();
    cudaFree(dev_lane_offsets);
    return 0;
}

The corresponding ncu output is:

    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    2
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    2
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    4
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    4
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    4
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    4
    l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum                                                                    8
7

I would guess (one has to test) that for 64 bit and 128 bit accesses, each lane still receives 32 bit per cycle, but in a different order per lane to reduce bank conflicts.

So with 128 bit access, e.g. float4, 8 lanes get x, 8 lanes get y, 8 z and 8 the w component. And they cycle for the 4 accesses.

Whether it is possible to use broadcasts would depend on which 8 lanes are read together with the same component.

I am not sure, whether the compiler or hardware at runtime can change those patterns to reduce the number of bank conflicts.

It could also be that we separately have to look into reading from shared memory and broadcasting to the lanes. Those components may have different bandwidth.

8

Thanks a lot for your answer, but I couldn’t figure out the root cause of the question I asked from it. I think I need to spend some time thinking it over more carefully.

9

I think what striker159 suggested would seem essentially correct: Unless you use a special broadcast pattern, LDS.128 gets processed in four phases, each taking 8 consecutive threads into account. You have a 2-way bank conflict in each phase, i.e. each set of 8 consecutive threads uses bank indices [0, 8, 16, 24, 0, 8, 16, 24]. This gets translated into two wavefronts for each phase, eight in total per LDS.128.

But it’s also clear that with the presence of broadcasting, something more complex is happening.
If you modify my example using the below patterns, you will see that the first one uses only 4 wavefronts (with 2 bank conflicts, confusingly), while the second one uses 8 wavefronts (with 4 bank conflicts). Note that my example uses int4, thus to get from offset to bank index you need to multiply by 4.

            { // 4 wavefronts
                    0, 2, 0, 2, 4, 6, 4, 6,
                    8, 10, 8, 10, 12, 14, 12, 14,
                    16, 18, 16, 18, 20, 22, 20, 22,
                    24, 26, 24, 26, 28, 30, 28, 30,
            },
            { // 8 wavefronts
                    0, 2, 4, 6, 8, 10, 12, 14,
                    16, 18, 20, 22, 24, 26, 28, 30,
                    0, 2, 4, 6, 8, 10, 12, 14,
                    16, 18, 20, 22, 24, 26, 28, 30,
            },
10

Hi, striker159

I found an answer on a forum about the wide-LDS Bank Conflict that is similar to your explanation( Problem about bank conflict test - #2 by Greg ). It seems like it could answer my question, but when I was testing, I ran into a new issue. If multiple threads among the 32 threads access the same bank at the same address (meaning when there’s broadcasting) while coexisting with LDS.64 or LDS.128, the explanation above doesn’t seem to work in that case. Below is the demo I used for my test.

#include <cuda_runtime.h>
#include <iostream>
#include <stdio.h>


// 错误处理宏
#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

static void HandleError(cudaError_t err, const char* file, int line) {
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}


template <int n>
__global__ void bank_test1(float* C) {
    int tid = threadIdx.x;

    int tCy = tid / 8; // 0~3
    int tC = tCy * 8; // 0、8、16、24
    __shared__ float smem[32];
    smem[tid] = tid;
    __syncthreads();
    for (int i = 0; i < n; i++) {
        C[tid] += smem[tC + i];
    }

}

template <int n>
__global__ void bank_test2(float* C) {
    int tid = threadIdx.x;

    __shared__ float smem[32 * n];
    for (int i = 0; i < n; i++) {
        smem[tid + 32 * n] = tid;
    }
    
    __syncthreads();
    for (int i = 0; i < n; i++) {
        C[tid] += smem[tid * n + i];
    }

}



void main() {
    int M = 32, K = 1;
    float* dev_C1, * dev_C2;

    HANDLE_ERROR(cudaMalloc((void**)&dev_C1, M * K * sizeof(float)));
    HANDLE_ERROR(cudaMalloc((void**)&dev_C2, M * K * sizeof(float)));

    bank_test1<2> << <1, 32 >> > (dev_C1);
    bank_test2<2> << <1, 32 >> > (dev_C2);
    bank_test1<4> << <1, 32 >> > (dev_C1);
    bank_test2<4> << <1, 32 >> > (dev_C2);

    cudaDeviceSynchronize();
    cudaError_t error = cudaGetLastError();
    if (error != cudaSuccess) {
        printf("CUDA Error: %s\n", cudaGetErrorString(error));
    }

    cudaFree(dev_C1);
    cudaFree(dev_C2);

}

For bank_test2: when n=2 and i is unrolled, it becomes an LDS.64 instruction. Each i, 32 threads access 16 banks, with each bank accessed at 2 addresses. Since threads 0~15 access 0~31 for i=0 and i=1, the compiler treats the access of threads 0~15 at i=0 and i=1 as a single wavefront, so in total there are 2 wavefronts.

When n=4 and i is unrolled, it’s an LDS.128 instruction. Each i, 32 threads access 16 banks, 2 addresses per bank. Since threads 0~7 access 0~31 for i=0,1,2,3, the compiler treats the access of threads 0~7 at i=0,1,2,3 as a single wavefront, totaling 4 wavefronts. I think the wavefront analysis of bank_test2 is consistent with the principles in the links mentioned before and your explanation.

But for bank_test1: when n=2 and i is unrolled, it’s an LDS.64 instruction. Each i, 32 threads access the same address in 4 banks, so it’s a broadcast. In theory, LDS.64 needs 2 wavefronts, but here there is actually only 1 wavefront.

When n=4 and i is unrolled, it’s an LDS.128 instruction. Each i, 32 threads access the same address in 4 banks, so it’s a broadcast. In theory, LDS.128 needs 4 wavefronts, but here there are actually only 2 wavefronts.

Do you know why this happens in the case of broadcasts?

image
image2643×1149 405 KB

11

Yes, It seems like broadcasting makes the wavefront analysis more complicated, but I haven’t found any related analysis or information so far.

12

I’ve done some more experiments, and I’ve so far found only two special cases for broadcasting LDS.64 and LDS.128:

{
        a, b, a, b, c, d, c, d, e, f, e, f, g, h, g, h,
        i, j, i, j, k, l, k, l, m, n, m, n, o, p, o, p,
},
{
        a, a, b, b, c, c, d, d, e, e, f, f, g, g, h, h,
        i, i, j, j, k, k, l, l, m, m, n, n, o, o, p, p,
},

If warp-lanes provide pairwise identical addresses according to these patterns, LDS.64 will have baseline of 1 wavefront (instead of 2) and LDS.128 will have a baseline of 2 wavefronts (instead of 4; consisting of threads 0-15 and 16-31). There are still bank conflicts possible, of course (e.g. if (a != b) && (bank(a) == bank(b))), which cause extra wavefronts. Identical addresses (e.g. b == f) do not cause extra wavefronts.

13

Great finding!

So broadcasting those will read 128B/cycle from shared memory and write 256B/cycle to registers?

Or 128B/cycle for both, but shared memory is only occupied half as many cycles.

Shared memory access is a MIO pipeline asynchronous operation with the short scoreboard to fill registers.

14

Interesting finding, I just tried it too, when broadcasting on bankidx, wavefronts come out differently for ‘aabb’ and ‘abab’ because of the different order.

15

Yeah, it seems that way, but there’s very little information online about the specific calculations of these NCU indicators, which is a bit frustrating.

16

Will

{
        a, b, a, b, c, d, c, d, e, f, e, f, g, h, g, h,
        i, j, k, l, m, n, o, p, q, r, s, t, u, v, w, x,
}

need 3 wavefronts for LDS.128 or 4 wavefronts?
The LDS.128 divides into groups of 8 or 16 lanes.
The first 16 lanes can be done with a single wavefront (and would according to the previous experiments). The second 16 lanes in all cases need 2 wavefronts.