50%的共享内存效率？
50%的共享内存效率？
提问于 2018-06-20 21:00:02
下面的代码使用共享内存执行平铺矩阵转置以提高性能。共享内存由1列填充，以避免32x32线程块的银行冲突。
__global__ void transpose_tiled_padded(float *A, float *B, int n)
{
int i_in = blockDim.x*blockIdx.x + threadIdx.x;
int j_in = blockDim.y*blockIdx.y + threadIdx.y;
int i_out = blockDim.x*blockIdx.y + threadIdx.x;
int j_out = blockDim.y*blockIdx.x + threadIdx.y;
extern __shared__ float tile[];
// coalesced read of A rows to (padded) shared tile column (transpose)
tile[threadIdx.y + threadIdx.x*(blockDim.y+1)] = A[i_in + j_in*n];
__syncthreads();
// coalesced write from (padded) shared tile column to B rows
B[i_out + j_out*n] = tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];
}运行这段代码，我将在NVIDIA视觉分析器中获得100%的共享内存效率，正如我所预期的那样。但是，当我使用16x16线程块运行它时，效率只有50%。为什么会这样呢？据我所知，在这个布局中，翘曲中的任何线程都不会从同一个银行读取。还是我搞错了？
回答 1
Stack Overflow用户
回答已采纳
发布于 2018-06-21 00:07:11
是的，你错了。
考虑到在16x16块中对warp 0的(读取)访问：
tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
"index"以下是经纱中每一个线程的相关计算：
warp lane: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 23 25 26 27 28 29 30 31
threadIdx.x: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
threadIdx.y: 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
"index": 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32
bank: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 0所以我们看到，对于这个偏差，第一个线程和最后一个线程都是从银行0读取的。这将导致双向银行冲突、双向序列化和50%的效率.
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/50956939
复制相关文章
相似问题
腾讯云开发者
