4 CUDA 环境搭建

4.1 简介

本章面向从未接触过CUDA的初学者。我们将依次介绍如何在不同操作系统上安装CUDA、有哪些可用的CUDA 工具以及CUDA如何编译代码，最后介绍应用程序接口提供的错误处理手段，并帮助读者识别CUDA代码和开发过程中必然碰到的应用程序接口报错。Windows、Mac 和 Linux 三大主流操作系统均支持CUDA。最易于使用和学习CUDA程序开发的操作系统，应该是你最熟悉的那个。对于零基础的初学者，Windows加上MicrosofVisual C++可能是最好选择。在 Windows和 Mac 上安装 CUDA 主要是一些点击操作，它们都提供了非常标准的集成开发环境，很适合CUDA程序开发。

4.2 在 Windows 下安装软件开发工具包

这里的安装以工具包 4.1版为例。在基于 Windows系统的个人计算机上安装CUDA需要一些组件，可以到英伟达开发者社区下载，入口在http://developer.nvidia.com/cudatoolkit-41。在本书付梓刊印之际，开发包5.0版已经处于待发布阶段。请到英伟达官网获取最新版本。

你需要事先安装好 Microsoft Visual Studio 2005、2008或者2010。接着首先要下载并安装对应于操作系统的英伟达开发驱动程序，下载地址同上。然后你还要依次下载并安装32位或 64位版本的 CUDA 工具包、GPU计算软件开发包及其软件开发包的示例程序。要确认你安装的程序版本号与你的操作系统是匹配的。建议的安装次序如下:

1)英伟达开发驱动程序

2)CUDA 工具包

3)CUDA软件开发工具包

4)GPU计算软件开发工具包

5)并行 Nsight 调试器

在 Windows7系统下，软件开发工具包把全部文件放置在ProgramData目录下。这个目录处于C盘，是隐藏的。为了查看其中的文件，可以通过桌面上的CUDA软件开发工具包图标进行浏览或者进入 Windows Folder Options(文件夹选项)对话框里进行设置，以允许查看隐藏文件，如图4-1所示，

图4-1 Folder Options 对话框里设置允许查看隐藏文件

4.3 Visual Studio

CUDA 支持 Visual Studio 版本的范围为 2005~2010,也支持大部分的学习版。学习版可以从Microsoft免费得到。专业版也能通过DreamSpark计划得到，需要在https:/kwww.dreamspark.com 网站注册为学生身份免费获得。

注册时，只需提供你的学校信息和身份编号。旦注册成功，就可以下载VisualStudio软件以及很多其他开发工具。这个计划不仅面向美国的学术机构:也涵盖全世界的大学生。

综合来看，Visual Studio 2008对CUDA的支持最好，它的编译速度比 Visual Studio 2010更快。但 Visual Studio 2010可以实现源代码的自动语法检查，这一特性非常实用。在使用一项未定义的类型时，它能使用红色下划线指明错误所在，与Microsoft Word 里提示拼写错误是一样的。对于明显的问题，这一特性极其有用，它将大大节省不必要的编译次数。因此建议使用 2010版本进行CUDA开发，特别是当你可以从 DreamSpark 获取免费版时。

4.3.1 工程

为了快速新建一个工程，你可以选择一个软件开发工具包示例作蓝本，然后移除其中不需要的工程文件，并插人自己的源文件。你的CUDA源代码，应该包含“.cu”扩展名，这样它的编译会采用英伟达编译器而不是 VisualC编译器。另一种新建工程的方式，可以通过工程模板向导，方便地建立一个基本的工程框架，细节将在后面看到。

4.3.2 64位用户

当使用 Windows 64位版本时，要注意一些工程文件默认设置为以 32位应用程序运行。因此，当尝试生成程序“时，你可能会收到以下错误消息:“致命错误LNK1181:无法打开输人文件'cutil32D.lib'”

这是因为没有安装该库的缘故。你很有可能只安装了对应64位Windows的64位版本软件开发工具包。要更正此问题，需要把目标平台从32位改为64位。可以使用VisuaStudio 的 Build(生成)菜单，然后把平台改变为X64，如图 4-2 所示。

图4-2 VisualC目标平台选择

当选择重新生成时，可能会提示你保存该工程。只需添加“X86”到工程名称之后并保存。该工程将生成为64位版本，并链接正确的库文件。链接时也可能碰到找不到库文件的问题，例如缺少“cutil32.lib”。在安装软件开发工具包的时候，它设置了一个环境变量,$(CUDA_LIB_PATH)。这个变量通常设置为:C:\ProgramFilesINVIDIA GPU Computing Toolkit\CUDAlv4.1VlibX64.

