《CUDA C编程权威指南》——2.1节CUDA编程模型概述

本节书摘来自华章社区《CUDA C编程权威指南》一书中的第2章，第2.1节CUDA编程模型概述，作者[美]　马克斯·格罗斯曼（Max Grossman），更多章节内容可以访问云栖社区“华章社区”公众号查看

2.1　CUDA编程模型概述
CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁。图2-1说明了程序和编程模型实现之间的抽象结构的重要。通信抽象是程序与编程模型实现之间的分界线，它通过专业的硬件原语和操作系统的编译器或库来实现。利用编程模型所编写的程序指定了程序的各组成部分是如何共享信息及相互协作的。编程模型从逻辑上提供了一个特定的计算机架构，通常它体现在编程语言或编程环境中。

除了与其他并行编程模型共有的抽象外，CUDA编程模型还利用GPU架构的计算能力提供了以下几个特有功能。
一种通过层次结构在GPU中组织线程的方法
一种通过层次结构在GPU中访问内存的方法
在本章和下一章你将重点学习第一个主题，而在第4章和第5章将学习第二个主题。
以程序员的角度可以从以下几个不同的层面来看待并行计算。
领域层
逻辑层
硬件层
在编程与算法设计的过程中，你最关心的应是在领域层如何解析数据和函数，以便在并行运行环境中能正确、高效地解决问题。当进入编程阶段，你的关注点应转向如何组织并发线程。在这个阶段，你需要从逻辑层面来思考，以确保你的线程和计算能正确地解决问题。在C语言并行编程中，需要使用pthreads或OpenMP技术来显式地管理线程。CUDA提出了一个线程层次结构抽象的概念，以允许控制线程行为。在阅读本书中的示例时，你会发现这个抽象为并行编程提供了良好的可扩展性。在硬件层，通过理解线程是如何映射到核心可以帮助提高其性能。CUDA线程模型在不强调较低级别细节的情况下提供了充足的信息，具体内容详见第3章。
2.1.1　CUDA编程结构
CUDA编程模型使用由C语言扩展生成的注释代码在异构计算系统中执行应用程序。在一个异构环境中包含多个CPU和GPU，每个GPU和CPU的内存都由一条PCI-Express总线分隔开。因此，需要注意区分以下内容。
主机：CPU及其内存（主机内存）
设备：GPU及其内存（设备内存）
为了清楚地指明不同的内存空间，在本书的示例代码中，主机内存中的变量名以h_为前缀，设备内存中的变量名以d_为前缀。
从CUDA 6.0开始，NVIDIA提出了名为“统一寻址”（Unified Memory）的编程模型的改进，它连接了主机内存和设备内存空间，可使用单个指针访问CPU和GPU内存，无须彼此之间手动拷贝数据。更多细节详见第4章。现在，重要的是应学会如何为主机和设备分配内存空间以及如何在CPU和GPU之间拷贝共享数据。这种程序员管理模式控制下的内存和数据可以优化应用程序并实现硬件系统利用率的最大化。
内核（kernel）是CUDA编程模型的一个重要组成部分，其代码在GPU上运行。作为一个开发人员，你可以串行执行核函数。在此背景下，CUDA的调度管理程序员在GPU线程上编写核函数。在主机上，基于应用程序数据以及GPU的性能定义如何让设备实现算法功能。这样做的目的是使你专注于算法的逻辑（通过编写串行代码），且在创建和管理大量的GPU线程时不必拘泥于细节。
多数情况下，主机可以独立地对设备进行操作。内核一旦被启动，管理权立刻返回给主机，释放CPU来执行由设备上运行的并行代码实现的额外的任务。CUDA编程模型主要是异步的，因此在GPU上进行的运算可以与主机-设备通信重叠。一个典型的CUDA程序包括由并行代码互补的串行代码。如图2-2所示，串行代码（及任务并行代码）在主机CPU上执行，而并行代码在GPU上执行。主机代码按照ANSI C标准进行编写，而设备代码使用CUDA C进行编写。你可以将所有的代码统一放在一个源文件中，也可以使用多个源文件来构建应用程序和库。NVIDIA 的C编译器（nvcc）为主机和设备生成可执行代码。
一个典型的CUDA程序实现流程遵循以下模式。
1.把数据从CPU内存拷贝到GPU内存。

1. 调用核函数对存储在GPU内存中的数据进行操作。
2. 将数据从GPU内存传送回到CPU内存。
   首先，你要学习的是内存管理及主机和设备之间的数据传输。在本章后面你将学到更多GPU核函数执行的细节内容。

2.1.2　内存管理
CUDA编程模型假设系统是由一个主机和一个设备组成的，而且各自拥有独立的内存。核函数是在设备上运行的。为使你拥有充分的控制权并使系统达到最佳性能，CUDA运行时负责分配与释放设备内存，并且在主机内存和设备内存之间传输数据。表2-1列出了标准的C函数以及相应地针对内存操作的CUDA C函数。
用于执行GPU内存分配的是cudaMalloc函数，其函数原型为：

https://yqfile.alicdn.com/26818766f14b3d2f15d0787586ca8b19619515e8.png" >

cudaGetErrorString函数和C语言中的strerror函数类似。
CUDA编程模型从GPU架构中抽象出一个内存层次结构。图2-3所示的是一个简化的GPU内存结构，它主要包含两部分：全局内存和共享内存。第4章和第5章详细介绍了GPU内存层次结构的内容。
内存层次结构
CUDA编程模型最显著的一个特点就是揭示了内存层次结构。每一个GPU设备都有用于不同用途的存储类型。在第4章和第5章将会详细介绍。

