原文链接：https://traxnet.wordpress.com/2011/07/18/understanding-modern-gpus-2/
理解现代GPU的第二章，我们将谈一谈driver，数据流以及他们涉及到的模块。前一篇文章中主要介绍了GPU的软件前端，从这篇文章开始的内容将会和硬件联系更加紧密。

用户空间和内核空间命令流

对于任何一种API，如OpenGL，DirectX，OpenCL，它们会在用户空间实现相应的功能。Windows 7/Vista的WDDM（Windows Display Driver Model）驱动程序模型给出了如下范式：
diagram-illustrating-the-windows-vista-and-later-display-driver-model-architecture1.png

每个使用图形 API 的应用程序都会将一组driver/API 功能加载到其专用进程内存中，这些功能不与其他进程共享。用户态driver（UMD）会将所有 API 调用按顺序转换和存储，得到命令缓冲区（command buffer），后者再被提交到内核态driver（KMD)。每个进程的command buffer取决于供应商，并且可能包含提供给 GPU 的实际命令。此外，shaders也会在运行时（即UMD）被编译，并且driver可能会注入自定义代码（用来修补着色器）以实现不能为硬件所用的特定功能。这部分取决于正在使用的GPU。

由于 GPU 是不同进程之间共享的资源，因此必须有一种机制来确保命令按顺序执行，GPU 可以被所有这些进程使用，并且没有数据corruption。这是由KMD完成的。为此KMD实现了一套调度程序，还有用于与 GPU 通信的终点：命令环形缓冲区（Command Ring Buffer）。和Command Ring Buffer相对应的是命令处理器（Command Processor，CP），它读取数据流，解码命令并将其喂给线程/流调度程序（Threads/Stream Scheduler）（我们将在其他文章中讨论这一点）。
cb.png

KMD scheduler从每个单独的（进程）command buffer读取数据，并将它们移动到DMA command buffer。事实上，这会产生资源浪费。现在的GPU 包含一个DMA控制器和一个MMU。前者允许GPU直接与主机RAM通信，以获取和写入数据，无需CPU干预。MMU对GPU/主机内存进行虚拟化，同时提供一些内存保护。例如，费米架构包含2个DMA引擎，以获得双向同时传输。另一种可能是在没有CPU干预的情况下在GPU之间传输数据。请务必注意，这些传输与当前的command buffer并行执行，从而增加了另一个级别的并行性。

我们向 GPU 发送一个命令，使用其DMA引擎从内存区域获取数据，而不是直接传输，这样我们就可以为每一个UMD创建不同的command buffer，并让 GPU 获取它们。

COMMAND/RING BUFFER

DirectX给出了如何设置command buffer。
基本上，通过command buffer，你可以在GPU中设置一些状态，将其设置为获取数据并发出执行命令。以前的用户API有一个很大的缺点，你必须通过命令直接向 API 指定基元，例如 glBegin，glEnd。这些有害的调用现在已从某些API中删除，如OpenGL ES，因为它们是现代显卡的性能杀手。你可以将CPU和GPU视为通过command ring buffer进行通信的两个线程。它是一个环（FIFO），由CPU填充并由GPU读取，直到耗尽为止。如果环为空（写入和读取指针相等），GPU 将停止并等待，直到有事情要做。在这种情况下的性能问题可能是CPU bounded。如果 CPU 填满了整个缓冲区，并且必须等待一些可用空间，则问题是GPU bounded。

cb1.png

上图显示了通过环形缓冲区（Ring Buffer，RB） 连接的主机和图形控制器（CP）。RB使用固定缓冲区大小进行初始化，并且写入指针和读取指针都设置为零（空缓冲区）。driver将数据包写入RB，并且更新设备内的写入指针寄存器。当设备读取数据时，更新读取指针。更新这两个指针会产生一些开销，缓解这个开销的方法是，只在使用了某些数据块时（通过将数据包分组为块）而不是为每个数据包执行此操作，才通过更新这些寄存器。这也需要更多的逻辑以避免在RB已满时写入（此处的更多信息 http://developer.amd.com/gpu_assets/R5xx_Acceleration_v1.2.pdf 似乎已经404...）。

Command stream导致了另一些必须注意的同步问题。想象一下：你创建了一个将要处理的巨大数据数组，但是一旦GPU从主内存获取到它，我们希望尽快用新数据更新它。CPU 如何知道某些命令已被处理，以便我们可以更新数组？GPU可以从内存中获取数据，但同时，GPU和CPU都可以并行获取数据。解决方案实际上非常简单。有一些特殊的命令类型被填充到到command stream中，其称为围栏（fence）（我发现了VIA关于此事的专利 http://www.patentgenius.com/patent/7755632.html）。这些fence由GPU读取，GPU会更新一些寄存器，以便CPU知道我们已经获取到了stream中的数据。

到目前为止，OpenCL kernel或顶点着色器会发生什么情况？内核代码由UMD编译为中间语言（例如，用于CUDA设备的PTX或用于AMD设备的AMD IL）。然后，代码由driver编译到特定硬件并传递到 GPU。由于某些 GPU 可能缺少功能或需要不同的步骤来计算某些功能，因此代码需要针对计算机上运行的硬件。例如，双浮点操作可能需要额外的步骤才能在缺乏专用双精度浮点处理器的硬件上获得所需的精度。一些GPU架构牺牲了IEEE合规性，而其他GPU架构则完全缺乏双精度。

我们在谈论的一些概念可以在OpenCL API中看到（请记住，我告诉过你，从逻辑角度来看，这个API和硬件的实际工作方式之间存在某种直接映射）。使用clCreateCommandQueue创建command buffer，使用clEnqueueReadBuffer从设备内存中读取数据进行排队，使用clEnqueueNDRangeKernel运行计算kernel，依此类推。

在下一章中，我们将讨论CP和一些设置逻辑。敬请期待！