cuda 本地内存使用_CUDA学习笔记（二）

一、CUDA代码的高效策略

1.高效公式

最大化计算强度：Math/Memory ，即:数学计算量/每个线程的内存最大化每个线程的计算量；

最小化每个线程的内存读取速度；每个线程读取的数据量少

每个线程读取的速度快内地内存>共享内存>>全局内存

合并全局内存

2.合并全局内存(按顺序读取的方式最好)

3.避免线程发散

线程发散：同一个线程块中的线程执行不同内容的代码

(1)kernel中做条件判断

(2)循环长度不一

二、Kernel加载方式

(1)查看本机参数

注：Kernel的加载中，自定义的线程数、线程块的数量等都不要超过系统本身的设定

(2)Kernel加载1D、2D、3D模式

(3)Kernel函数的关键字

三、CUDA中的各种内存的代码使用

1.本地内存(变量)

2.全局内存(变量)

3.共享内存(变量)

四、CUDA同步操作

1.原子操作

(1)解决的问题

对于有很多线程需要同时读取或写入相同的内存时，保证同一时间只有一个线程能进行操作

(2)特点只支持某些运算(加、减、最小值、异或运算等，不支持求余和求幂等)和数据类型(整型)；

运行顺序不定；

安排不当，会使速度很慢(因为内部是个串行的运行)。

2.同步函数

(1)_syncthreads()

线程块内线程同步；保证线程块内所有线程都执行到统一位置。

(2)_threadfence()

一个线程调用_threadfence后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对grid在的所有线程可见。

(3)_threadfence_bolck()

一个线程调用_threadfence_block后，该线程在该语句前对全局存储器或共享存储器的访问已经全部完成，执行结果对block在的所有线程可见。

注：(2)(3)两个函数的作用：及时通知其他线程，全局内存或共享内存内的结果已经读入或写入完成了。

3.CPU/GPU同步 cudaStreamSynchronize()/cudaEventSynchronize()主机端代码中使用cudaThreadSynchronize()函数实现CPU与GPU线程同步运行

kernel启动后控制权将异步返回，利用该函数可以确定所有设备端线程均已运行结束

五、并行化高效策略(一)

1.规约Reduce

2.扫描Scan

3.直方图第一步：并行计算局部直方图

第二步：把所有局部直方图每个分组bin使用Reduction(规约)并行累加起来形成一个总的直方图

4.压缩与分配

(1)压缩：在输出中，对真值输入分配1，对假值输入分配0

压缩的步骤：通过一个并行判断把需要进行分析的线程先通过一个判断全部筛选出来

把筛选好的线程做并行分析

(2)分配：输出项数可以动态的从每一个输入项计算出

5.分段扫描

6.排序

(1)奇偶排序(转排序、冒泡算法)

通过比较数组中相邻的(奇-偶)位置数字对，如果该奇偶对是错误的顺序(第一个大于第二个)，则交换。下一步重复该操作，针对所有的(奇-偶)位置数字对，如此交替进行下去。

(2)归并排序

将两个已经排序的序列合并成一个序列的操作

(3)排序网