有时，可能在默认工程文件的路径设置中没有$(CUDA_LIB_PATH)这一项。要添加它，可以单击该工程，然后选择Projed(工程)→Properties(属性)菜单。这将弹出如图 4-3所示的对话框。

图 4-3 附加库路径

单击在最右侧的“…”按钮，会弹出一个对话框，可以在此处添加库路径(如图4-4所示)。只需添加“S(CUDA_LIB_PATH)”作为一个新行，该工程即可正常链接。

如果准备构建64位和32位两种CUDA应用程序，则需要提前安装32位和64位的CUDA软件开发工具包。对于软件开发工具包里的样例，要生成32位和64位两个版本的程序，也需要安装两种版本的软件开发工具包。

进入以下目录并生成解决方案文件，就可以建立所需的库文件:

C:\ProgramData\NVIDIA Corporation\NVIDIA GPUComputing SDK 4.1\C\common

C:\ProgramData\NVIDIA Corporation\NVIDIA GPUComputing SDK 4.1\shared

可以在 C:PrgramDataNNVIDIA Corporation\NVIDIA GPU Computing SDK4.1\Ccommonlibx64下找到所需的库文件。也可以手动添加它们到任何需要的工程中。遗憾的是，对于软件开发包里的样例，无法直接生成，所以它们无法自动生成所需的库文件，而二进制的库文件也没有随身附带，所以要生成软件开发包样例程序，会有些小麻烦。

4.3.3 创建工程

要创建支持 CUDA 的新应用程序，只需通过File(文件)→NEW(新建)→Project Wizard(工程向导)创建 CUDA应用程序，如图4-5所示。然后工程向导将创建包含“kernel.cu”文件的单个工程。“kemmel.cu”文件里混合了两种不同的代码，一种在 CPU端执行，另一种要在GPU端执行。GPU代码包含在 addKernel函数里。该函数的人口参数有三个，一个为指向目的数组的指针c，另外一对指针分别指向输人数组a和b。该函数对数组a和b的内容执行加法操作，结果存到目标数组c中。这个简单示例把执行CUDA 程序的框架展示了出来。

图4-5 CUDA 工程向导

其余的代码是一些基础代码，包括复制数据到GPU设备、调用内核函数、将数据从设备传回主机等。在起步阶段尝试这个简单工程，是大有裨益的，它可以让你实际感受一下在CUDA下编译工程的感觉。后文将涵盖CUDA程序运转所需的标准框架。研究这些代码很有好处，如果可能，尽量理解它们。即便这些代码对你来说难以理解，也没有关系。本书将逐步讲解如何编写CUDA程序。

4.4 Linux

多个 Linux发行版都支持CUDA。这些发行版的具体版本号取决于所安装CUDA 工具包的版本。支持CUDA 的 Linux 发行版如下:

• Fedora 14
• lRedhat6.0和5.5/CentOS 6.2(Redhat 的免费版)
• Ubuntu 11.04OpenSUSE 11.2

在 Linux平台上安装CUDA的第一步是确保你拥有最新的内核软件。在终端窗口里使用如下的命令达到这一目的:

sudo yum update

sudo 命令允许你以管理员身份登录，yum命令是标准的LinuxRPM软件包的安装工具。本条命令行只检查所有已安装的软件包有无可用的更新。这将确保你在安装任何驱动程序之前，系统都是最新的。许多基于图形用户界面的安装程序，不再使用旧版的命令行，而是自带图形用户界面的软件更新程序。

一旦内核更新到最新，运行如下命令

sudo yum install gcc-c++ kernel-devel

这将安装标准的 GNU C++环境以及重新编译内核所需的内核源代码。要注意，软件包名称是区分大小写的。这会提示你下载内容大约21MB，安装过程需要数分钟。同样，你也可以通过操作系统的图形用户界面安装程序来安装软件包。最后，如果你要绘制一些图形输出，就需要一个OpenGL开发环境。这一环境的安装用下面的命令:

sudo yum install freeglut-devel 1ibXi-devel libXmu-devel

现在，你已经准备好安装 CUDA驱动程序。确保你安装的CUDA工具包至少是4.1版。有许多方法来安装更新后的英伟达驱动程序。英伟达没有发布驱动程序的源代码，所以大多数 Linux 发行版，默认安装一个很基本的图形驱动程序。

