《CUDA并行程序设计-GPU编程指南》读书笔记--(2)CUDA内存处理

本文详细介绍了CUDA中的内存处理,包括高速缓存、寄存器、共享内存、常量内存和全局内存的使用。强调了GPU如何通过多线程隐藏内存延迟,并解释了内存访问优化的重要性,如寄存器使用、合并访问和对齐内存分配。
摘要由CSDN通过智能技术生成

CUDA内存处理

高速缓存

这里写图片描述


不仅要思考如何高效地访问全局内存,也要时刻想办法减少对全局内存的访问次数,尤其在数据会被重复利用的时候。


CPU与GPU架构的一个主要区别就是CPU与GPU映射寄存器的方式。CPU通过使用寄存器重命名和栈来执行多线程。为了运行一个新任务,CPU需要进行上下文切换,将当前所有寄存器的状态保存到栈(系统内存)上,然后从栈中恢复当前需要执行的新线程上次的执行状态。这些操作通常需要花费上百个CPU时钟周期。如果在CPU上开启过多的线程,时间几乎都将花费在上下文切换过程中寄存器内容的换进/换出操作上。因此,如果在CPU开启过多的线程,有效工作的吞吐量将会快速降低。

然而,GPU却恰恰相反。GPU利用多线程隐藏了内存获取与指令执行带来的延迟。因此,在GPU上开启过少的线程反而会因为等待内存事务使GPU处于闲置状态。此外,GPU也不使用寄存器重命名的机制,而是致力于为每一个线程都分配真实的寄存器。因此,当需要上下文切换时,所需要的操作就是将指向当前寄存器组的选择器(或指针)更新,以指向下一个执行的线程束的寄存器组,因此几乎是零开销。


如果一个内核函数中的每个线程需要的寄存器过多,在每个SM中GPU能够调度的线程块的数量就会受到限制,因此总的可以执行的线程数量也会受到限制。


for(i=0;i<31;<i++){
    packed_result |= (pack_array[i]<<i);
}

如果变量packed_result存于内存中,则需要做32次读/写内存的操作。但如果将变量packed_result设置为局部变量,编译器会将其放入寄存器中,在寄存器中而不是在主内存中做操作,最后再将结果写回主内存中,因此可节省31次内存读/写的操作。

寄存器版本:

__global__ void tes
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值