【cuda by example学习笔记】-2023.7.31

常量内存

用于保存在核函数执行期间不会发生变化的数据。某些情况用常量内存替换全局内存能有效减少带宽。

修饰符:__constant__

需要使用cudaMemcpyToSymbol()将数据复制到常量内存

常量内存中读取数据可以节省带宽的两个原因:

        1.对常量内存的单独读操作会将数据广播到其他的“邻近”线程,会节约15次读操作。

        2.常量内存的数据将缓存起来,对相同地址的数据进行连续读操作不会产生额外内存通信量。

“邻近”是什么意思?我们需要先了解以下线程束。

线程束

它是一个线程集合,程序的每一行,线程束中的每个线程都在不同数据上执行相同的指令。处理常量内存时,NVIDIA硬件将单次读写内存操作广播到每个半线程束,如果半线程束中每个线程都要读取常量内存中相同地址的数据,那么GPU只会产生一次读取请求,将数据广播到半线程束的每个线程,这样产生的内存流量只是使用全局内存的1/16(约6%)。实际使用中会减少更多带宽,因为硬件会主动把常量内存数据缓存在GPU上,这样当其他半线程束需要数据时,将命中缓存。同时,也可能会带来负面影响。如果半线程束上的每个线程需要访问常量内存上不同的数据,那么读取操作将被串行化,但在全局内存中读取,这些请求会被同时发出。那么怎么可以知道,产生的影响是正面的还是负面的?我们可以使用GPU的事件API。

事件

cuda中的事件本质是一个GPU时间戳。获得一个时间戳需要两个步骤:

        1.创建一个事件。cudaEventCreate()

        2.记录一个事件。cudaEventRecord()

要统计代码执行时间,我们要创建开始事件和结束事件。但是存在一个问题,对于一些异步函数,GPU执行完之前,CPU会执行程序中的下一行代码,这将使计时工作更加复杂。采用的解决办法就是一行代码:cudaEventSynchronize(),当该函数返回结果时,可以安全的读取结束的时间戳。

由于cuda事件是在GPU上执行的,因此它不适用于在CPU和GPU上混合执行的代码。

  • 1
    点赞
  • 1
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值