内核驱动程序的安装(CentOS、Ubuntu 10.4)

每个CUDA发行版本，应使用特定的一组开发驱动程序。安装驱动程序的方式如果与这里列出的不一致，可能会导致CUDA不能正常工作。对特定版本的CUDA 工具包，请确认操作系统的版本是否支持它。某些Linux最新发行版未必支持;使用旧版本的发行版也极可能无法正常工作。因此，安装的第一步是根据你的Linux系统版本采用相应版本的驱动程序，替换现有的驱动程序，如图4-6所示。

图4-6截至2012年9月，支持的Linux下载和支持的驱动程序版本

一旦下载工作完成，你需要引导Linux进入纯文本模式。在 Windows下的安装，始终处于图形用户界面模式，而与此不同，在Linux下安装驱动程序需要文本模式。打开终端窗口(通常在图形用户界面的Systems菜单中)，使用下面的命令，可以让大多数系统在启动时进人文本模式:

sudo init 3

这将重启 Linux机器，并把它带人文本模式。如果要还原到图形模式，可以使用sudo init 5。

如果收到“User <user_name>is not in sudoers fle”的错误，请使用su命令，登录为root用户。编辑“/etc/sudoers”文件，添加如下命令:

your_user_name ALL=(ALL) ALL

注意要把上述命令中的your_user_name更换为你的登录用户名。对某些发行版(例如Ubuntu)，init模式对它失效，仍会引导进入图形用户界面。这里给出一种解决方法。在文本窗口里编辑GRUB启动文件:

sudo chmod +w /etc/default/grub

sudo nano /etc/default/grub

更改以下行:

GRUB_CMDLINE_LINUX_DEFAULT="quiet splash”

GRUB_CMDLINE LINUX DEFAULT=""

为:

#GRUB_CMDLINE_LINUX_DEFAULT="quiet splash"

GRUB_CMDLINE_LINUX_DEFAULT="text"

请使用如下命令更新 GRUB:

sudo update-grub

最后，重启机器，它应该可以运行在纯文本模式。成功安装驱动程序之后，使用之前的命令重新启动到图形用户界面。

下面，请定位到“.nm”文件所在地点，这个文件是你从英伟达网站下载得到的。然后键入:

sudo ^sh NVIDIA-Linux-x86 64-285.05.33.run

下载的驱动程序一定有不同的版本。安装开始之际，将要求你同意英伟达许可证，随后的安装过程需要几分钟。在此期间，安装程序会尝试使用所需英伟达驱动程序，替换默认的Nouveau驱动程序。如果安装程序询问是否继续，请选择“是”(Yes)。并不是每个版本的安装都能轻松自如，整个过程是很容易出错的。如果英伟达安装程序无法删除 Nouveau驱动那么就要把它列入黑名单，以便英伟达安装程序可以安装正确的驱动程序。当你正确安装了英伟达驱动程序，请输人:

sudo init 5

随后的机器重启会恢复到正常的图形模式，请参见前文针对Ubuntu 的说明。接下来的任务是安装工具包。根据操作系统发行版是Fedora、RedHat、Ubuntu、OpenSUSE还是SUSE，选择需要的工具包。与前面类似，只需定位到你安装软件开发包的路径，运行如下命令:

sudo sh<sdk_version>.run

其中的<sak_version>是你下载的文件。随后，它会安装所需的所有工具，并返回安装成功的信息。接着，它会提醒你必须更新PATH和LD_LIBRARY_PATH环境变量，这些修改需要你手工来做。为了达到上述目的，需要编辑“/etc/profle”启动文件。添加以下几行:

export PATH=/usr/local/cuda/bin:$PATH

export LD_LIBRARY_PATH=/usr/1ocal/cuda/ib:$LD_LIBRARY_PATH

请注意，你必须拥有修改该文件的权限。如果需要，请使用“sudo chmod +w/etc/profle赋予修改权限。你可以使用你喜欢的编辑器，达到类似“sudonano/etc/profle”命令的效果。请注销并重新登录，然后键人:

env

这将列出所有当前的环境变量设置。请检查刚加入的两条新条目。至此，CUDA已经安装到了“/usr/local/bin”文件夹下。

接着我们需要 GNU C++ 编译器。安装“g++”包到你的系统上，即可完成任务。“g++”包可以系统中在很多正在使用的软件安装程序中找到。

