赞
踩
CUDA存储器模型:
GPU片内:register,shared memory;
板载显存:local memory,constant memory, texture memory ,global memory;
host 内存: host memory, pinned memory.
1)register: 访问延迟极低;
每个线程占有的register有限,编程时不要为其分配过多私有变量;
对每个线程来说,寄存器都是线程私有的–这与CPU中一样。如果寄存器被消耗完,数据将被存储在本地存储器(localmemory)。
Localmemory对每个线程也是私有的,但是localmemory中的数据是被保存在显存中,而不是片内的寄存器或者缓存中,因此速度很慢。
线程的输入和中间输出变量将被保存在寄存器或者本地存储器中。
2)local memory:寄存器被使用完毕,数据将被存储在局部存储器中;
3)shared memory:访问速度与寄存器相似;
实现线程间通信的延迟最小;
保存公用的计数器或者block的公用结果;
用于线程间通信的共享存储器。共享存储器是可以被同一block中的所有thread访问的可读写存储器。
访问共享存储器几乎和访问寄存器一样快速,是实现线程间通信的延迟最小的方法。
共享存储器可以实现许多不同的功能,如用于保存共用的计数器(例如计算循环次数)或者block的公用结果。
4)global memory:存在于显存中,也称为线性内存(显存可以被定义为线性存储器或者CUDA数组);
使用的是普通的显存,无缓存,可读写,速度慢
整个网格中的任意线程都能读写全局存储器的任意位置,并且既可以从CPU访问,也可以从GPU访问。
cudaMalloc()函数分配,cudaFree()函数释放,cudaMemcpy()进行主机端与设备端的数据传输;
Cuda执行模型
SIMT和SIMD架构:
相似性:两者都将相同的指令广播给多个执行单元来实现并行。
区别:SIMD要求同一向量的所有元素在一个同步组中一起执行,而SIMT则允许同一线程束的多个线程独立执行。
尽管一个线程束中所有线程在相同的程序地址同时开始执行,但是单独的线程仍有可能有不同的行为。SIMT确保可以编写独立的线程级并行代码、标量线程、以及用于协调线程的数据并行代码。
SIMT模型包含3个SIMD所不具备的关键特征:
每个线程都有自己的指令地址计数器;
每个线程都有自己的寄存器状态;
每个线程都可以有一个独立的执行路径
Fermi架构,Kepler架构
GPU架构:GPU架构是围绕一个流式多处理器(SM)的扩展阵列搭建的
线程组织结构
网格 grid
kernel在device上执行时,实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间。grid是线程结构的第一层次。
线程块 block
网格又可以分为很多线程块(block),一个block里面包含很多线程。各block是并行执行的,block间无法通信,也没有执行顺序。block的数量限制为不超过65535(硬件限制)。第二层次。
grid和block都是定义为dim3类型的变量,dim3可以看成是包含三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1。grid和block可以灵活地定义为1-dim,2-dim以及3-dim结构。
CUDA中,每一个线程都要执行核函数,每一个线程需要kernel的两个内置坐标变量(blockIdx,threadIdx)来唯一标识,其中blockIdx指明线程所在grid中的位置,threaIdx指明线程所在block中的位置。它们都是dim3类型变量。
一个线程在block中的全局ID,必须还要知道block的组织结构,这是通过线程的内置变量blockDim来获得。它获取block各个维度的大小。对于一个2-dim的block(D_x,D_y),线程 (x,y) 的ID值为(x+y∗D_x),如果是3-dim的block (D_x,D_y,D_z),线程(x,y,z)的ID值为(x+y∗D_x+z∗D_x∗D_y) 。另外线程还有内置变量gridDim,用于获得grid各个维度的大小。
每个block有包含共享内存(Shared Memory),可以被线程块中所有线程共享,其生命周期与线程块一致。
每个thread有自己的私有本地内存(Local Memory)。此外,所有的线程都可以访问全局内存(Global Memory),还可以访问一些只读内存块:常量内存(Constant Memory)和纹理内存(Texture Memory)。
线程 thread
一个CUDA的并行程序会被以许多个threads来执行。数个threads会被群组成一个block,同一个block中的threads可以同步,也可以通过shared memory通信。
线程束 warp
GPU执行程序时的调度单位,SM的基本执行单元。目前在CUDA架构中,warp是一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。同一个warp中的每个线程都将以不同数据资源执行相同的指令,这就是所谓 SIMT架构(Single-Instruction, Multiple-Thread,单指令多线程)。一个warp中的线程必然在同一个block中,如果block所含线程数目不是warp大小的整数倍,那么多出的那些thread所在的warp中,会剩余一些inactive(待用)的thread
Warp divergence:因为所有同一个warp中的thread必须执行相同的指令,那么如果这些线程在遇到控制流语句时,如果进入不同的分支,那么同一时刻除了正在执行的分支之外,其余分支都被阻塞了,十分影响性能。这类问题就是warp divergence。
SM(Streaming Multiprocessor): 流式多处理器,每个SM最多占用1536个thread,8个block
Grid:由一个kernel启动所产生的所有线程统称为一个线程网格(Grid)。
同一线程网格中的所有线程共享同全局内存空间。一个网格有多个线程块(Block)构成,一个线程块包含一组线程,同一线程块内的线程协同可以通过“同步”和“共享内存”的方式来实现。不同线程块内的线程不能协作。
在一个网格中,我们通过以下两个坐标变量来定位一个线程,
(1)blockIdx:线程块在线程网格中ID号
(2)threadIdx:线程在线程块内的ID号
这些坐标变量是kernel函数中需要预初始化的内置变量。
当执行有一个核函数时,CUDA Runtime 为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标,我们将数据分配到不同的GPU线程上,然后并行处理所有的数据。
坐标变量blocIdx和threadIdx都是基于unit3定义的CUDA内置的向量类型,分别包含3个无符号的整数结构,可以通过x,y,z三个元素来进行索引。
任务/数据的映射关系
这个我没写
在2里面
现代DRAM系统设计为始终可访问在burst模式下。突发字节被传输到处理器,但当访问不到连续位置时,将被丢弃。
当warp的所有线程执行load指令时,如果所有访问的位置都属于同一个burst段,则只会发出一个DRAM请求,并且访问完全合并。
当访问的位置分布在突发段时边界:凝聚失败,发出多个DRAM请求,通道未完全合并。
原子操作是指不会被线程调度机制打断的操作;这种操作一旦开始,就一直运行到结束,
原子操作的关键概念
-由单个硬件指令对存储器位置地址执行的读-修改-写操作
-读取旧值,计算新值,并将新值写入该位置
-硬件确保在当前原子操作完成之前,没有其他线程可以在同一位置执行另一个读-修改-写操作
-试图在同一位置执行原子操作的任何其他线程通常会被保留在队列中
-所有线程在同一位置连续执行原子操作
并发度=吞吐量∗延迟
直方图
吞吐量=1/延迟
占用率是指每个多处理器(Streaming Multiprocessor,SM)的活动线程束(warps)数量与实际的活动warps数量的比率。
占用率=每个SM中活跃线程束的数量/每个SM中最大的线程束的数量
CPU:强大的算术逻辑单元,大型的缓存,复杂的控制逻辑,低延时,适合处理复杂的逻辑分支判断运算
GPU:简单大量的计算单元,小型的缓存,简单的控制逻辑,吞吐量大,适合处理大规模的数据并发运算
向量加法host code
矩阵乘法
直方图
卷积
并行归约
扫描算法
Copyright © 2003-2013 www.wpsshop.cn 版权所有,并保留所有权利。