CUDA编程基础与实践学习笔记（八）

最新推荐文章于 2024-11-18 17:41:44 发布

zkxhlbt

最新推荐文章于 2024-11-18 17:41:44 发布

阅读量392

点赞数

分类专栏： CUDA 编程文章标签： pytorch cv c++

本文链接：https://blog.csdn.net/zkxhlbt/article/details/124333400

版权

本文介绍了CUDA编程中的共享内存概念，它能提升全局内存访问效率，用于线程块内的高效通信。文章探讨了数字规约算法在并行计算中的应用，并讨论了静态与动态共享内存的差异，包括内存bank的工作原理和bank冲突问题。

摘要由CSDN通过智能技术生成

共享内存
共享内存可以被程序员直接操控，减少对全局内存的访问，实现高效的线程块内部通信，也可以提高全局内存访问的合并度。

数字规约算法。可以并行快速算出数组的和。
共享内存限定符__shared__ 变量名可以以s_开头，提醒自己这是共享内存变量。所有线程都有共享内存变量的副本，每个副本不一样但共用一个变量名。核函数对于共享内存的操作作用在所有线程上。
因为还没有在实际中用过，所以有一个忧虑的点：其实程序一开始就要将数据从全局变量中移到共享内存中，这一步的耗时如何？如果算术强度不大或者不对共享内存进行多次访问，那么其实是有点亏的。。。但一般应该会多次访问。有一个好处是不会再运行过程中改变全局变量的内容。
分为静态共享内存和动态共享内存.<<<grid_size, block_size, sizeof(real) * block_size>>> 定义核函数的时候，第三个参数其实是共享内存的大小，被默认设为了0。动态共享内存在性能上几乎没有差别，而且可以提高程序的可维护性。定义：extern __shared__ real s_[] 注意这里是数组，不能定义成指针，否则无法完成编译。指针和数组并不等价。
共享内存在物理上被分为32个同样宽度、能被同时访问的内存bank，每个bank的宽度为4B(只有kepler架构为8B)，地址模128的结果相同的属于同一层.

如果不同线程同时对同一个bank的不同层访问，就会触发多次内存事务(memory transaction)，从而导致bank冲突，同时访问同一bank的n层则称为n路bank冲突。如果不同时访问，则只会触发一次内存事务。