下一步是安装软件开发包样例代码。这些代码是我们生成和测试程序的素材。从英伟达网站下载和运行它们，同样使用sh sdk_version.run命令(sdk_version 改为你实际下载的文件名称)。请不要以root身份运行此软件开发包，否则后面对任何样例的生成将要求以 root 身份才能进行。

默认情况下，软件开发包将安装到用户账户区的一个子目录下。它可能会提示找不到 CUDA安装路径，并将使用默认目录(与之前 CUDA 安装的目录相同)。你可以安全地忽略此消息。一日 GPU计算软件开发包安装完成，你需要去“Common”子目录，运行make 以创建一组库文件。

这时，你可以生成软件开发包的样例程序。你将在Linux下执行你的第一个CUDA程序，并测试驱动程序是否工作正常。

4.5Mac

与其他平台一样，Macintosh 版本的工具也可从http://developer.nvidia.com/cudatoolkit-41下载。只需下载并按以下顺序安装软件包:

·开发驱动程序

·CUDA 工具包

·CUDA 工具软件开发包和代码样例

CUDA 4.1版本要求 Mac 系统的发行版本为10.6.8(雪豹)或更高。最新版本(10.7.x)或狮子版本可从苹果商店下载或从苹果单独购买。

软件开发包安装到“GPUComputing”目录下，其高层目录为“Developer”。只需浏览“ Developer/GPU Computing/C/bin/darwin/release”目录，你会发现预编译的可执行文件。运行 deviceQuery，这样有助于验证驱动程序安装和运行时环境是否正确。

为了编译样例程序，需要事先安装XCode。这是Mac下等价于GCC(GNUC编译器)的工具。XCode可以从苹果商店下载。这个产品是收费的，但对苹果开发者计划的用户是免费的。苹果开发者计划既包括 Macintosh的开发，也包括iPhone/iPad 应用程序的开发。在狮子版操作系统发布不久，XCode就允许狮子版系统拥有者免费下载。

当Xcode 安装完成，请打开一个终端窗口。先定位到Finder，打开 Utilities，然后双击终端窗口，键入以下命令:

cd /Developer/'GPU Computing/C/src/project

make-i

替换其中的 project为需要编译的某个软件开发包应用程序的名称。如果遇到编译错误可能是你没有下载XCode软件包或者目前的版本过旧。

4.6 安装调试器

在 Windows平台上，CUDA 提供了 Parallel Nsight 调试环境。它支持调试 CPU 和 GPU代码，并加亮那些运行不够高效的代码段。这也非常利于调试多线程应用程序。Nsight是极为有用的免费工具。它只要求你是CUDA开发者社区用户，注册过程也是完全免费的。一旦注册成功，你将能够从英伟达网站下载该工具。

请注意，你必须安装了VisualStudio2008或更高版本(不能是学习版)，同时要安装Service Pack1(SP1)。Nsight的发行说明中给出了SP1的下载链接。Parallel Nsight包含两个部分。一个是集成到 Visual Studio中的应用程序，如图 4-7所示;另一个是独立的监视程序。监视程序与主应用程序进行协作。监视程序一般与VsualStudio环境驻留在同一台机器上，但它们并不是必须在同一台机器上。

Parallel Nsight运行的最佳环境是配备两个支持CUDA的 GPU，一个GPU专门运行代码，另一个专门执行常规的显示。因此，运行着目标代码的GPU无法再用来显示。由于多数GPU卡具有双显示器输出，所以你可以通过配置双显示器设置，在显示卡上连接两个显示器。需要指出，在最新的2.2版本中，Parallel Nsight不再需要两个 GPU。

图4-7 Nsight集成到Microsoft Visual Studio 环境

它也可以通过配置调试工具从远程GPU获得数据。然而，在大多数情况下，额外购买-个低端GPU，并安装到个人计算机或者工作站上是更简单的做法。Windows平台下配置Parallel Nsight的第一步是禁用TDR功能(参见图4-8)。超时检测和恢复(Timeout Detection andRecovery，TDR)是 Windows系统的一种机制，用来检测底层驱动程序代码的异常崩溃情形如果驱动程序停止响应事件，Windows则会重置此驱动程序。鉴于在程序断点处，驱动程序将暂停响应，所以为了防止出现重置操作，TDR 功能需要关掉。

要设置此值，只需运行监视程序并单击程序对话框右下方的“ Nsight Monitor Options ”超链 接。这将弹出如图4-8所示的对话框。设置“WDDM TDR enabled”将通过修改注册表来禁用此功能。重启计算机，ParallelNsight将不再给出TDR 的启用警告。

图4-8 关闭Windows 内核的超时功能

