CUDA 中的线程组织

明朝那些事中有一句话：我之所以写徐霞客是想告诉你，所谓千秋霸业万古流芳，与一件事相比，其实都算不了什么，这件事情就是——用你喜欢的方式度过一生。

我们以最简单的 CUDA 程序：从 GPU 中输出 Hello World! 字符串开始 CUDA 编程的学习。

经典的 Hello World 程序几乎是学习任何一门新编程语言的出发点。

学会了 HelloWorld 程序的开发过程，就对一个新的编程语言有了一个初步的认识。

本书的所有范例都是基于 Linux 操作系统开发的，但大部分也在 Windows 操作系统中使用 Command Prompt 命令行通过测试。因此，读者需要掌握基本的 Linux 或 Windows 命令行操作知识。

C++ 语言中的 Hello World 程序

学习 CUDA C++ 编程需要读者比较熟练地掌握 C++ 编程的基础。

虽然 CUDA 支持很多 C++ 的特征，但作者写的 C++ 程序有很多 C 程序的痕迹，而且本书基本上不涉及 C++ 中的类和模板等编程特征。

我们先回顾一下 C++ 中 Hello World 程序的开发过程。在 C++ 语言中开发一个程序的大致过程如下：

• 1. 用文本编辑器写一个源代码（source code）。
• 2. 用编译器对源代码进行预处理、编译、汇编并链接必要的目标文件得到可执行文件（executable）。这些步骤往往可由一个命令完成。
• 3. 运行可执行文件得到结果。

1 #include <stdio.h>
2
3 int main(void)
4 {
5 printf("Hello World!\n");
6 return 0;
7 }

首先，让我们用编辑器写下 Listing 2.1 中的源代码。然后，将程序的文件命名为 hello.cpp，并用 g++ 编译（如上所述，此处及后面所说的编译其实包含了预处理、编译、汇编、链接等步骤）：

首先，让我们用编辑器写下 Listing 2.1 中的源代码。然后，将程序的文件命名为 hello.cpp，并用 g++ 编译（如上所述，此处及后面所说的编译其实包含了预处理、编译、汇编、链接等步骤）：

$ g++ hello.cpp

编译通过后，将得到一个名为 a.out 的可执行文件。用如下命令执行该文件：

$ ./a.out

接着，就可以看到屏幕上打印出如下文字：

Hello World!

也可以在编译时指定二进制文件的名字。例如，用如下命令：

$ g++ hello.cpp -o hello

将得到一个名为 hello 的可执行文件，可以用如下命令运行它：

$ ./hello

以上假定使用了 GCC 编译器套装。如果使用 Windows 下的 MSVC 编译器套装，则可用 cl 编译程序：

$ cl hello.cpp

这将产生一个名为 hello.exe 的可执行文件。

CUDA 中的 Hello World 程序

在复习了 C++ 语言中的 Hello World 程序之后，我们接着介绍 CUDA 中的 Hello World 程序。

只有主机函数的 CUDA 程序

其实，我们已经写好了一个 CUDA 中的 Hello World 程序。这是因为，CUDA 程序的编译器驱动（compiler driver）nvcc 支持编译纯粹的 C++ 代码。

一般来说，一个标准的 CUDA 程序中既有纯粹的 C++ 代码，也有不属于 C++ 的真正的 CUDA 代码。CUDA 程序的编译器

驱动 nvcc 在编译一个 CUDA 程序时，会将纯粹的 C++ 代码交给 C++ 的编译器（如前面提到的 g++ 或 cl）去处理，它自己则负责编译剩下的部分。CUDA 程序源文件的后缀名默认是 .cu，所以我们可以将上面写好的源文件更名为 hello1.cu，然后用 nvcc 编译：

$ nvcc hello1.cu

编译好之后即可运行。运行结果与 C++ 程序的运行结果一样。关于 CUDA 程序的编译过程，将在本章最后一节及后续的某些章节详细讨论，现在只要知道可以用 nvcc 编译 CUDA 程序即可。

使用核函数的 CUDA 程序

虽然上面的第一个版本是由 CUDA 的编译器编译的，但程序中根本没有使用 GPU。下面来介绍一个使用 GPU 的 Hello World 程序。

首先，我们要知道，GPU 只是一个设备，要它工作的话还需要有一个主机给它下达命令。这个主机就是 CPU。

所以，一个真正利用了 GPU 的 CUDA 程序既有主机代码（在程序 hello1.cu 中的所有代码都是主机代码），也有设备代码（可以理解为需要设备执行的代码）。

主机对设备的调用是通过核函数（kernel function）来实现的。所以，一个典型的、简单的 CUDA 程序的结构具有下面的形式：

int main(void)
{
主机代码
核函数的调用
主机代码
return 0;
}

CUDA 中的核函数与 C++ 中的函数是类似的，但一个显著的差别是：它必须被限定词（qualifier）global 修饰。

其中 global 前后是双下划线。另外，核函数的返回类型必须是空类型，即 void。

这两个要求读者先记住即可。关于核函数的更多细节，以后再逐步深入介绍。遵循这两个要求，我们先写一个打印字符串的核函数：

__global__ void hello_from_gpu()
{
	printf("Hello World from the GPU!\n");
}

限定符 global 和 void 的次序可随意。也就是说，上述核函数也可以写为如下形式：

这里是引用

void __global__ hello_from_gpu()
{
	printf("Hello World from the GPU!\n");
}