GPU高性能编程CUDA实战-第3章.pdf

第3章
CUDAC简介

如果你学习了第1章的内容，那么我们认为你已经了解了图形处理器的
强大计算能力，并且需要在程序中使用这种能力。此外，如果你学习了第2
章的内容，那么就应该配置好了编译和运行CUDAC代码的开发环境。如
果你跳过了这两章的内容，例如只是想浏览代码示例，或者只是在书店里
随机地打开这本书翻到这页，或者可能迫不及待地想立即开始动手编写和
运行代码，那么也没关系。无论是何种情况，你都应该准备好了编写第一
个代码示例，那么让我们开始吧。
16GPU高性能编程CUDA实战

3.1本章目标

通过本章的学习，你可以：

•编写第一段CUDAC代码。
•了解为主机（Host）编写的代码与为设备（Device）编写的代码之间的区别。
•如何从主机上运行设备代码。
•了解如何在支持CUDA的设备上使用设备内存。
•了解如何查询系统中支持CUDA的设备的信息。

3.2第一个程序

我们希望通过示例来学习CUDAC，因此来看第一个CUDAC示例。为了保持计算机编程
书籍的行文风格，我们首先给出的是一个“Hello,World!”示例。

3.2.1Hello,World!

#include"../common/book.h"

intmain(void){
printf("Hello,World!\n");
return0;
}

当看到这段代码时，你肯定在怀疑本书是不是一个骗局。这不就是C吗？CUDAC是不是
真的存在？这些问题的答案都是肯定的。当然，本书也不是一个骗局。这个简单的“Hello，
World！”示例只是为了说明，CUDAC与你熟悉的标准C在很大程度上是没有区别的。

这个示例很简单，它能够完全在主机上运行。然而，这个示例引出了本书的一个重要区分：
我们将CPU以及系统的内存称为主机，而将GPU及其内存称为设备。这个示例程序与你编写过
的代码非常相似，因为它并不考虑主机之外的任何计算设备。

为了避免使你产生一无所获的感觉，我们将逐渐完善这个简单示例。我们来看看如何使用
GPU（这就是一个设备）来执行代码。在GPU设备上执行的函数通常称为核函数（Kernel）。

3.2.2核函数调用

现在，我们在示例程序中添加一些代码，这些代码比最初的“Hello，World！”程序看上
去会陌生一些。
第3章CUDAC简介17

#include<iostream>

_global_voidkernel(void){
}

intmain(void){
kernel<<<1,1>>>();
printf("Hello,World!\n");
return0;
}

这个程序与最初的“Hello，World！”相比，多了两个值得注意的地方：

•一个空的函数kernel()，并且带有修饰符__global__。
•对这个空函数的调用，并且带有修饰字符<<<1,1>>>。

在上一节中看到，代码默认是由系统的标准C编译器来编译的。例如，在Linux操作系统上
用GNUgcc来编译主机代码，而在Windows系统上用MicrosoftVisualC来编译主机代码。
NVIDIA工具只是将代码交给主机编译器，它表现出的行为就好像CUDA不存在一样。

现在，我们看到了CUDAC为标准C增加的__global__修饰符。这个修饰符将告诉编译器，
函数应该编译为在设备而不是主机上运行。在这个简单的示例中，函数kernel()将被交给编译设
备代码的编译器，而main()函数将被交给主机编译器（与上一个例子一样）。

那么，kernel()的调用究竟代表着什么含义，并且为什么必须加上尖括号和两个数值？注意，
这正是使用CUDAC的地方。

我们已经看到，CUDAC需要通过某种语法方法将一个函数标记为“设备代码（Device
Code）”。这并没有什么特别之处，而只是一种简单的表示方法，表示将主机代码发送到一个编
译器，而将设备代码发送到另一个编译器。事实上，这里的关键在于如何在主机代码中调用设
备代码。CUDAC的优势之一在于，它提供了与C在语言级别上的集成，因此这个设备函数调
用看上去非常像主机函数调用。在后面将详细介绍在这个函数调用背后发生的动作，但就目前
而言，只需知道CUDA编译器和运行时将负责实现从主机代码中调用设备代码。

因此，这个看上去有些奇怪的函数调用实际上表示调用设备代码，但为什么要使用尖括号
和数字？尖括号表示要将一些参数传递给运行时系统。这些参数并不是传递给设备代码的参数，
而是告诉运行时如何启动设备代码。在第4章中，我们将了解这些参数对运行时的作用。传递
给设备代码本身的参数是放在圆括号中传递的，就像标准的函数调用一样。

3.2.3传递参数

前面提到过可以将参数传递给核函数，现在就来看一个示