实现算法并行是第一步，但是后续的优化会大幅度影响算法运行速度，这里讲一些可以后续优化的点。GPU优化比较麻烦的是，不一定优化会起作用，因此要把所有方法都试一遍，才能知道是（优化前，优化A，优化B ...）那种最适合自己的问题

参考学习视频

链接：https://pan.baidu.com/s/1NlFZS2EHdY86z09AlR41Jg 
提取码：1234

1.warm up

第一次运行程序可能由于GPU没有完全启动，导致时间过慢，可以在程序开始中加一个空转函数，这个函数不计时，让GPU完全运转。

如果是科研的话，可以取第二次以后的实验结果，基本比较稳定。

2.内存速度

常量内存constant > 共享内存shared > 全局内存global

常量内存类似于c++ const，但是太小了，很多时候要经常访问的数据量很大，放不进去，有时候放进去，有时候放进去了，但是速度也没有明显提升。

这里说的明显是多次运行程序没有明显提升，例如优化前5次：10.9s，10.1s，10.3s，10.2s，10.4s。优化后5次：10.8s，10.2s，10.1s，10.0s，10.1s。第一次时间长是没有warm up，应该从第2次计算。这个优化效果其实不太明显，如果放在科研里，不能算加速，但是工程中可能算大概5%的加速吧

下面一个简单的例子，展示共享内存的简单用法，x拷贝到共享内存中，做乘2处理，实际可能更加复杂，然后拷贝给y。共享内存还可以动态分配，在<<<>>>的第三个参数中。速度有可能会更快，可以试试，也是不一定。

__global__ func(int *x, int *y){
    __shared__ int s[10];
    int id = threadIdx.x;
    if(id >= N)
        return;
    s[id] = x[id];
    //一些处理后，将s拷贝给y
    s[id] = s[id] * 2;
    __syncthreads ();
    y[id] = s[id];
}

局部内存local > 共享内存shared > 全局内存global

__global__ func(int *x, int *y, int *z){
    __shared__ int a, b, c;
    int m, n, p;
    n = p;//1
    a = b;//2
    a = *x;//3
    *x = *y;//4
}

运行速度为1，2，3，4代码中比较直观，不再赘述。

全局内存的问题：内存合并

 每次读取全局内存的时候，是安装一个内存块读的，例如图中虚线框，如果图中连续4个红色点，在一个内存读取块中（这里也假设一个内存读取块大小为5），所以只需要读取一次，然后后边黑色点距离比较远，要读取两次，那显然第一个速度要更快。

但是实际遇到问题中，还必须用第二种的方式计算，所以有时候也没办法，或者有更好的优化方法要学习。

原子操作

原子操作会保证正确性，但降低运行速度，但有时又不一定会降低运行速度。

例如计算出的结果中0比较多，也就是结果比较稀疏，如果最后几十万个数据归约，有时候效果并不一定好。这时判断if(num > 0) 原子加，可能更好，虽然理论上if会导致条件分支（如下图所示，条件分支导致一些线程要等待其他线程），导致性能变差，但是实际效果可能会不错，要试试。还是那句话，GPU优化要试，不一定那种比较好，《cuda c编程》也说了类似观点。

扫描算法，柱状图，排序

扫描算法实现比较复杂，可以直接用thrust库，thrust库还有一些例如排序算法，如果数据不大，例如几万以内的就不要用扫描排序了，可能还会更慢，这里慎重使用GPU优化，一般数据量真的没几千万，几个亿的，提升效果其实不明显的 。

柱状图类似扫描。

流水线

三个操作

H2D：CPU传到GPU

KER：GPU上的计算

D2H：GPU传到CPU

Stream1：H2D -> KER -> D2H

Stream2：             H2D -> KER -> D2H

Stream3：                          H2D -> KER -> D2H

<<<>>>中第四个参数，应用于频繁需要CPU，GPU之间交流的程序，如果需要CPU串行的计算很多，则需要使用流水线，但是如果串行不多，可以把数据全部放到GPU处理，虽然GPU处理串行能力较差，但是如果数据不多，影响就会很小，而且减少了CPU和GPU的数据交换和同步，运行时间也有大幅度缩减。