我在学习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()造成的,目前原因不明。
祝您好运~