CUDA连载

本文不是科普，是自我学习道路上的一个笔记，一个总结。

一，术语。

CUDA是NVIDIA公司的通用并行计算架构，全名为 Computer Unified Device Architecture。它目前支持C编程语言。GPU是一个多核处理器，每个核心称为multiprocessor (MP)。

SP(Streaming Processor)是线程最终执行的单位，它包括ALU（逻辑运算单元），一个FPU（浮点运算单元）以及一个Register File(寄存器堆)。SM包括8个SP。SM内包含有一个Instruction Unit、一个Constant Memory、一个Texture Memory，8192个Register、一个16KB的Share Memory、8个Stream Processor(SP)和两个Special Function Units（SFU）。warp愿意为丝线，在CUDA里面它的专有意义为同时运行的一组线程。CUDA程序执行的时候，就是以 warp为单位的，为什么以warp为单位呢？我们假设线程数为512个，如果SM(Streaming Multi-Processor)以一个线程一个线程地去切换，那么效率过低！而且由于SM本身就是一个并行的结构，所以，它可以一次切换多个线程，那么，这个“多个”到底是多少呢？我们就把它定义为warp，这个值一般为32。属于一个warp的线程们的线程ID是连续的。那么，block又是什么呢？这个要从SM说起了，一个GPU有多个SM，它们都会去接受某些任务，以warp的形式去调度。那么，分配到一个SM上的任务就属于一个block，也可以理解为多个warp组成一个block，它们共享存储。从这里可以看出，block的size(线程个数)一定是warp的size的整数倍，这个值一般取64，128，256等等。一个block最多包括32个warp,同一个block里面的线程在同一个MP上执行。而多个block对应一个grid，即一个实际的任务。一个grid里面的block共享全局内存。一个完整的GPU程序称为一个kernel(核心)。如果GPU资源足够，则一次可以运行多个kernel。

关于共享内存的问题，我们可以从CUDA定义的变量存储关键字来区分：

•__device__

–储存于GPU上的globalmemory空间

–和应用程序具有相同的生命期(lifetime)

–可被grid中所有线程存取,CPU代码通过runtime函数存取

•__constant__

–储存于GPU上的constantmemory空间

–和应用程序具有相同的生命期(lifetime)

–可被grid中所有线程存取,CPU代码通过runtime函数存取

•__shared__

–储存于GPU上threadblock内的共享存储器

–和threadblock具有相同的生命期(lifetime)

–只能被threadblock内的线程存取

•无修饰（Local变量）

–储存于SM内的寄存器和localmemory

–和具有相同的生命期(lifetime)

Thread私有

...不定期更新