我在学习Mark harris 的博文
文中对memory copy的例子中，一个kernel是直接拷贝，读取和写入都是coalesced的，另外一个kernel在其中加入了share memory，发现有share memory的kernel在K20上带宽提高了。
kernel 1：无share memory
global void copy(float *odata, const float *idata){
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j < TILE_DIM; j+= BLOCK_ROWS)
odata[(y+j)*width + x = idata[(y+j)*width + x];
}
kernel 2：有share memory
global void copySharedMem(float *odata, const float *idata)
{
shared float tile[TILE_DIM * TILE_DIM];
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x = idata[(y+j)*width + x];
__syncthreads();
for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
odata[(y+j)*width + x = tile[(threadIdx.y+j)*TILE_DIM + threadIdx.x];
}
其结果如下[table=1]
[tr][td][/td][td=2,1]
Effective Bandwidth (GB/s, ECC enabled)
[/td][/tr]
[tr][td]Routine[/td][td]Tesla M2050[/td][td]Tesla K20c[/td][/tr]
[tr][td]copy[/td][td]105.2[/td][td]136.0[/td][/tr]
[tr][td]copySharedMem[/td][td]104.6[/td][td]152.3[/td][/tr]
[/table]我的问题有2个
（1）为什么在k20上带宽会有提高，而在M2050上却下降了（我的理解是两个都会下降啊）？
（2）这里带宽提高了是不是计算时间也缩短了？如果是的话，对于类似于copy这样的kernel（不同于矩阵乘法和矩阵转置的kernel）是不是也可以使用share memory可以提高性能呢？
不能发链接 我把博文的名字发上来 An Efficient Matrix Transpose in CUDA C/C++
LZ您好：
我来先帮您发一下链接：
http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/
大致参看了该博文的内容，大致是将shared memory在矩阵转置的算法中的使用。
原始的copy的例子是用来衡量两款GPU的copy速度的。
然后使用了没有任何优化的转置kernel，这里面有大量的非合并访问，得到了极低的访存带宽。
之后使用了存在bank conflict的shared memory版本的kernel，取得了一定的提速。
再之后为了验证此时低于直接copy的速度是否已经达到了使用shared memory 的极限，用一个使用了shared memory的copy kernel作为对比，此时发现差异还很大。
最后使用了消除bank conflict的转置kernel，此时达到了接近shared memory版本copy kernel的访存带宽。
原文作者以这样的思维流程阐述了如何使用shared memory来解决非合并访问（其实就是将shared memory作为缓冲，利用线程协作合并读入或者写出。），并消除bank conflict。
但是如您所述，文中shared memory版本的copy kernel和最终优化版的转置 kernel的访存带宽在K20c上都是高于一开始的直接copy的带宽的，这一点十分费解，而且作者也并没有解释原因。
我目前并不知道此处的问题何在，暂无法答复您。
请其他人予以补充。
关于您的两个问题：
1：我不清楚为什么会这样，手头尚无K20卡做验证，倘若真的如此，那么可能和kepler架构对长延迟操作的调度有关，但目前尚无任何信息。
2：因为文中的kernel主要就是copy和转置，并且这个等效的带宽就是按照时间计算的，所以这里按照带宽来衡量了效能。但如1：中所言，尚不知道为何会这样，所以无法继续向下分析。
大致如此，祝您好运~
谢谢ice深夜解答哈
不客气的，还请其他感兴趣的网友/斑竹/原厂支持继续讨论~
LZ您好：
经过和横扫千军斑竹的讨论和测试，基本上在kepler架构下重现了您发现的问题。
以及，进一步研究发现，实际上并不需要使用shared memory，仅仅在读入和写出的部分之间添加一个__syncthreads()，（先写到kernel里面的临时变量，然后同步，然后写出），就可以达到类似于文中那样比直接copy提升10%以上的访存带宽的效果。
以及，并无任何官方资料对此现象负责，仅能推测这个和GPU内部的调度情况有关，加入的同步恰好优化了GPU内部的调度。
再无其他具体解释了。
祝您好运~
ice 版主你好，你的测试结果是K20上的吗？
我在GTX680上测试 share memory copy 带宽82GB/s copy的带宽为135GB/s 并没提升，难道只是针对k20的卡才有提升？
[
我们是在SM30的GT650Ti上测试的，照说应该和您的680有一致性，您的情况，我不清楚为何了。
以及，您的CUDA版本是？
另外，您的shared memory copy的带宽似乎偏低。您没有误用带有bankconflict的转置kernel吧？或者您是debug模式测试的？
以及补充下，楼主应当仔细看下7#的详细说法，
我们是直接读取到寄存器，然后同步，然后写入的。
并无直接读写shared memory的过程，实际上根据实践，直接同步一次就能取得速度提高。
请您详细看下。
ice 不好意思，我不仔细，我是在debug下跑的，在release版本下，确实稍有提高。谢谢哈
嗯嗯，不客气的，之前见过有人反映说使用shared memory的矩阵乘法要比不使用还慢，后来发现是在debug模式下，所以此次也如此猜测了。
以及，实际上这个提高是__syncthreads()造成的，目前原因不明。
祝您好运~