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

本文介绍了如何在CUDA中安排2D共享内存的处理方法，对大家解决问题具有一定的参考价值，需要的朋友们下面随着小编来一起学习吧！ 问题描述 29岁程序员，3月因学历无情被辞！ 我一直使用线性共享内存（加载，存储，访问邻居），但我做了一个简单的测试在2D中研究银行冲突，结果困惑了我。 下一个代码将数据从一维全局内存数组读取到共享内存，并将其从共享内存复制回全局内存。 __ global__ void update（int * gIn，int * gOut，int w）{ //共享内存空间 __shared__ int shData [16] [16]; //从threadIdx / BlockIdx到数据位置的映射 int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y; //计算一维数组中的全局id int gid = x + y * w; //加载共享内存 shData [threadIdx.x] [threadIdx.y] = gIn [gid]; //同步线程不是真的需要，但为了方便 __syncthreads（）; //将数据写回全局内存 gOut [gid] = shData [threadIdx.x] [threadIdx.y];可视化分析工具报告共享内存冲突。 下一个代码避免冲突（只显示差异） //加载共享内存 shData [threadIdx.y ] [threadIdx.x] = gIn [gid]; //将数据写回全局内存 gOut [gid] = shData [threadIdx.y] [threadIdx.x]; 这种行为让我困惑，因为在大规模并行处理器编程。我们可以阅读以下实践方法：根据行主要约定，将C和CUDA中的矩阵元素放置到线性寻址的位置中：这与共享内存段？或线程索引？也许我错过了什么？ 内核配置如下： // kernel configuration dim3 dimBlock = dim3（16，16，1）; dim3 dimGrid = dim3（64，64）; //使用16x16线程启动64x64块的网格 - > 1048576 threads update<<< dimgrid，dimBlock>>>>（d_input，d_output，1024）; 提前感谢。方案是的，共享内存按照您所期望的顺序排列。所以你的[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 是在pre-Fermi硬件上的16个32位共享存储器组，每个列中的每个整数条目映射到一个共享存储器组。那么，如何与您选择的索引方案相互作用？ 要记住的一点是，一个块中的线程编号为等于列主顺序结构的x维度是最快的变化，其次是y，其次是z）。因此，当您使用此索引方案时： shData [threadIdx.x] [threadIdx.y] shData [threadIdx.y] [threadIdx.x] 线程将从相同的行读取，这意味着从16个不同的共享存储器组中的每一个读取，发生冲突。 I've always worked with linear shared memory (load, store, access neighbours) but I've made a simple test in 2D to study bank conflicts which results have confused me.The next code read data from one dimensional global memory array to shared memory and copy it back from shared memory to global memory.__global__ void update(int* gIn, int* gOut, int w) {// shared memory space__shared__ int shData[16][16];// map from threadIdx/BlockIdx to data positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;// calculate the global id into the one dimensional arrayint gid = x + y * w;// load shared memoryshData[threadIdx.x][threadIdx.y] = gIn[gid];// synchronize threads not really needed but keep it for convenience__syncthreads();// write data back to global memorygOut[gid] = shData[threadIdx.x][threadIdx.y];}The visual profiler reported conflicts in shared memory. The next code avoid thouse conflicts (only show the differences)// load shared memoryshData[threadIdx.y][threadIdx.x] = gIn[gid];// write data back to global memorygOut[gid] = shData[threadIdx.y][threadIdx.x];This behavior has confused me because in Programming Massively Parallel Processors. A Hands-on approach we can read:Is this related to shared memory arragment? or with threads indexes? Maybe am I missing something?The kernel configuration is as follow:// kernel configurationdim3 dimBlock = dim3 ( 16, 16, 1 );dim3 dimGrid = dim3 ( 64, 64 );// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threadsupdate<<<dimGrid, dimBlock>>>(d_input, d_output, 1024);Thanks in advance. 解决方案 Yes, shared memory is arranged in row-major order as you expected. So your [16][16] array is stored row wise, something like this: bank0 .... bank15row 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 15Because there are 16 32 bit shared memory banks on pre-Fermi hardware, every integer entry in each column maps onto one shared memory bank. So how does that interact with your choice of indexing scheme?The thing to keep in mind is that threads within a block are numbered in the equivalent of column major order (technically the x dimension of the structure is the fastest varying, followed by y, followed by z). So when you use this indexing scheme:shData[threadIdx.x][threadIdx.y]threads within a half-warp will be reading from the same column, which implies reading from the same shared memory bank, and bank conflicts will occur. When you use the opposite scheme:shData[threadIdx.y][threadIdx.x]threads within the same half-warp will be reading from the same row, which implies reading from each of the 16 different shared memory banks, no conflicts occur. 这篇关于如何在CUDA中安排2D共享内存的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持！ 上岸，阿里云！