cuda优化策略.doc

1. 将串行代码并行化。这里说简单也简单，说难也难，比如，对比for循环，就可以很简单的直接拆开，并行。但是，如果是迭代算法，比如Gauss-Seidel迭代求解，那么，就需要整理Gauss-Seidel算法，提起共同项，然后并行。 
conflict是不一样的。，多个thread访问同一个bank，就会出现bank conflict，half-wrap内所有thread访问同一个bank除外。但是，，多个thread访问同一个bank已经不再是bank conflict了。比如：
__shared__ char Sdata[32];
char  data = Sdata[BaseIndex+tid];
conflict，因为，0~3thread访问同一个bank，4~7访问同一个bank，类推，这种情况属于4-way bank conflict。但是，，这种情况已经不是bank conflict了，，牛吧，哈哈。这里要多看看矩阵乘积和矩阵转置例子中的share memory的使用，如何保证memory coalescing和避免bank conflict的。
9. texture memory是有cache的，但是，如果同一个wrap内的thread的访问地址很近的话，那么性能更高。 
 
以下是要注意的：
1. ，L1 cache比texture cache具有更高的数据带宽。所以，看着使用哈。
2. 对global memory的访问，，容易造成memory uncoalescing，，容易造成bandwidth waste。 ，，除了多了L1 cache，没有其他的特别之处。
3.  采用-maxrregcount=N阻止complier分配过多的register。
4.  occupancy是每个multiprocessor中active wrap的数目与可能active wrap的最大数目的比值。higher occupancy并不意味着higher performance，因为毕竟有一个点，超过这个点，再高的occupancy也不再提高性能了。
5. 影响occupancy的一个因素，就是register的使用量。比如，，每个multiprocessor最多有8192个register，而最多的simultaneous thread个数为768个，那么对于一个multiprocessor，如果occupancy达到100%的话，每个thread最多可以分配10个register。另外，，一个kernel里面的一个block有128个thread，每个thread使用register个数为12，那么，occupancy为83%，这是因为一个block有128个thread，则，由于multiprocessor里面最大的simultaneous thread为768，根据这个数目计算，最多同时有6个active