合并的全局内存访问中读写的理解-优快云博客

__global__ void transposeSmem(float *out, float *in, int nx, int ny) {
    // static shared memory
    __shared__ float tile[BDIMY][BDIMX];

    // coordinate in original matrix
    unsigned int ix, iy, ti, to;
    ix = blockIdx.x * blockDim.x + threadIdx.x;
    iy = blockIdx.y * blockDim.y + threadIdx.y;

    // linear global memory index for original matrix
    ti = iy * nx + ix;

    // thread index in transposed block

    unsigned int bidx, irow, icol;
    bidx = threadIdx.y * blockDim.x + threadIdx.x;
    irow = bidx / blockDim.y;
    icol = bidx % blockDim.y;

    // coordinate in transposed matrix
    unsigned int ox = blockIdx.y * blockDim.y + icol;
    unsigned int oy = blockIdx.x * blockDim.x + irow;

    // linear global memory index for transposed matrix
    to = oy * ny + ox;

    // transpose with boundary test
    if (ix < nx && iy < ny)
    {
        // load data from global memory to shared memory
        tile[threadIdx.y][threadIdx.x] = in[ti];

        // thread synchronization
        __syncthreads();

        // store data to global memory from shared memory
        out[to] = tile[icol][irow];
    }
}

上面是书上的源码修改了一下，添加独立字段ox oy。

个人理解：
第一个连续读取任务到__syncthreads()就结束了，接下来单独考虑第二个任务的连续写问题。
如果直接用threadIdx交换xy来写，那么因为原始线程是x方向上连续的，那么写就不会连续。
因为此时 ox = iy oy = ix, 而x是连续的，那么oy变成连续的，ox不是连续的，这就给写带来交叉访问，我们希望ox上连续。

如果直接用threadIdx交换xy来写，那么就是最初的版本：

out[ix * ny + iy] = in[iy * nx + ix];

AI总结：
__syncthreads() 把"读进来"和"写出去"切成两个独立阶段；装载阶段靠 tx 天然合并；卸载阶段因为 out 的行主序要求"ox 必须跟着线程号快变"，而线程身份 (tx,ty) 无法满足（会退化成 oy 快变 = 交叉访问），所以必须重新派发任务——icol = bidx % BDIMY 就是为此造出的那个"快变的 ox"