为了在远程计算机上使用 ParallelNsight，只需在远程计算机的 Windows系统安装监视程序。当你第一次运行监视程序时，将收到 Windows 防火墙已阻止“公共网络”(基于互联网的)接入到监视程序的警告。这完全是你所希望的。但是，该工具需要接入局域网，所以你需要把这一例外加入到安装监视器程序的机器中已建立的防火墙规则里。针对局域节点，必须修正 TDR 问题并重启一次。

    接下来的步骤是在主机端运行 Visual Studio，并选择一个新的分析活动。你会看到窗口的顶部有一块类似图 4-9所示的区域。注意，“Connection Name”里填写的是localhost，意为本地机器。打开 Windows资源管理器查看局域网络上远程待调试计算机的计算机名称。用Windows 资源管理器中显示的名称替换localhost。然后按“Connect”(连接)按钮。你应该看到两个连接已经建立的证据，如图4-10所示。

图4-9 ParallelNsight远程连接设置窗口

图4-10 ParallelNsight远程已连接

证据一是“Connect”(连接)按钮变为“Disconnect”(未连接)。证据二是“Connection Status”(连接状态)对话框变成绿色，并显示目标机器上的所有GPU参见图4-11)。图中所示情况表明，我们连接到的测试计算机中装有五个GTX470 GPU 卡。  

    紧挨着“Connection Status"(连接状态)面板有一个”Application Control”(应用程序控制)面板，单击上面的“Launch”(启动)按钮，将远程启动目标机器上的应用程序。在此之前，所有必要的文件需要复制到远程机器上。这个过程是自动完成的，可能需要几秒钟。整体而言，分析和调试远程应用程序的这一方式非常简单。

    假设你有一台笔记本电脑，并希望调试，或者干脆在GPU服务器上远程运行应用程序，那么上述配置 Parallel Nsight的方式正是你需要的。这种用法包括多人在不同时刻共享使用GPU服务器，或者在课堂教学中使用。也可能远程开发者需要在特设的测试服务器上运行代码，而这些服务器存有大量数据，把这些数据传回本地开发机器反而不明智。这同时意味着你不需要在每个远程服务器上安装 Visual C++。

针对 Linux和 Mac 平台的调试环境是 CUDA-GDB，它提供了一个扩展的 GNU 调试程序包。与 Parallel Nsight一样，它允许调试主机代码和 CUDA代码，可以在 CUDA代码中设置断点、单步跟踪和选择调试线程等。安装软件开发包时，CUDA-GDB和VisualProfler 工具是默认同时安装的，而不像 Parallel Nsight需要单独下载。截至 2012年，适用于Linux平台下 Eclipse环境的 Parallel Nsight 也发布了。

Windows和 Mac/Linux平台间的主要区别在于对分析工具(profling tool)的支持。Parallel Nsight 工具在这方面大大优于 Visual Profler。Visual Profler 也有 Windows 版本。它针对代码中的问题，提供了比较抽象的概述和建议，比较适合CUDA初学者。相比之下Parallel Nsight则针对更高级的用户。后续章节会分别介绍Parallel Nsight和 Visual Profler。然而，本书主要使用 ParallelNsight作为GPU开发的调试/分析工具。

对于高级的 CUDA开发任务，强烈建议使用 Parallel Nsight进行调试和分析。对于大多数 CUDA 新手来说，结合 Visual Profler 和 CUDA-GDB 也足以满足开发需要。

4.7 编译模型

英伟达编译器 NVCC运行在后台，当CUDA源文件请求编译时调用。CUDA源文件或常规源文件的不同类型和含义在表 4-1给出。文件扩展名决定编译时是使用NVCC还是主机编译器。

生成的可执行文件或胖二进制文件，针对不同代次的GPU会包含一个或多个二进制可执行映像。它还包含了PTX映像，允许CUDA运行时进行实时(Just-In-Tine，JIT)编译。这跟Java 字节码非常相似，它针对一个虚拟的架构，在实际目标硬件上的编译环节在运行程序时进行。PTX的实时编译仅在可执行文件不包含当前运行GPU的二进制映像时发生。因此，所有未来的架构与基本的虚拟架构都是向后兼容的。即使没有在某个GPU上执行编译的程序，通过在运行时编译嵌入在可执行文件中的PTX代码，也可以合法运行在该GPU上。

