3.2_第一个程序

3.2 第一个程序

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

3.2.1 Hello, World!

include"../common/book.h" int main(void）{ printf("Hello，World！\n"); return0;   
1

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

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

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

3.2.2 核函数调用

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

include<iostream>   
.globalvoid kernel(void）{   
}   
int main(void）{ kernel<<1,1>>(); printf("Hello，World！\n"); return0;   
}

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

  • 一个空的函数kernel()，并且带有修饰符__global__。

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

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

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

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

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

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

3.2.3 传递参数

前面提到过可以将参数传递给核函数，现在就来看一个示例。考虑下面对“Hello，

World！”应用程序的修改：

include<iostream>   
#include"book.h"   
__global__void add(int a,int b,int \*c）{  $\star c = a + b$  1   
}   
int main(void）{ intc; int\*dev_c; HANDLE_ERROR(cudaMalloc(void\*\*)&dev_c,sizeof(int)）); add<<<1,1>>>(2,7，dev_c);   
HANDLE_ERROR(cudaMemcpy(&c, dev_c, sizeof(int),udaMemcpyDeviceToHost））; printf("2+7=%d\n",c); JudaFree(dev_c); return0;   
}

注意这里增加了多行代码，在这些代码中包含两个概念：

  • 可以像调用C函数那样将参数传递给核函数。

  • 当设备执行任何有用的操作时，都需要分配内存，例如将计算值返回给主机。

在将参数传递给核函数的过程中没有任何特别之处。除了尖括号语法之外，核函数的外表和行为看上去与标准C中的任何函数调用一样。运行时系统负责处理将参数从主机传递给设备的过程中的所有复杂操作。

更需要注意的地方在于通过CUDAAlloc()来分配内存。这个函数调用的行为非常类似于标准的C函数malloc(), 但该函数的作用是告诉CUDA运行时在设备上分配内存。第一个参数是一个指针, 指向用于保存新分配内存地址的变量, 第二个参数是分配内存的大小。除了分配内存的指针不是作为函数的返回值外, 这个函数的行为与malloc()是相同的, 并且返回类型为void*。函数调用外层的HANDLE_ERROR()是我们定义的一个宏, 作为本书辅助代码的一部分。这个宏只是判断函数调用是否返回了一个错误值, 如果是的话, 那么将输出相应的错误消息, 退出应用程序并将退出码设置为EXIT_FAILURE。虽然你也可以在自己的应用程序中使用这个错误处理码, 但这种做法在产品级的代码中很可能是不够的。

这段代码引出了一个微妙但却重要的问题。CUDA C的简单性及其强大功能在很大程度上都是来源于它淡化了主机代码和设备代码之间的差异。然而，程序员一定不能在主机代码中对CUDA malloc()返回的指针进行解引用（Dereference）。主机代码可以将这个指针作为参数传递，对其执行算术运算，甚至可以将其转换为另一种不同的类型。但是，绝对不可以使用这个指针来读取或者写入内存。

遗憾的是，编译器无法防止这种错误的发生。如果能够在主机代码中对设备指针进行解引用，那么CUDA C将非常完美，因为这看上去就与程序中其他的指针完全一样了。我们可以将设备指针的使用限制总结如下：

可以将cudaMalloc()分配的指针传递给在设备上执行的函数。

可以在设备代码中使用CUDAAlloc()分配的指针进行内存读/写操作。

可以将CUDAAlloc()分配的指针传递给在主机上执行的函数。

不能在主机代码中使用CUDAAlloc()分配的指针进行内存读/写操作。

如果你仔细阅读了前面的内容，那么可以得出以下推论：不能使用标准C的free()函数来释放CUDAAlloc()分配的内存。要释放CUDAAlloc()分配的内存，需要调用CUDAFree()，这个函数的行为与free()的行为非常相似。

我们已经看到了如何在设备上分配内存和释放内存，同时也清楚地看到，在主机上不能对这块内存做任何修改。在示例程序中剩下来的两行代码给出了访问设备内存的两种最常见方法——在设备代码中使用设备指针以及调用CUDAMemcpy()。

设备指针的使用方式与标准C中指针的使用方式完全一样。语句*c = a + b的含义同样非常简单：将参数a和b相加，并将结果保存在c指向的内存中。这个计算过程非常简单，甚至吸引不了我们的兴趣。

在前面列出了在设备代码和主机代码中可以/不可以使用设备指针的各种情形。在主机指针的使用上有着类似的限制。虽然可以将主机指针传递给设备代码，但如果想通过主机指针来访问设备代码中的内存，那么同样会出现问题。总的来说，主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。

前面曾提到过，在主机代码中可以通过调用cudaMemcpy()来访问设备上的内存。这个函数调用的行为类似于标准C中的memcpy()，只不过多了一个参数来指定设备内存指针究竟是源指针还是目标指针。在这个示例中，注意cudaMemcpy()的最后一个参数为cudaMemcpyDeviceToHost，这个参数将告诉运行时源指针是一个设备指针，而目标指针是一个主机指针。

显然，cudaMemcpyHostToDevice将告诉运行时相反的含义，即源指针位于主机上，而目标指针是位于设备上。此外还可以通过传递参数cudaMemcpyDeviceToDevice来告诉运行时这

两个指针都是位于设备上。如果源指针和目标指针都位于主机上，那么可以直接调用标准C的memcpy()函数。

前の章3.1_本章目标
次の章3.3_查询设备