CUDA避免Bank Conflict

什么是Bank

为了提高内存读写带宽，共享内存被分割成了32个等大小的内存块，即Bank，因为一个Warp有32个线程，这样相当于一个线程对应一个内存Bank。在计算能力为5.x或更高的设备上，每个存储体在每个时钟周期都有32位的带宽，并且连续的32位被分配给连续的存储体，即bank有4个字节。

如 4 字节的浮点数，给共享内存数据分配 bank 时，数据 1、2、3…… 依次分给 bank1、bank2、bank3…… 第 32 个数据分给 bank1，如此 32 个一循环。

在访问shared memory时，因多个线程读写同一个Bank中的不同数据地址时，导致shared memory 并发读写 退化成顺序读写的现象叫做Bank Conflict；

避免Bank conflict

在与Shared memory的数据交换中，最好每次32个线程读写32个连续的word。如shared merory为32个word(4字节)，则在循环中，第一次，thread0 读0word，thread1 读1word，… ；第二次，thread0 读32word，thread1 读33word，… 。也可以是不规则的：如线程0读写bank1，线程1读写bank0。这种同一个时刻每个bank只被最多1个线程访问的情况下不会出现Bank conflict。

特殊情况如果有多个线程同时访问同一个bank的同一个地址的时候也不会产生Bank conflict，即broadcast。但当多个线程同时访问同一个bank不同地址时，Bank conflict就产生了。例如线程0访问地址0，而线程1访问地址32，由于它们在同一个bank，就导致了这种冲突。

改变数据存储方式是一种避免Bank conflict的方式

例如一个线程块有32个线程，每个线程有一个长度为6的float类型（4-bytes）数组（向量）。则这个数组可以有3种声明的方式：

方式1：__shared__ int Vector1[32*6];
方式2：__shared__ int Vector1[32][6];，行优先
方式3：__shared__ int Vector1[6][32];，列优先

方式1的情况下，需要对这个一维数组进行分割，如果每6个连续的字为一个向量，则结果和方式2存储方式相同。方式2，3比较直观，但是可能具有不同的性能，他们形状如下图：

图1 方式2（左图），方式3（右图）

方式 2：__shared__ int Vector1[32][6];这种声明方式是按照行优先（Row - major）的方式进行存储布局。如果把这个二维数组看作是一个表格，那么数据是一行一行地连续存储在共享内存中的。

方式 3：__shared__ int Vector1[6][32]; 这种声明方式是按照列优先（Column - major）的方式进行存储布局。可以想象成数据是一列一列地连续存储在共享内存中的，但地址却变成了离散的。

一个warp，32个线程，共分配了32个向量，各个线程访问自己向量第一个元素时，共32个bank，则：

方式2每个线程访问数据，访问的是第一列，地址是threadIdx.x*6+0，对应Bank的索引是threadIdx.x*6+0)%32，可以发现线程0和16（下图绿色），1和17（下图红色）带入计算的bank索引相同，一个bank同时被多个线程访问，出现Bank conflict。如下图

图2 方式2 在共享内存中bank划分

方式3每个线程访问数据，访问的是第一行，地址是threadIdx.x，对应Bank的索引是threadIdx.x%32，每个向量都存储在不同的bank中，对bank的访问是规则连续的。因而改变数据存储方式可以避免Bank conflict。如下图所，

图3 方式3 在共享内存中bank划分

方式3存储的缺点：数组元素被离散存储了，即图1（右）第一列，元素0到5的地址是不连续的，某些情况下对编程造成了很大的不便，

而方式2的优点正在于每个数组的元素0到5的地址，都是连续存储的，特别是该数组是一个矩阵时，会带来巨大的便利。在上面的例子中，如果Vector0其实是一个2行3列的矩阵，则对于方式2存储时，可以用指针将一维数组转换为二维数组的访问：

int (*ptr)[3] = (int (*)[3])Vector0;

然后就可以使用诸如ptr[1][2]的方式来替代Vector0进行访问，这将简化一维数组的索引计算问题。

因此方式2能避免Bank conflict会更方便。下面优化方式2

上面的例子中向量的长度为6，是一个偶数，只要长度为偶数，按照方式2存储就会引入Bank conflict，而只要是奇数，则并不会导致这种冲突。因而当数组长度为偶数时，只需要将共享内存的数组长度增加1变为奇数，然后只使用前面的偶数个元素即可：

 __shared__ int Vector1[32][6+1];

方式2数组为奇数时，如下图：绿色，红色错开，且其他黄色都错开了，不在同一个bank。

图4 优化的方式2 在共享内存中bank划分