在GPU内存层次结构中，最主要的两种内存是全局内存和共享内存。全局类似于CPU的系统内存，而共享内存类似于CPU的缓存。然而GPU的共享内存可以由CUDA C的内核直接控制。

而不是用：

应用程序在运行时将会崩溃。
为了避免这类错误，CUDA 6.0提出了统一寻址，使用一个指针来访问CPU和GPU的内存。有关统一寻址的内容详见第4章。

2.1.3　线程管理
当核函数在主机端启动时，它的执行会移动到设备上，此时设备中会产生大量的线程并且每个线程都执行由核函数指定的语句。了解如何组织线程是CUDA编程的一个关键部分。CUDA明确了线程层次抽象的概念以便于你组织线程。这是一个两层的线程层次结构，由线程块和线程块网格构成，如图2-5所示。

由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一线程块内的线程协作可以通过以下方式来实现。
同步
共享内存
不同块内的线程不能协作。
线程依靠以下两个坐标变量来区分彼此。
blockIdx（线程块在线程格内的索引）
threadIdx（块内的线程索引）
这些变量是核函数中需要预初始化的内置变量。当执行一个核函数时，CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标，你可以将部分数据分配给不同的线程。
该坐标变量是基于uint3定义的CUDA内置的向量类型，是一个包含3个无符号整数的结构，可以通过x、y、z三个字段来指定。

网格和线程块的维度
通常，一个线程格会被组织成线程块的二维数组形式，一个线程块会被组织成线程的三维数组形式。
线程格和线程块均使用3个dim3类型的无符号整型字段，而未使用的字段将被初始化为1且忽略不计。

在CUDA程序中有两组不同的网格和块变量：手动定义的dim3数据类型和预定义的uint3数据类型。在主机端，作为内核调用的一部分，你可以使用dim3数据类型定义一个网格和块的维度。当执行核函数时，CUDA运行时会生成相应的内置预初始化的网格、块和线程变量，它们在核函数内均可被访问到且为unit3类型。手动定义的dim3类型的网格和块变量仅在主机端可见，而unit3类型的内置预初始化的网格和块变量仅在设备端可见。
你可以通过代码清单2-2来验证这些变量如何使用。首先，定义程序所用的数据大小，为了对此进行说明，我们定义一个较小的数据。

把代码合并保存成名为checkDimension.cu的文件，如代码清单2-2所示。

对于一个给定的数据大小，确定网格和块尺寸的一般步骤为：
确定块的大小
在已知数据大小和块大小的基础上计算网格维度
要确定块尺寸，通常需要考虑：
内核的性能特性
GPU资源的限制
本书的后续章节会对以上几点因素进行详细介绍。代码清单2-3使用了一个一维网格和一个一维块来说明当块的大小改变时，网格的尺寸也会随之改变。

网格和块的维度存在几个限制因素，对于块大小的一个主要限制因素就是可利用的计算资源，如寄存器，共享内存等。某些限制可以通过查询GPU设备撤回。
网格和块从逻辑上代表了一个核函数的线程层次结构。在第3章中，你会发现这种线程组织方式能使你在不同的设备上有效地执行相同的程序代码，而且每一个线程组织具有不同数量的计算和内存资源。

2.1.4　启动一个CUDA核函数
你应该对下列C语言函数调用语句很熟悉：

之前所有的核函数调用完成后开始拷贝数据。当拷贝完成后，控制权立刻返回给主机端。
异步行为
不同于C语言的函数调用，所有的CUDA核函数的启动都是异步的。CUDA内核调用完成后，控制权立刻返回给CPU。

2.1.5　编写核函数
核函数是在设备端执行的代码。在核函数中，需要为一个线程规定要进行的计算以及要进行的数据访问。当核函数被调用时，许多不同的CUDA线程并行执行同一个计算任务。以下是用__global__声明定义核函数：

考虑一个简单的例子：将两个大小为N的向量A和B相加，主机端的向量加法的C代码如下：

2.1.6　验证核函数
既然你已经编写了核函数，你如何能知道它是否正确运行？你需要一个主机函数来验证核函数的结果。

验证核函数代码
除了许多可用的调试工具外，还有两个非常简单实用的方法可以验证核函数。
首先，你可以在Fermi及更高版本的设备端的核函数中使用printf函数。
其次，可以将执行参数设置为<<<1，1>>>，因此强制用一个块和一个线程执行核函数，这模拟了串行执行程序。这对于调试和验证结果是否正确是非常有用的，而且，如果你遇到了运算次序的问题，这有助于你对比验证数值结果是否是按位精确的。

2.1.7　处理错误
由于许多CUDA调用是异步的，所以有时可能很难确定某个错误是由哪一步程序引起的。定义一个错误处理宏封装所有的CUDA API调用，这简化了错误检查过程：

CHECK(cudaDeviceSynchronize())会阻塞主机端线程的运行直到设备端所有的请求任务都结束，并确保最后的核函数启动部分不会出错。以上仅是以调试为目的的，因为在核函数启动后添加这个检查点会阻塞主机端线程，使该检查点成为全局屏障。
2.1.8　编译和执行
现在把所有的代码放在一个文件名为sumArraysOnGPU-small-case.cu的文件中，如代码清单2-4所示。
代码清单2-4　基于GPU的向量加法（sumArraysOnGPU-small-case.cu）

在这段代码中，向量大小被设置为32，如下所示：

你需要确保一般情况下进行更改所产生结果的正确性。