用户
 找回密码
 立即注册
quanzhang100 该用户已被删除
发表于 2013-12-1 22:00:51
2086512
本帖最后由 quanzhang100 于 2013-12-1 22:04 编辑

我在学习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];      
}
其结果如下
Effective Bandwidth (GB/s, ECC enabled)
RoutineTesla M2050Tesla K20c
copy105.2136.0
copySharedMem104.6152.3
我的问题有2个
(1)为什么在k20上带宽会有提高,而在M2050上却下降了(我的理解是两个都会下降啊)?
(2)这里带宽提高了是不是计算时间也缩短了?如果是的话,对于类似于copy这样的kernel(不同于矩阵乘法和矩阵转置的kernel)是不是也可以使用share memory可以提高性能呢?
使用道具 举报 回复
发表于 2013-12-1 23:46:49
quanzhang100 发表于 2013-12-1 22:07
不能发链接 我把博文的名字发上来  An Efficient Matrix Transpose in CUDA C/C++

LZ您好:

我来先帮您发一下链接:
http://devblogs.nvidia.com/paral ... -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 反对 0
发表于 2013-12-1 22:07:58
不能发链接 我把博文的名字发上来  An Efficient Matrix Transpose in CUDA C/C++
使用道具 举报 回复 支持 反对
发表于 2013-12-1 23:56:16
关于您的两个问题:

1:我不清楚为什么会这样,手头尚无K20卡做验证,倘若真的如此,那么可能和kepler架构对长延迟操作的调度有关,但目前尚无任何信息。

2:因为文中的kernel主要就是copy和转置,并且这个等效的带宽就是按照时间计算的,所以这里按照带宽来衡量了效能。但如1:中所言,尚不知道为何会这样,所以无法继续向下分析。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-12-2 09:59:58
ice 发表于 2013-12-1 23:46
LZ您好:

我来先帮您发一下链接:

谢谢ice深夜解答哈
使用道具 举报 回复 支持 反对
发表于 2013-12-2 11:22:55
quanzhang100 发表于 2013-12-2 09:59
谢谢ice深夜解答哈

不客气的,还请其他感兴趣的网友/斑竹/原厂支持继续讨论~
使用道具 举报 回复 支持 反对
发表于 2013-12-2 13:12:33
LZ您好:

经过和横扫千军斑竹的讨论和测试,基本上在kepler架构下重现了您发现的问题。

以及,进一步研究发现,实际上并不需要使用shared memory,仅仅在读入和写出的部分之间添加一个__syncthreads(),(先写到kernel里面的临时变量,然后同步,然后写出),就可以达到类似于文中那样比直接copy提升10%以上的访存带宽的效果。

以及,并无任何官方资料对此现象负责,仅能推测这个和GPU内部的调度情况有关,加入的同步恰好优化了GPU内部的调度。

再无其他具体解释了。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-12-2 22:17:10
ice 发表于 2013-12-2 13:12
LZ您好:

经过和横扫千军斑竹的讨论和测试,基本上在kepler架构下重现了您发现的问题。

ice 版主你好,你的测试结果是K20上的吗?
我在GTX680上测试 share memory copy 带宽82GB/s    copy的带宽为135GB/s 并没提升,难道只是针对k20的卡才有提升?
使用道具 举报 回复 支持 反对
ice 来自手机
9#
发表于 2013-12-3 00:32:07
quanzhang100 发表于 2013-12-2 22:17
ice 版主你好,你的测试结果是K20上的吗?
我在GTX680上测试 share memory copy 带宽82GB/s    copy的带 ...

我们是在SM30的GT650Ti上测试的,照说应该和您的680有一致性,您的情况,我不清楚为何了。

以及,您的CUDA版本是?
使用道具 举报 回复 支持 反对
ice 来自手机
10#
发表于 2013-12-3 00:35:51
另外,您的shared memory copy的带宽似乎偏低。您没有误用带有bankconflict的转置kernel吧?或者您是debug模式测试的?
使用道具 举报 回复 支持 反对
12下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册