1. 将串行代码并行化这里说简单也简单，说难也难，比如，对比 for 循环， 就可以很简单的直接拆开，并行但是，如果是迭代算法，比如 Gauss-Seidel 迭代求解，那么，就需要整理 Gauss-Seidel 算法，提起共同项，然后并行 2. 尽量减少 Host 和 Device 之间的数据拷贝拷贝一次，尽量多用3. 在配置 kernel 的时候，分配合理的 thread 个数和 block 个数，使得 device 的使用达到最大化，充分利用硬件资源切记：thread 不是越多越好4. 保证 global memory 在访问的时候，memory coalesced5. 尽可能的用 share memory 取代用 global memory 对数据进行访问这个道 理，你懂的6. 同一个 wrap 中，注意尽量减少分支同一 wrap 内的 thread 尽量做同样的 事情7. memory coalescing，保证内存融合因为 global memory 在 CC 为 1.x 上是 按照 half wrap 进行访问读写的，而在 2.x 上是按照 wrap 进行访问读写的在 显存中，有多个存储器控制器，负责对显存的读写，因此，一定要注意存储器 控制器的负载均衡问题。
每一个存储器控制器所控制的那片显存中的地址空间 称为一个分区连续的 256Byte 数据位于同一个分区，相邻的另一组 256Byte 数据位于另一个分区访问 global memory 就是要让所有的分区同时工作合 并访问就是要求同一 half-wrap 中的 thread 按照一定 byte 长度访问对齐的段 在 1.0 和 1.1 上，half-wrap 中的第 k 个 thread 必须访问段里的第 k 个字，并且half-wrap 访问的首地址必须是字长的 16 倍，这是因为 1.0 和 1.1 按照 half- wrap 进行访问 global memory，如果访问的是 32bit 字，比如说一个 float，那 么 half-wrap 总共访问就需要 16 个 float 长，因此，每个 half-wrap 的访问首地 址必须是字长的 16 倍1.0 和 1.x 只支持对 32bit、64bit 和 128bit 的合并访问， 如果不能合并访问，就会串行 16 次1.2 和 1.3 改进了 1.0 和 1.1 的访问要求， 引进了断长的概念，与 1.0 和 1.1 上的端对齐长度概念不同，支持 8bit-段长 32Byte、16bit-段长 64Byte、32bit-64bit-128bit-段长 128Byte 的合并访问。
对 1.2 和 1.3 而言，只要 half-wrap 访问的数据在同一段中，就是合并访问，不再 像 1.0 和 1.1 那样，非要按照顺序一次访问才算合并访问如果访问的数据首 地址没有按照段长对齐，那么 half-wrap 的数据访问会分两次进行访问，多访 问的数据会被丢弃掉所以，下面的情况就很容易理解：对 1.0 和 1.1，如果 thread 的 ID 与访问的数据地址不是顺序对应的，而是存在交叉访问，即：没有 与段对齐，那么，就会 16 次串行访问，而对 1.2 和 1.3 来讲，会判断这 half- wrap 所访问的数据是不是在同一个 128Byte 的段上，如果是，则一次访问即可， 否则，如果 half-wrap 访问地址连续，但横跨两个 128Byte，则会产生两次 传输，一个 64Byte，一个 32Byte当然，有时还要考虑 wrap 的 ID 的奇偶性1.2 和 1.3 放宽了对合并访问的条件，最快的情况下的带宽是最好的情况下的带宽 的 1/2，然而，如果 half-wrap 中的连续 thread 访问的显存地址相互间有一定的 间隔时，性能就会灰常差比如，half-wrap 按列访问矩阵元素，如果 thread 的 id 访问 2*id 的地址空间数据，那么，半个 wrap 访问的数据刚好是 128Byte，一次访问可以搞定，但是，有一半数据会丢失，所以，也表示浪费 了带宽，这一点一定要注意。
如果不是 2 倍，而是 3 倍、4 倍，那么，有效带 宽继续下降在程序优化时，可以使用 share memory 来避免间隔访问显存8. bank conflict，bank 冲突先说一下，share memory 在没有 bank conflict 情况下，访问速度是 global 和 local 的 100 倍呢，你懂的类似 global memory 的分区，share memory 进行了 bank 划分如果 half-wrap 内的很多 thread 同时要求访问同一个 bank，那么就是 bank conflict，这时，硬件就会将 这些访问请求划分为独立的请求，然后再执行访问但是，如果 half-wrap 内 所有 thread 都访问同一个 bank，那么会产生一次 broadcast 广播，只需要一 次就可以相应所有访问的请求每个 bank 宽度长度为 32bit对于 1.x 来讲， 一个 SM 中的 share memory 被划分为 16 个 bank，而 2.x 是 32 个 bank1.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 分配过多的 register4. 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 block，但是 6 个 active block，就会导致总共 thread 个数为 128*6*12 个，严重超过了 8192， 所以不能为 6，得为 5，因为 128*5>log2(n)), (i%n)=(icase 1:break;...case 31:break;}上面这个例子，则不会发生 divergence，因为控制条件刚好和 wrap 里面的 thread 相对应其实，有时，compiler 会采用 branch predication 分支预测来打开 loop 循环或 者优化 if 和 switch 语句， 这时，wrap 就不会出现 divergence 了。
在写 code 时，我们也可以自己采用#pragma uroll 来打开 loop 循环在使用 branch predication 时，所有指令都将会执行，其实，只有预测正确的真正的执行了， 而预测错误的，其实就是 thread，不会去读取该 instruction 的地址和数据，也 根本不会写结果其实，编译器做分制预测，是有条件的，只有分支条件下的 指令 instruction 的个数小于等于某个阈值的时候，才会做分支预测 branch predication如果编译器觉得可能会产生多个 divergent wrap，那么阈值为 7， 否则为 4这里很不理解 7 和 4 是怎么来的）10. 在 loop 循环的 counter，尽量用 signed integer，不要用 unsigned integer比如：for(i = 0; i < n; i++) {out[i] = in[offset+stride*i];} 这里呢， stride*i 可以会超过 32 位 integer 的范围，如果 i 被声明为 unsigned，那么 stride*i 这个溢出语句就会阻止编译器做一些优化，比如 strength reduction。
相 反，如果声明为 signed，也没有溢出语句时，编译器会对很多地方做优化所 以，loop counter 尽量设置为 int，而不是 unsigned int11. 在 1.3 及其以上的 device 上，才支持 double-precision floating-point values，即：64 位双精度浮点运算当使用 double 时，在编译器选项里面添 加：-arch=sm_1312. 还有一点需要注意，如果 A、B、C 都是 float，那么 A+(B+C)并不一定等于(A+B)+C13. 先看下面两个语句：float a; a = a * 1.02;对于 1.2 及其以下的 device 来讲，或者 1.3 及其以上 device，但是没有打开支 持 double 运算的选项，那么，由于不支持 double，所以，1.02*a 这个乘积是 一个 float；对于 1.3 及其以上的 device 来讲，如果打开了支持 double 运算的选项，那么， a*1.02 是一个 double，而将乘积赋值给 a，这个结果是 float，所以，是先做了 从 float 到 doub。