跟 Java类似，CUDA编译支持代码托管。定义环境变量CUDA DEVCODECACHE 指向一个目录，该目录将指示运行时把编译的二进制文件保存以备后用。通常情况下，每次编译需要把 PTX 代码根据未知 GPU的类型进行转换。采用代码托管，这种启动延时就可以避免了。

后面的章节将介绍如何查看真正的目标汇编代码，这是PTX翻译为目标代码的结果。

表 4-1 不同的 CUDA 文件类型

4.8 错误处理

CUDA的错误处理，乃至C语言的错误处理，都是不够好的。它们仅会执行少数运行时检查，对于很多欠考虑的事情，运行时通常也不会阻止。这会导致GPU程序奇怪的退出。幸运的时候，你会收到类似编译错误的错误信息。随着这方面经验的增多，你才能学会错误信息的具体含义。

几乎所有的 CUDA函数调用，都会返回类型为cudaError_t的整数值。cudaSuccess以外的任何值将代表致命错误。这通常是由于你的程序在使用前没有正确设置，或使用了已经销毁的对象。这也可能是源于Microsoft Windows中的GPU内核超时。如果你安装的 ParallelNsight 工具没有禁用TDR，当内核函数运行时间超过数秒钟，就会导致GPU内核超时(参见前文介绍)。内存越界访问也可能产生异常，输出各种错误信息到stderr(标准错误输出)。

每个函数返回一个错误代码时，每次函数调用必须进行错误检查，并编写处理函数。这使得错误处理很麻烦，并导致高度缩进的程序代码。例如:

if(cudaMal1oc(...)==cudaSuccess)

{

if(cudaEventCreate(&event)==cudaSucess)

{

        …

}

}

else

{

…

}

为了避免这种重复性的编程，本书始终使用下面的宏定义来调用CUDA应用程序接口:

#define CUDA CALL(x){const cudaError_t a=(x):if(a!=cudasuccess){printf("\nCuDAError:%s(err_num=%d)\n",cudaGetErrorString(a),a):cudaDeviceReset();assert(0);}}

这个宏允许你指定x为一些函数调用，例如:

CUDA_CALL(cudaEventCreate(&kernel_start));

接着创建一个临时变量a并把cudaError_t类型的函数返回值赋予它。然后检查这个值是否等于 cudasuccess，如果不等于则说明函数调用碰到了错误。如果检测到错误，它会在屏幕上打印返回的错误外加简短的错误解释。它还使用assert宏，用来识别发生错误时所在的源文件及代码行，因此有助于你轻松地定位检测到的错误位置。

这种技术适用于除了内核函数调用之外的所有CUDA调用。内核函数是在GPU设备上运行的程序。它们的执行使用<<<和>>>，如下所示:

my_kernel<<<num_blocks, num_threads>>>(paraml, param2,...);

对于内核的错误检查，我们将使用下面的函数:

__host__ void cuda_error_check(const char * prefix, const char * postfix)

{

if(cudaPeekAtLastError()!=cudaSuccess)

{

printf("\n%s%s%s", prefix, cudaGetErrorString(cudaGetLastError()), postfix);

cudaDeviceReset();

wait_exit();

exit(1);

}

}

此函数应该在内核函数调用后立即执行。它会检查任何即时错误。一旦发现错误，则输出错误信息、复位 GPU，其中的 wait_exit 函数会等待用户的任一按键操作，然后退出程序。

请注意，内核函数调用与CPU代码是异步执行的，所以上述方法并非万无一失。异步执行表示当我们调用cudaPeekAtLastError时，GPU代码正在后台运行。如果此时没有检测到错误，则不会输出错误，函数继续执行下面的代码行。通常情况下，下一条代码行是从GPU内存将数据复制到CPU内存的操作。内核函数的错误可能会导致随后的应用程序接口调用失败，一般应用程序接口的调用是紧跟内核函数调用的。针对全部的应用程序接口调用，均使用CUDA_CALL包裹，这种错误就会被标记出来。

也可以强制在内核函数完成后再进行错误检查，只需在cudaPeekAtLastError调用之前加人 cudaDeviceSynchronize调用即可。然而，这一强制行为只能在调试版程序或者想让CPU在GPU占满时处于闲置状态时使用。学完本书，你就能明白这种同步操作适合进行调试，但会影响性能。所以，如果这些调用只是为了调试，请务必在产品代码中删去。

4 9 本章小结

现在你成功安装了CUDA软件开发包，其中包括GPU计算软件开发包样例和调试环境。你学会了生成一个简单的GPU计算软件开发包样例，比如deviceQuery工程，并在运行时利用它识别系统中的 GPU。