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而言的。

存储优化，优化实列：矩阵转置

未优化的矩阵转置

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

每个block负责矩阵的一个子块（tile），所有子块并行执行。TILE_DIM=blockDIm.x。

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

所以：

使用shared memory优化的矩阵转置

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

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

实现：

__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冲突，如何解决。

解决上述过程中的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冲突。相同的颜色在不同的行和列。