CUDA-存储优化-例-矩阵转置

存储优化

重叠内存传输和计算,组团传输减少小块数据的频繁传输。

现代GPU对访存做了优化,使得不是严格的Coalescing Access也是可以接受的。但是

  • 避免凌乱无规律的访存
  • 避免一个线程访问连续的一段空间。

将global memory中的数据放入shared memory中,使得数据的位置相邻后写入global memory,此时相邻的threads就可以访问相邻的地址,满足Coalescing Access。shared memory的设计目的之一就是通过对它的编程,来规则化访存模式。

回忆:shared memory的架构是连续32 bits(即4 Bytes)的地址被分配到连续的bank中。见下图:

原图来自这里。里面有不少其他资源可参考。

对于bank冲突,注意理解地址的编号和bank的编号的不同

注意

Coalescing access是就global memory而言的。

存储优化,优化实列:矩阵转置

  1. 未优化的矩阵转置

    对global memory的读和写总有一个,对global存储的地址访问不是连续的,这就不能最大化coalescing access。看下图:

    每个block负责矩阵的一个子块(tile),所有子块并行执行。TILE_DIM=blockDIm.x

    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    __global__ void transpose(float* a, float* b){
    // 每个线程的id
    int idx = threadIdx.x + TILE_DIM * blockIdx.x;
    int idy = threadIdx.y + TILE_DIM * blockIdx.y;

    // 每个数据转置前后,在矩阵中的索引
    int index_a = idx + WIDTH * idy;
    int index_b = idy + HEIGHT * idx;

    // 转置操作
    b[index_b] = a[index_a];
    }

    所以:

  2. 使用shared memory优化的矩阵转置

    在写入global memory之前,先从global中将所读取的元素存入shared memory中,当shared memory中有了所有元素,将此时的shared memory转置一下,最后将结果再写(合并地)入global memory。

    如此一来对global的读和写,地址都是被连续访问的。看下图:

    实现:

    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    __global_ void transpose(float* a, float* b){

    __shared__ float tile[TILE_DIM][TILE_DIM];

    // 每个线程的id
    int idx = threadIdx.x + TILE_DIM * blockIdx.x;
    int idy = threadIdx.y + TILE_DIM * blockIdx.y;
    // idx 与 idy 线性组合 得到原矩阵的Index
    int index_in = idx + idy * WIDTH;

    // 每个数据转置前后,在矩阵中的索引
    idx = threadIdx.x + TILE_DIM * blockIdx.y;
    idy = threadIdx.y + TILE_DIM * blockIdx.x;
    // // idx 与 idy 线性组合 得到转置后矩阵的Index
    int index_out = idx + idy * HEIGHT;

    // 从a中写入到bank中,会产生冲突。(马上解决)
    tile[threadIdx.y][threadIdx.x] = a[index_in];

    // 等待这个block对应的tile中的 元素都有了之后,再执行下后续操作。
    __syncthreads();

    // 转置操作
    b[index_out] = tile[threadIdx.x][threadIdx.y];
    }

    在Shared memory中存在bank冲突,如何解决。

  3. 解决上述过程中的bank冲突

    对于上述实现过程中tile[threadIdx.y][threadIdx.x] = a[index_in];,如果tile的大小是16x16,那么这一句会产生16路的bank冲突。 !!!如何理解bank的编号!!!

    如何解决bank冲突,看下图:

    实现中只需更tile的定义为__shared__ float tile[TILE_DIM][TILE_DIM+1];,列中的数据存于相同的bank。

    如此一来,不论是对tile行或列的访存,都不会产生bank冲突。相同的颜色在不同的行和列。