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

《GPU高性能编程CUDA实战-第3章.pdf》由会员上传分享，免费在线阅读，更多相关内容在行业资料-天天文库。

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

2、（Device）编写的代码之间的区别。•如何从主机上运行设备代码。•了解如何在支持CUDA的设备上使用设备内存。•了解如何查询系统中支持CUDA的设备的信息。3.2第一个程序我们希望通过示例来学习CUDAC，因此来看第一个CUDAC示例。为了保持计算机编程书籍的行文风格，我们首先给出的是一个“Hello,World!”示例。3.2.1Hello,World!#include"../common/book.h"intmain(void){printf("Hello,World!");return0;}当看到这段代码时，你肯定在怀疑本书是不是一个骗局。这不就是C吗？CU

3、DAC是不是真的存在？这些问题的答案都是肯定的。当然，本书也不是一个骗局。这个简单的“Hello，World！”示例只是为了说明，CUDAC与你熟悉的标准C在很大程度上是没有区别的。这个示例很简单，它能够完全在主机上运行。然而，这个示例引出了本书的一个重要区分：我们将CPU以及系统的内存称为主机，而将GPU及其内存称为设备。这个示例程序与你编写过的代码非常相似，因为它并不考虑主机之外的任何计算设备。为了避免使你产生一无所获的感觉，我们将逐渐完善这个简单示例。我们来看看如何使用GPU（这就是一个设备）来执行代码。在GPU设备上执行的函数通常称为核函数（Kernel）。3.

4、2.2核函数调用现在，我们在示例程序中添加一些代码，这些代码比最初的“Hello，World！”程序看上去会陌生一些。第3章　CUDAC简介17#include_global_voidkernel(void){}intmain(void){kernel<<<1,1>>>();printf("Hello,World!");return0;}这个程序与最初的“Hello，World！”相比，多了两个值得注意的地方：•一个空的函数kernel()，并且带有修饰符__global__。•对这个空函数的调用，并且带有修饰字符<<<1,1>>>。在上一节中看

5、到，代码默认是由系统的标准C编译器来编译的。例如，在Linux操作系统上用GNUgcc来编译主机代码，而在Windows系统上用MicrosoftVisualC来编译主机代码。NVIDIA工具只是将代码交给主机编译器，它表现出的行为就好像CUDA不存在一样。现在，我们看到了CUDAC为标准C增加的__global__修饰符。这个修饰符将告诉编译器，函数应该编译为在设备而不是主机上运行。在这个简单的示例中，函数kernel()将被交给编译设备代码的编译器，而main()函数将被交给主机编译器（与上一个例子一样）。那么，kernel()的调用究竟代表着什么含义，并且为什么必

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

7、函数调用实际上表示调用设备代码，但为什么要使用尖括号和数字？尖括号表示要将一些参数传递给运行时系统。这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备代码。在第4章中，我们将了解这些参数对运行时的作用。传递给设备代码本身的参数是放在圆括号中传递的，就像标准的函数调用一样。3.2.3传递参数前面提到过可以将参数传递给核函数，现在就来看一个示例。考虑下面对“Hello，18GPU高性能编程CUDA实战World！”应用程序的修改：#include#include"book.h"_global_voidadd(i