分类：
2009-08-20 11:39:20
------------------------------------------------
本文系本站原创,欢迎转载!
一直以来对GPU和CPU的协同工作非常感兴趣，但是一直没去弄，前几天，教研室的一个项目组做了一个OPENCL的相关的东西，于是就看了下CUDA架构下的GPU运行机制。
原来一直以为GPU就是以SFU（协处理器）的模式和CPU交互，协同，看了资料后才发现，GPU和CPU的协同工作同,ARM和DSP的协同工作模式几乎一样，就是host,device模式，当然dsp也可自启动，不清楚GPU有这个没有。Host,device模式就是，CPU将GPU要执行的代码复制到GPU的RAM，然后发送相关的控制信息使GPU运行。
这个就是典型的CUDA的伪代码。
unsigned int numBytes = N * sizeof(float)
float* h_A = (float*) malloc(numBytes);
// allocate device memory
float* d_A = 0;
cudaMalloc((void**)&d_A, numbytes);
// copy data from host to device
cudaMemcpy(d_A, h_A, numBytes, cudaMemcpyHostToDevice);
// execute the kernel
Increment_gpu<<< N/blockSize, blockSize>>>(d_A, b);
//Copy data from device back to host
cudaMemcpy(h_A, d_A, numBytes, cudaMemcpyDeviceToHost);
// free device memory
cudaFree(d_A);
先在主机（cpu）分配一段内存，将这段内存的数据复制到GPU(device)的ram ,然后GPU开始运行代码(kernel)，完成后GPU将处理好的数据复制到指定的cpu的ram ，因为一般的cpu的内存有可能搬移，这一过程一般是通过DMA的方式进行。
CUDA架构下的GPU的多线程处理能力可谓是异常强大。
讲到线程：
就先来讲下我对线程的理解：
线程在操作系统中是指：
执行同一段代码，且共享地址空间的两个不同代码运行体：
从上面可以看出线程有三个关键点：
同一代码：
共享地址空间
不同运行体
不同运行体现在运行环境，在操作系统中体现为线程的私有堆栈，堆栈保存着线程的寄存器等信息，而在GPU中体现为不同ThreadID,thread的寄存器，local memory,这些是每个线程拥有的，私有的。
正因为共享地址空间这一原因，线程之间通信异常方便，带来的性能损耗比进程间通信更少，同时线程切换也相对容易很多。
但是在操作系统中多线程是可以并发执行，但是不能并行执行，因为他只有一个执行体，如果在多核cpu,应该就可以达到并行执行，但是这也有一个前提：那就是线程间完全独立，这一点估计大部分不好满足。
可在CUDA架构下的GPU算是真正的并行执行，且并行执行的线程数还非常大。
CUDA下的GPU是如何做到的?
先看下下面这张硬件架构图。
从上面的这张架构图我们可以看到，NVIDIA的Shader核心是由若干个材质处理单元（TPC）的群组所构成的。例如一个8800GTX，它内部就有8个群组，而8800GTS内部就有6个群组。事实上，每一个群组都是由1个材质处理单元和2个流多重处理器构成的。而处理器又由一个前端的读取/解码单元，一个指令发送单元，一个由八个计算单元组成的组，和2个SFU超级功能单元所组成，他们在处理指令方面都属于SIMD单指令多数据流。同样这类指令也适用于所有warp中的线程。NVIDIA这种并行模式叫做SIMT单指令多线程执行单元.
上面的instuction fetch/dispatch就是达到了各个核心执行相同的代码，这一点很好理解，就是相当于有一个专门的取指解释部件，他然后给八个执行processor发控制命令（这个可以看成cpu模型里的微命令），这样processor就做同样的事，只不过他们有着各自的环境（threadid,这个其实体现在他寄存器上或者local memory），所以执行的结果就不样（数据来源，数据结果）,这也和cpu多核的并行执行不一样，多核cpu，各个processor拥有独立的instuction fetch/dispatch部件。
那他们是如何做到完全独立的呢，这个就是编程人员要做的。就是让他们访存的不要冲突。这个主要在share memory的confision（冲突）上，即一个stream mutiprocessor同时执行的16trhead不能存取输入同一bank的内存。
目前 CUDA 装置中，每个 multiprocessor 有 16KB 的 shared memory。Shared memory 分成 16 个 bank。如果同时每个 thread 是存取不同的 bank，就不会产生任何问题，存取 shared memory 的速度和存取寄存器相同。不过，如果同时有两个（或更多个） threads 存取同一个 bank 的数据，就会发生 bank conflict，这些 threads 就必须照顺序去存取，而无法同时存取 shared memory 了，这个也好理解，一个bank只有一组地址线，数据，当然不允许多个processor同时访问。
比如如下一个程序：
__shared__ int data[128];
int number = data[base + tid];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行，因此分属于不同的 half warp 的 threads，不会造成 bank conflict。
而如果是
int number = data[base +4* tid];
0,4线程同时访问bank0,1,5同时访问bank1,就会出现 share memory 的bank conflict时，这时冲突的线程就会阻塞，这样就会影响并行行。
在 CUDA 架构下，显示芯片执行时的最小单位是 thread。数个 thread 可以组成一个 block。一个 block 中的 thread 能存取同一块共享的内存，而且可以快速进行同步的动作。
每一个 block 所能包含的 thread 数目是有限的。不过，执行相同程序的 block，可以组成 grid。不同 block 中的 thread 无法存取同一个共享的内存，因此无法直接互通或进行同步。因此，不同 block 中的 thread 能合作的程度是比较低的。不过，利用这个模式，可以让程序不用担心显示芯片实际上能同时执行的 thread 数目限制。例如，一个具有很少量执行单元的显示芯片，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同grid 则可以执行不同的程序（即 kernel）。
Grid、block 和 thread 的关系，如下图所示：
每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread 则有共享的一份 share memory。此外，所有的 thread（包括不同 block 的 thread）都共享一份 global memory、constant memory、和 texture memory。不同的 grid 则有各自的 global memory、constant memory 和 texture memory
Share memory为block块中的线程提供了一种信息沟通的途径。它最重要的作用就是强调限制级别。所有的线程都保证被封装到一个block块中，这就保证了多重处理器可以有效的处理这些任务。反过来说，指派block块到不同的多重处理器中，是非常不确定的。这就意味着，在执行的时候来自不同block块中的线程是无法进行通信的，使用这种存储区是一件非常复杂的事儿。但是设计这些复杂的存储区，对于整体架构来说也是值得的。因为某些特殊的线程，可能会贸然访问主内存，这也许会产生很多冲突，而这种共享存储区可以快速的链接寄存器。
整个关系图如下：
block块内部架构
这些共享的存储区并不是多重处理器唯一可以访问的存储设备。很显然，他们可以使用显存，但是相对于共享存储区来说，显存的带宽和速度都不如前者。因此，这种机制可以抑制对内存过于频繁的访问。NVIDIA也提供了多重处理器的高速缓存，可以存储常量和纹理，大致相当于每个多重处理器能分配到8KB的空间。
多重处理器内部架构
同时，多重处理器也具备8192个寄存器。在多重处理器中所有激活的block块中所有的线程，都可以通过这些寄存器共享信息。不过激活的warps总数被限制在24个，也就是768个线程。例如8800GTX最多可在同一时间处理12288个线程。在这方面加以限制，是为了不让GPU的资源消耗的太多，可以更合理的分配计算任务。
针对CUDA优化过的程序从本质上讲就是在block块的数量与他们的尺寸之间找到一种平衡。在一个block块中加入更多的线程，有利于提高内存潜伏期的效率，但是与此同时可以使用的寄存器数量就要少了。如果block块中的线程数太多了，比如达到512个线程的水平，那么整个流水线的执行效能就大大降低了。这仅仅够喂饱一个多重处理器的，浪费了256个线程的处理能力。因此NVIDIA建议每个block块使用128至256个线程，这是最为折中的办法，它会在内核的潜伏期与寄存器数量之间。
我们以学校的一次体检测试为例来讲一下cuda架构下的GPU的工作机制。
张无能是“曾哥小学”的校长，一天他接到上级领导批示“为了更好的促进小学生身体健康发展，你要对全校学生进行一次全面的体检，体检的具体细节表我已经发到你们校长室了”，张无能回到校长室一看体检说明表，说了声“都啥子东东啊，这么多啊，我还得去打高尔夫球呢，看来直接交给医院负责算了”。于是他就将医院院长叫来说，“这里是体检说明书，你按照这个给我们全校的学生体检，完成了后将全校学生的结果放到校长办公室，然后通知我，我会到校长办公室去看的”。
医院院长一看，“靠，两千多人，一个一个体检得要多久啊，干脆让机器人给学生体检得了“。由于这个机器人的智能化程度不高，必须要由熟练的操作人员语音控制，于是他叫来了医院的机器人体检操作师，将体检说明书递给对说：”你按照这个说明书指导机器人体检学生，然后将体检报告送到院长办公室，并通知我，你现在就到体检中心去，开始体检时我会通知你“。然后院长将全校学生召集到医院外面的广场，然后给每个人发一个编号和一张体检表。由于医院体检机器人只有100台，于是，他就打算分组进行，每次100人。于是他叫前100人进入医院，然后通知体检师可以体检了。
体检师于是拿着扩音器按照体检说明书对体检机器人大喊：
1. 请从学生手中拿体检表。
2．请为他们测量身高，并将身高结果写到在体检表上。
3．如果是女生，请在体检表上填写女。
于是机器人开始识别男女，发现是女学生的机器人的都动笔写上了“女“字，而发现是男学生的机器人一动不动。
4．如果是男生，请在体检表上填写男。
于是机器人开始识别男女，发现是男学生的机器人的都动笔写上了“女“字，而发现是女学生的机器人一动不动。
5. 请为他们测量他们的体重，并将体重结果写到在体检表上。。
………..
…………..
7.将体检表交到院长办公室。
8.将学生带到到医院后广场等待其他同学。
体检师通知院长这一百人体检完了。
然后院长就开始叫下一100个学生执行上面的流程。
上面体检说明书就是CUDA程序，学生是线程，学校就是cpu,校长办公室是cpu的ram(容纳体检结果),医院是GPU,院长办公室是GPU的ram(用来暂放体检表结果),校长是运行在cpu上的操作系统，院长是GPU的调度系统和资源分配系统，体检师是流多重处理器stream mutiprocessor中的前端的读取/解码单元 instruction fetch/patch,机器人就是流处理单元stream processor,校长给院长发通知，就是cpu将GPU要执行的代码复制到GPU的RAM。
同时上面的3，4步骤说明GPU中同时运行的Thread是完全一致，即使在条件语句中，每个thread都会执行完条件的两个分支，只是条件不满足的thread执行空操作，而不是去执行条件的另一分支。
CUDA runtime相当于JAVA虚拟,.NET Framework。Opencl, c for CUDA程序最终都要生成PTX中间语言。
这个中间语言然后由CUDA runtime解释为具体GPU能够识别的语言，并且添加具体的cpu平台的一些准备代码，比如拷贝数据到gpu ram，这个对于不同cpu,肯定是不一样。
那为什么要一个一个CUDA Runtime,其实就是为了跨平台，就像java虚拟机，.net平台，都是独立出一个中间平台，这个中间平台是和操作系统相关，而平台上面的就是跨平台的了，对于不同的cpu平台，操作系统平台，只需使用其对应的CUDA Runtime,而之上的CUDA程序不用改变，就可以运行。
