编译CUDA Samples工程 cuda c编程权威指南

第2章

CUDA编程模型

本章内容：

• 写一个CUDA程序
• 执行一个核函数
• 用网格和线程块组织线程
• GPU性能测试

CUDA是一种通用的并行计算平台和编程模型，是在C语言基础上扩展的。借助于CUDA，你可以像编写C语言程序一样实现并行算法。你可以在NVIDIA的GPU平台上用CUDA为多种系统编写应用程序，范围从嵌入式设备、平板电脑、笔记本电脑、台式机、工作站到HPC集群（高性能计算集群）。熟悉C语言编程工具有助于在整个项目周期中编写、调试和分析你的CUDA程序。在本章中，我们将通过向量加法和矩阵加法这两个简单的例子来学习如何编写一个CUDA程序。

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内存。
2. 调用核函数对存储在GPU内存中的数据进行操作。
3. 将数据从GPU内存传送回到CPU内存。

首先，你要学习的是内存管理及主机和设备之间的数据传输。在本章后面你将学到更多GPU核函数执行的细节内容。

2.1.2　内存管理

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

用于执行GPU内存分配的是cudaMalloc函数，其函数原型为：

该函数负责向设备分配一定字节的线性内存，并以devPtr的形式返回指向所分配内存的指针。cudaMalloc与标准C语言中的malloc函数几乎一样，只是此函数在GPU的内存里分配内存。通过充分保持与标准C语言运行库中的接口一致性，可以实现CUDA应用程序的轻松接入。

cudaMemcpy函数负责主机和设备之间的数据传输，其函数原型为：

此函数从src指向的源存储区复制一定数量的字节到dst指向的目标存储区。复制方向由kind指定，其中的kind有以下几种。

这个函数以同步方式执行，因为在cudaMemcpy函数返回以及传输操作完成之前主机应用程序是阻塞的。除了内核启动之外的CUDA调用都会返回一个错误的枚举类型cuda Error_t。如果GPU内存分配成功，函数返回：

否则返回：

可以使用以下CUDA运行时函数将错误代码转化为可读的错误消息：

cudaGetErrorString函数和C语言中的strerror函数类似。

CUDA编程模型从GPU架构中抽象出一个内存层次结构。图2-3所示的是一个简化的GPU内存结构，它主要包含两部分：全局内存和共享内存。第4章和第5章详细介绍了GPU内存层次结构的内容。

这是一个纯C语言编写的程序，你可以用C语言编译器进行编译，也可以像下面这样用nvcc进行编译。

nvcc封装了几种内部编译工具，CUDA编译器允许通过命令行选项在不同阶段启动不同的工具完成编译工作。-Xcompiler用于指定命令行选项是指向C编译器还是预处理器。在前面的例子中，将-std=c99传递给编译器，因为这里的C程序是按照C99标准编写的。你可以在CUDA编译器文件中找到编译器选项（http：//docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html）。

现在，你可以在GPU上修改代码来进行数组加法运算，用cudaMalloc在GPU上申请内存。

使用cudaMemcpy函数把数据从主机内存拷贝到GPU的全局内存中，参数cudaMemc-pyHostToDevice指定数据拷贝方向。

当数据被转移到GPU的全局内存后，主机端调用核函数在GPU上进行数组求和。一旦内核被调用，控制权立刻被传回主机，这样的话，当核函数在GPU上运行时，主机可以执行其他函数。因此，内核与主机是异步的。

当内核在GPU上完成了对所有数组元素的处理后，其结果将以数组d_C的形式存储在GPU的全局内存中，然后用cudaMemcpy函数把结果从GPU复制回到主机的数组gpuRef中。

cudaMemcpy的调用会导致主机运行阻塞。cudaMemcpyDeviceToHost的作用就是将存储在GPU上的数组d_c中的结果复制到gpuRef中。最后，调用cudaFree释放GPU的内存。

2.1.3　线程管理

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

由一个内核启动所产生的所有线程统称为一个网格。同一网格中的所有线程共享相同的全局内存空间。一个网格由多个线程块构成，一个线程块包含一组线程，同一线程块内的线程协作可以通过以下方式来实现。

• 同步
• 共享内存

不同块内的线程不能协作。

线程依靠以下两个坐标变量来区分彼此。

• blockIdx（线程块在线程格内的索引）
• threadIdx（块内的线程索引）

这些变量是核函数中需要预初始化的内置变量。当执行一个核函数时，CUDA运行时为每个线程分配坐标变量blockIdx和threadIdx。基于这些坐标，你可以将部分数据分配给不同的线程。

该坐标变量是基于uint3定义的CUDA内置的向量类型，是一个包含3个无符号整数的结构，可以通过x、y、z三个字段来指定。

CUDA可以组织三维的网格和块。图2-5展示了一个线程层次结构的示例，其结构是一个包含二维块的二维网格。网格和块的维度由下列两个内置变量指定。

• blockDim（线程块的维度，用每个线程块中的线程数来表示）
• gridDim（线程格的维度，用每个线程格中的线程数来表示）

它们是dim3类型的变量，是基于uint3定义的整数型向量，用来表示维度。当定义一个dim3类型的变量时，所有未指定的元素都被初始化为1。dim3类型变量中的每个组件可以通过它的x、y、z字段获得。如下所示。

