CUDA学习

CUDA相关知识点

CUDA相关知识点

近期在学习CUDA相关东西，记录些零碎知识。

关于cudaDeviceReset();

这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。

关于报错

nvcc warning : The ‘compute_20’, ‘sm_20’, and ‘sm_21’ architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
sum_arrays.cu
���ڴ����� sum_arrays.lib �Ͷ��� sum_arrays.exp

原因是，Makefile中采用了CUDA的compute capability 2.0和2.1，这是两种计算能力。安装的CUDA版本是8.0，但 从CUDA 8.0开始compute capability 2.0和2.1被弃用了，所以可以将-gencode arch=compute_20,code=sm_20 和-gencode arch=compute_20,code=sm_21这两行删除即可。

nvprof的使用

终端目录下切换到当前文件夹下执行.exe文件 例如nvprof ./sum_arrays_timer.exe
获取各部分执行时间，nvprof这个强大的工具给了我们优化的目标，分析数据可以得出我们重点工作要集中在哪部分。
箭头所指为核函数执行时间

优化方面

调整线程以及网格块大小，32整数倍。
规约，交错配对，避免分支，比较几种方法的不同。
循环展开
用一个线程块处理多块数据
最后的32个线程展开
对齐与合并访问
对齐合并内存访问，以减少带宽的浪费
足够的并发内存操作，以隐藏内存延迟
增加每个线程中执行独立内存操作的数量
对核函数启动的执行配置进行试验，已充分体现每个SM的并行性
内存带宽调整

内存

共享内存：
共享内存是块内线程可见的，有竞争问题的存在，为了避免内存竞争，可以使用同步语句：
void __syncthreads();
此语句相当于在线程块执行时各个线程的一个障碍点，当块内所有线程都执行到本障碍点的时候才能进行下一步的计算，这样可以设计出避免内存竞争的共享内存使用程序、
注意，__syncthreads();频繁使用会影响内核执行效率。

常量内存：
常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用：
__constant__
修饰，常量内存在核函数外，全局范围内声明，对于所有设备，只可以声明64k的常量内存，常量内存静态声明，并对同一编译单元中的所有核函数可见。
常量内存不能被修改的，这里不能被修改指的是被核函数修改，主机端代码是可以初始化常量内存的，不然这个内存谁都不能改就没有什么使用意义了，常量内存，被主机端初始化后不能被核函数修改，初始化函数如下：
常量内存的最优访问模式是线程束所有线程访问一个位置

cudaError_t cudaMemcpyToSymbol(const void* symbol,const void *src,size_t count);

同 cudaMemcpy的参数列表相似，从src复制count个字节的内存到symbol里面，也就是设备端的常量内存。多数情况下此函数是同步的，也就是会马上被执行。当线程束中所有线程都从相同的地址取数据时，常量内存表现较好，比如执行某一个多项式计算，系数都存在常量内存里效率会非常高，但是如果不同的线程取不同地址的数据，常量内存就不那么好了，因为常量内存的读取机制是：
一次读取会广播给所有线程束内的线程。