杭州网站建设专家,做产地证的网站,制作自己的网站教程,seo免费优化1. 引言
前序博客#xff1a;
CUDA简介——基本概念CUDA简介——编程模式CUDA简介——For循环并行化CUDA简介——Grid和Block内Thread索引 CUDA内存模式#xff0c;采用分层设计#xff0c;是CUDA程序与正常C程序的最大不同之处#xff1a;
Thread-Memory Correspondenc…1. 引言
前序博客
CUDA简介——基本概念CUDA简介——编程模式CUDA简介——For循环并行化CUDA简介——Grid和Block内Thread索引 CUDA内存模式采用分层设计是CUDA程序与正常C程序的最大不同之处
Thread-Memory CorrespondenceBlock-Memory CorrespondenceGrid-Memory Correspondence
总体的CUDA内存模式为
1Registers Local Memory声明在Kernel内的常规变量。小空间变量存储于Register中大空间变量存储于Local Memory中。因Local Memory操作速度慢应尽量避免使用Local Memory。2Shared Memory允许同一Block内的Threads间相互通信。3Constant Memory用于存储Kernel执行期间不变的数据。会缓存到On-Chip memory中可大大降低Kernel执行期间Global Memory的通信吞吐量。4Global Memory用于存储与Host交互的数据。
2. Thread-Memory Correspondence Thread-Memory Correspondence即Threads等价为Local Memory (and Registers)
其范围为对相应Thread是private的对其它Threads来说是不可访问的。其生命周期为Thread。当Thread执行完成与该Thread相关的任何local memory (and Registers) 都将自动释放。不过local memory和registers存在完全不同的性能特性。
3. Block-Memory Correspondence Block-Memory Correspondence即Blocks等价为Shared Memory
其范围为同一Block内的每个Thread均可访问。其生命周期为Block。当Block执行完成其shared memory内容都将自动释放。
4. Grid-Memory Correspondence Grid-Memory Correspondence即Grids等价为Global Memory
其范围为所有 Grids内的每个Thread均可访问。其生命周期为Host代码内的整个main()程序。或可通过在Host代码中手工调用cudaFree(...)来释放。
5. Device内存模式 Host由CPU及机器内存等组成。Device由GPU及其DRAM等组成。【Device图上的绿色方格表示的是一个CUDA core。】 Device的DRAM中有
Global Memory物理空间Local Memory物理空间。需注意此处的Local并不是指其物理位置此处的Local是指该内存空间的scope范围和 lifetime生命周期。
而Device的GPU中有
Registers物理空间Shared Memory物理空间
Device图上的绿色方格表示的是一个CUDA core。Device上的CUDA cores组合在一起成为streaming Multiprocessor简称为SM。 Device图上的黄色方格表示SM。黄色方格组合在一起为CUDA cores集合。
位于SM上的内存称为
“On-Chip” Device memory。因此Registers和Shared Memory均对应为 “On-Chip” Device memory。因Registers和Shared Memory均物理存在与GPU的streaming Multiprocessor中。
非SM上的内存称为
“Off-Chip” Device memory。因其并不存在与GPU之上。对应Global Memory和Local Memory均为“Off-Chip” Device memory。也即Device上的DRAM为 “Off-Chip” Device memory。
以NVIDIA GPU Geforce Titan 物理布局为例
上图蓝色框所示为Device的DRAM均为 “Off-Chip” Device memory。上图绿色框即为实际的GPU对应为“On-Chip” Device memory。
理解Blocks如何映射到SM是设计kernel的基本要求从而获得优化的GPU计算性能。
5.1 内存速度
不同的内存空间其带宽和延迟各不相同
on-chip memory操作速度要快于off-chip memory。
5.2 Global Memory访问 Global Memory访问方式有
cudaMalloc()cudaMemset()cudaMemCopy()cudaFree()
无法避免不使用Global Memory因必须使用Global Memory空间来将数据由host传递到device。不过应尽量减少Global Memory通讯量因为其速度很慢。
Global Memory的优势在于
其通常很大。如计算机内存为8GB或16GB而Titan和Tesla k40均有6GB的global DRAM。
5.3 Registers and Local Memory Kernel中声明的变量存储于Register中
对应为On-Chip Device memory。为最快的内存形式。
太大不适于Register的数组将存储在Local Memory中
对应为Off-Chip Device memory。由编译器控制。“Local”是指范围而不是位置。此处“Local”是指相对每个Thread的Local Memory。 每个Thread均有其自己的private local memory和registers对其它Thread不可访问。 应尽量避免因Local Memory为最慢的内存形式之一。Registers为最快的内存形式。 因此设计目标为避免将更多信息存储于local变量中以免超过register的存储空间。register space为稀缺硬件资源。应更好地规划充分利用。
5.4 Shared Memory 借助Shared Memory
支持同一Block内的Threads之间相互通信 同步方式通信 Shared Memory为非常特殊的内存对实现计算性能和正确性至关重要 Shared Memory处理速度快仅次于Registers。因其为On-Chip device memory。支持Block内的Threads相互通信可将其看成是用户定义的L1 Cache可用作“scratch-pad高速暂存存储器”内存。后续将介绍Shared Memory和L1 Cache关系密切。 使用__shared___关键字来表示分配的为shared memory
5.5 Constant Memory
Constant Memory为Device Memory的特殊区域
用于存储Kernel执行过程中不变的数据。对Kernel来说是只读的。Constant Memory为Off-Chip Device Memory。Constant Memory 积极缓存到 On-Chip Memory中。 Constant Memory的思想在于
GPU没有很大的cache。从而可使用constant memory来实现很简单的cache类型。Constant Memory可以很大因其实际位于Device DRAM中。所有Threads均可访问Constant Memory但其为只读内存。对于需频繁访问但Kernel执行过程中不变的数据可使用Constant Memory。Constant Memory是在off-chain DRAM硬件中实现的但实际其内容会积极缓存到On-Chip Memory中。因此使用Constant Memory可大大降低Kernel执行期间Global Memory的通信吞吐量。
参考资料
[1] Intro to CUDA (part 5): Memory Model
