下载此文档

cuda优化策略.doc


文档分类:IT计算机 | 页数:约8页 举报非法文档有奖
1/8
下载提示
  • 1.该资料是网友上传的,本站提供全文预览,预览什么样,下载就什么样。
  • 2.下载该文档所得收入归上传者、原创者。
  • 3.下载的文档,不会出现我们的网址水印。
1/8 下载此文档
文档列表 文档介绍
1. 将串行代码并行化。这里说简单也简单,说难也难,比如,对比for循环,就可以很简单的直接拆开,并行。但是,如果是迭代算法,比如Gauss-Seidel迭代求解,那么,就需要整理Gauss-Seidel算法,提起共同项,然后并行。 
21.x的bank conflict和2.x的bank conflict是不一样的。对1.x来讲,多个thread访问同一个bank,就会出现bank conflict,half-wrap内所有thread访问同一个bank除外。但是,对2.x来讲,多个thread访问同一个bank已经不再是bank conflict了。比如:
__shared__ char Sdata[32];
char  data = Sdata[BaseIndex+tid];
在1.x上属于bank conflict,因为,0~3thread访问同一个bank,4~7访问同一个bank,类推,这种情况属于4-way bank conflict。但是,对于2.x来讲,这种情况已经不是bank conflict了,以为2.x采用了broadcast机制,牛吧,哈哈。这里要多看看矩阵乘积和矩阵转置例子中的share memory的使用,如何保证memory coalescing和避免bank conflict的。
9. texture memory是有cache的,但是,如果同一个wrap内的thread的访问地址很近的话,那么性能更高。 
 
以下是要注意的:
1. 在2.x的CC上,L1 cache比texture cache具有更高的数据带宽。所以,看着使用哈。
2. 对global memory的访问,1.0和1.1的设备,容易造成memory uncoalescing,而1.2和1.3的设备,容易造成bandwidth waste。 而对2.x的设备而言,相比1.2和1.3,除了多了L1 cache,没有其他的特别之处。
3.  采用-maxrregcount=N阻止complier分配过多的register。
4.  occupancy是每个multiprocessor中active wrap的数目与可能active wrap的最大数目的比值。higher occupancy并不意味着higher performance,因为毕竟有一个点,超过这个点,再高的occupancy也不再提高性能了。
5. 影响occupancy的一个因素,就是register的使用量。比如,对于1.0和1.1的device来讲,每个multiprocessor最多有8192个register,而最多的simultaneous thread个数为768个,那么对于一个multiprocessor,如果occupancy达到100%的话,每个thread最多可以分配10个register。另外,如果在1.0和1.1上,一个kernel里面的一个block有128个thread,每个thread使用register个数为12,那么,occupancy为83%,这是因为一个block有128个thread,则,由于multiprocessor里面最大的simultaneous thread为768,根据这个数目计算,最多同时有6个active

内容来自淘豆网www.taodocs.com转载请标明出处.