这样当每个线程同时访问自己向量的第一个元素时，按方式2存储则每个线程访问的字地址将为：threadIdx.x*7+0，对应bank为(threadIdx.x*7+0)%32，就不会出现引入Bank conflict的问题。缺点是浪费了32个字的共享内存空间。

未优化的方式2，核函数：

// 行优先
__global__ void BankConflictRowMajor(const int M, const int N, float *A)
{
    __shared__ float shareA[32][6]; // 共享内存声明
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[threadIdx.x][i] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[threadIdx.x][i] = shareA[threadIdx.x][i] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[threadIdx.x][i];
    }
}

Nsight Compute测试结果：

This kernel has uncoalesced shared accesses resulting in a total of 16 excessive wavefronts (40% of the total 40 wavefronts). Check the L1 Wavefronts Shared Excessive table for the primary source locations

可知有16个bank冲突，正好是黄色红色绿色的总数量。

方式3的核函数：

// 列优先
__global__ void BankConflictColMajor(const int M, const int N, float *A)
{
    __shared__ float shareA[6][32]; // 共享内存声明
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[i][threadIdx.x] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[i][threadIdx.x] = shareA[i][threadIdx.x] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[i][threadIdx.x];
    }
}

结果得到：derived__memory_l1_wavefronts_shared_excessive 值为0，即没有bank conflict

方式2优化后的核函数：

// 行优先 padding
__global__ void BankConflictRowMajorPadding(const int M, const int N, float *A)
{
    __shared__ float shareA[32][6 + 1]; // 共享内存声明 padding 1 
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[threadIdx.x][i] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[threadIdx.x][i] = shareA[threadIdx.x][i] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[threadIdx.x][i];
    }
}

结果得到：derived__memory_l1_wavefronts_shared_excessive 值也为0，解决了bank conflict

总代码

demoBankConflict1D.cu

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

// 行优先
__global__ void BankConflictRowMajor(const int M, const int N, float *A)
{
    __shared__ float shareA[32][6]; // 共享内存声明
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[threadIdx.x][i] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[threadIdx.x][i] = shareA[threadIdx.x][i] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[threadIdx.x][i];
    }
}
// 行优先 padding
__global__ void BankConflictRowMajorPadding(const int M, const int N, float *A)
{
    __shared__ float shareA[32][6 + 1]; // 共享内存声明 padding 1 
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[threadIdx.x][i] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[threadIdx.x][i] = shareA[threadIdx.x][i] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[threadIdx.x][i];
    }
}
// 列优先
__global__ void BankConflictColMajor(const int M, const int N, float *A)
{
    __shared__ float shareA[6][32]; // 共享内存声明
    for (int i = 0; i < N; i++)     // 从global内存读取数据到share内存
    {
        shareA[i][threadIdx.x] = A[threadIdx.x * N + i];
    }
    for (int i = 0; i < N; i++) // 使用share内存的数据进行运算
    {
        shareA[i][threadIdx.x] = shareA[i][threadIdx.x] * 2;
    }
    for (int i = 0; i < N; i++)
    {
        A[threadIdx.x * N + i] = shareA[i][threadIdx.x];
    }
}

void viewMat(const float *p, const int r, const int c)
{
    for (int i = 0; i < r; ++i)
    {
        for (int j = 0; j < c; ++j)
        {
            std::cout << p[i * c + j] << " ";
        }
        std::cout << std::endl;
    }
}

int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);

    // setp1: 设置GPU设备
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("Using Device %d: %s\n", dev, deviceProp.name); // 设备信息
    cudaSetDevice(dev);

    // setp2: 初始化矩阵
    int M = 32;
    int N = 6;
    size_t nbytes = sizeof(float) * M * N;

    // 分配主机内存
    float *A;
    A = (float *)malloc(nbytes);

    // 在host上初始化数据
    for (int i = 0; i < M * N; ++i)
    {
        A[i] = i % 6;
    }

    std::cout << "Matrix A:" << std::endl;
    viewMat(A, M, N);

    // setp4: 分配GPU内存
    float *MatA;
    cudaMalloc((void **)&MatA, nbytes);
    // 将数据从host传入设备上
    cudaMemcpy(MatA, A, nbytes, cudaMemcpyHostToDevice);

    dim3 block(M);
    dim3 grid(1);
    BankConflictRowMajorPadding<<<grid, block>>>(M, N, MatA);
    cudaDeviceSynchronize();

    // setp6: 在主机中获取计算结果
    cudaMemcpy(A, MatA, nbytes, cudaMemcpyDeviceToHost);

    // 打印结果矩阵 C
    std::cout << "Result matrix A:" << std::endl;
    viewMat(A, M, N);

    // 释放设备全局内存
    cudaFree(MatA);
    // 释放host内存
    free(A);
    // 重置设备
    cudaDeviceReset();

    return 0;
}