在CUDA程序中有两组不同的网格和块变量：手动定义的dim3数据类型和预定义的uint3数据类型。在主机端，作为内核调用的一部分，你可以使用dim3数据类型定义一个网格和块的维度。当执行核函数时，CUDA运行时会生成相应的内置预初始化的网格、块和线程变量，它们在核函数内均可被访问到且为unit3类型。手动定义的dim3类型的网格和块变量仅在主机端可见，而unit3类型的内置预初始化的网格和块变量仅在设备端可见。

你可以通过代码清单2-2来验证这些变量如何使用。首先，定义程序所用的数据大小，为了对此进行说明，我们定义一个较小的数据。

接下来，定义块的尺寸并基于块和数据的大小计算网格尺寸。在下面的例子中，定义了一个包含3个线程的一维线程块，以及一个基于块和数据大小定义的一定数量线程块的一维线程网格。

你会发现网格大小是块大小的倍数。在下一章中你会了解必须这样计算网格大小的原因。以下主机端上的程序段用来检查网格和块维度。

在核函数中，每个线程都输出自己的线程索引、块索引、块维度和网格维度。

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

现在开始编译和运行这段程序：

因为printf函数只支持Fermi及以上版本的GPU架构，所以必须添加-arch=sm_20编译器选项。默认情况下，nvcc会产生支持最低版本GPU架构的代码。这个应用程序的运行结果如下。可以看到，每个线程都有自己的坐标，所有的线程都有相同的块维度和网格维度。

对于一个给定的数据大小，确定网格和块尺寸的一般步骤为：

• 确定块的大小
• 在已知数据大小和块大小的基础上计算网格维度

要确定块尺寸，通常需要考虑：

• 内核的性能特性
• GPU资源的限制

本书的后续章节会对以上几点因素进行详细介绍。代码清单2-3使用了一个一维网格和一个一维块来说明当块的大小改变时，网格的尺寸也会随之改变。

用下列命令编译和运行这段程序：

下面是一个输出示例。由于应用程序中的数据大小是固定的，因此当块的大小发生改变时，相应的网格尺寸也会发生改变。

2.1.4　启动一个CUDA核函数

你应该对下列C语言函数调用语句很熟悉：

CUDA内核调用是对C语言函数调用语句的延伸，<<<>>>运算符内是核函数的执行配置。

正如上一节所述，CUDA编程模型揭示了线程层次结构。利用执行配置可以指定线程在GPU上调度运行的方式。执行配置的第一个值是网格维度，也就是启动块的数目。第二个值是块维度，也就是每个块中线程的数目。通过指定网格和块的维度，你可以进行以下配置：

• 内核中线程的数目
• 内核中使用的线程布局

同一个块中的线程之间可以相互协作，不同块内的线程不能协作。对于一个给定的问题，可以使用不同的网格和块布局来组织你的线程。例如，假设你有32个数据元素用于计算，每8个元素一个块，需要启动4个块：

图2-6表明了上述配置下的线程布局。

由于数据在全局内存中是线性存储的，因此可以用变量blockIdx.x和threadId.x来进行以下操作。

• 在网格中标识一个唯一的线程
• 建立线程和数据元素之间的映射关系

如果把所有32个元素放到一个块里，那么只会得到一个块：

如果每个块只含有一个元素，那么会有32个块：

核函数的调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端。你可以调用以下函数来强制主机端程序等待所有的核函数执行结束：

一些CUDA运行时API在主机和设备之间是隐式同步的。当使用cudaMemcpy函数在主机和设备之间拷贝数据时，主机端隐式同步，即主机端程序必须等待数据拷贝完成后才能继续执行程序。

之前所有的核函数调用完成后开始拷贝数据。当拷贝完成后，控制权立刻返回给主机端。

2.1.5　编写核函数

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

核函数必须有一个void返回类型。

表2-2总结了CUDA C程序中的函数类型限定符。函数类型限定符指定一个函数在主机上执行还是在设备上执行，以及可被主机调用还是被设备调用。

__device__和__host__限定符可以一齐使用，这样函数可以同时在主机和设备端进行编译。

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

这是一个迭代N次的串行程序，循环结束后将产生以下核函数：

C函数和核函数之间有什么不同？你可能已经注意到循环体消失了，内置的线程坐标变量替换了数组索引，由于N是被隐式定义用来启动N个线程的，所以N没有什么参考价值。

假设有一个长度为32个元素的向量，你可以按以下方法用32个线程来调用核函数：

2.1.6　验证核函数

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

2.1.7　处理错误

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

例如，你可以在以下代码中使用宏：

如果内存拷贝或之前的异步操作产生了错误，这个宏会报告错误代码，并输出一个可读信息，然后停止程序。也可以用下述方法，在核函数调用后检查核函数错误：

CHECK(cudaDeviceSynchronize())会阻塞主机端线程的运行直到设备端所有的请求任务都结束，并确保最后的核函数启动部分不会出错。以上仅是以调试为目的的，因为在核函数启动后添加这个检查点会阻塞主机端线程，使该检查点成为全局屏障。

2.1.8　编译和执行

现在把所有的代码放在一个文件名为sumArraysOnGPU-small-case.cu的文件中，如代码清单2-4所示。

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

执行配置被放入一个块内，其中包含32个元素：

使用以下命令编译和执行该代码：

系统报告结果如下：

如果你将执行配置重新定义为32个块，每个块只有一个元素，如下所示：

那么就需要在代码清单2-4中对核函数sumArraysOnGPU进行修改：

一般情况下，可以基于给定的一维网格和块的信息来计算全局数据访问的唯一索引：

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