GPU编程个人入门（中）

实现算法并行是第一步，但是后续的优化会大幅度影响算法运行速度，这里讲一些可以后续优化的点。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拷贝给ys[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;//1a = b;//2a = *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的数据交换和同步，运行时间也有大幅度缩减。但是目前没有实际操作过这个。

shared线程发散

这个实际使用真的不好解决，除了归约，目前没见过其他应用，有待进一步学习

并行那些部分，需要很多优化么？

先用计时器计算算法每个模块的运行时间，并行运行时间最长的几个模块，如果有需要则全部并行，然后再去想优化，因为优化幅度20%，30%，可能不如多并行几个耗时长的模块来的有效。就算是O(n*n)->O(logn)也是可以的

block,thread分配数量？

以1080ti为例，每个块中线程数量最多1024，block的分配数量可以非常多，因此，每个block中的块测试32，64，128，256，512，1024分别测试，不一定哪一个好，但是会有一个最好的结果，还有就是一维度，二维度，都实验。例如并行2048个计算使用block（8），thread（16*16）。