如何在CUDA中安排2D共享内存

Question

我一直在使用线性共享内存(加载,存储,访问邻居),但我在2D中进行了一个简单的测试来研究银行冲突,结果让我感到困惑.

下一个代码将数据从一维全局存储器阵列读取到共享存储器,并将其从共享存储器复制回全局存储器.

__global__ void update(int* gIn, int* gOut, int w) {

// shared memory space
__shared__ int shData[16][16];
// map from threadIdx/BlockIdx to data position
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
// calculate the global id into the one dimensional array
int gid = x + y * w;

// load shared memory
shData[threadIdx.x][threadIdx.y] = gIn[gid];
// synchronize threads not really needed but keep it for convenience
__syncthreads();
// write data back to global memory
gOut[gid] = shData[threadIdx.x][threadIdx.y];
}

可视化分析器报告共享内存中的冲突.下一个代码避免了冲突(只显示差异)

// load shared memory
shData[threadIdx.y][threadIdx.x] = gIn[gid];

// write data back to global memory
gOut[gid] = shData[threadIdx.y][threadIdx.x];

这种行为让我很困惑,因为在编程大规模并行处理器中.我们可以阅读的实践方法:

根据行主要约定,将C和CUDA中的矩阵元素放入线性寻址的位置.也就是说,首先将矩阵的行0的元素按顺序放置到连续的位置中.

这与共享内存碎片有关吗？还是与线程索引？也许我错过了什么？

内核配置如下:

// kernel configuration
dim3 dimBlock  = dim3 ( 16, 16, 1 );
dim3 dimGrid   = dim3 ( 64, 64 );
// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threads
update<<<dimGrid, dimBlock>>>(d_input, d_output, 1024);

提前致谢.

Answer 1

是的,共享内存按行按主要顺序排列.所以你的[16] [16]数组是按行存储的,如下所示:

       bank0 .... bank15
row 0  [ 0   .... 15  ]
    1  [ 16  .... 31  ]
    2  [ 32  .... 47  ]
    3  [ 48  .... 63  ]
    4  [ 64  .... 79  ]
    5  [ 80  .... 95  ]
    6  [ 96  .... 111 ]
    7  [ 112 .... 127 ]
    8  [ 128 .... 143 ]
    9  [ 144 .... 159 ]
    10 [ 160 .... 175 ]
    11 [ 176 .... 191 ]
    12 [ 192 .... 207 ]
    13 [ 208 .... 223 ]
    14 [ 224 .... 239 ]
    15 [ 240 .... 255 ]
       col 0 .... col 15

由于前费米硬件上有16个32位共享存储区,因此每列中的每个整数条目都映射到一个共享存储区.那么它如何与您选择的索引方案相互作用？

需要记住的是,块中的线程编号与列主要顺序相当(技术上,结构的x维度变化最快,后跟y,后跟z).所以当你使用这个索引方案时:

shData[threadIdx.x][threadIdx.y]

半warp中的线程将从同一列读取,这意味着从同一共享存储库读取,并且将发生存储体冲突.当您使用相反的方案时:

shData[threadIdx.y][threadIdx.x]

同一半经线内的线程将从同一行读取,这意味着从16个不同的共享存储体中的每一个读取,不会发生冲突.