基于CUDA的GPU优化方法

初学CUDA，往往拿到代码无从下手，也没有什么明确的思路。我想有必要把前人的经验总结拿出来，便于后来者更快掌握这门技术。

对于block和thread的分配问题，有这么一个技巧，每个block里面的thread个数最好是32的倍数，因为，这样可以让计算效率更高，促进memory coalescing。其实，每个grid里面block的dimension维度和size数量，以及每个block里面的thread的dimension维度和size数量，都是很重要的。采用合适的维度可以更方便的将并行问题映射到CUDA架构上，但是，对性能不会有太大改进。所以，size才是最重要的。其实，访问延迟latency和occupancy占有率，都依赖于每个multiprocessor中的active wrap的数量，而active wrap的数量，又依赖于register和share memory的使用情况。首先，grid中block的数目要大于multiprocessor的数目，以保证每个multiprocessor里面最少有一个block在执行，而且，最好有几个active block，使得blocks不要等着__syncthreads()，而是占用了hardware。其次，block里面的thread的数目也很重要。对于1.0和1.1的设备来讲，如果一个kernel里面block的大小为512个thread，那么，occupancy为512/768=66%，并且一个multiprocessor中只有一个active block，然而，如果block里面的thread为256个thread，那么，768/256=3，是整数，因此，occupancy