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