How to calculate wavefronts that have both broadcasts and bank conflicts

Developer Tools Nsight Compute
1

Hi, everyone,

I have a question about Bank Conflict that I want to ask. If a warp of 32 threads accesses an address that has both broadcast and bank conflict, how many wavefronts would be generated? Or in this case, is it an 8-way conflict or a 4-way conflict?

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;
}

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 first 16 threads and the last 16 threads access the same data address, and the shared memory variable has dimensions of 8*128), 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 there are only 4 banks. I’m not sure if the broadcast mechanism will work here, meaning each bank will have 4 addresses, and each address is accessed by two threads. I’m not sure if this situation will generate 8 wavefronts or 4 wavefronts.

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

Best,

LiTian