异构计算(一)

3.1 本章目标

- 编写第一段CUDA代码

- 了解为主机(Host)编写的代码与为设备(Deivce)编写的代码之间的区别

- 如何从主机上运行设备代码

- 了解如何在支持CUDA的设备上使用设备内存

- 了解如何查询系统中支持CUDA的设备信息

3.2 第一个程序

 

 

这个例子很简单，完全可以在主机上运行，我们将CPU以及系统的内存称为主机，而将GPU及其内存称为设备，这个例子与普通代码非常相似，因为不考虑主机之外的任何计算设备。

GPU是一个设备，可以执行代码，在GPU设备上执行的函数通常称为核函数。

3.2.2 核函数调用

 

添加一些代码，多了两个值得注意的地方

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

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

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

CUDA C需要通过某种语法方法将一个函数标记为“设备代码”，这是一种简单的表示方法，表示将主机代码发送到一个编译器，而将设备代码发送到另一个编译器。这里的关键在于如何咋主机代码中调用设备代码。CUDA C的优势在于，它提供了C在语言级别上的集成，因此这个设备函数看上去十分像主机函数调用。CUDA编译器和运行时将负责实现从主机代码中调用设备代码。

所以这个看上去十分奇怪的函数调用实际上表示调用设备代码，使用尖括号表示要将一些参数传递给运行时系统。这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备代码。传递给设备代码本身的参数是放在圆括号中传递，就像标准的函数调用一样。

3.2.3 传递参数

参数可以传递给核函数，继续对程序进行修改：

 

增加了多行代码，包含两个概念：

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

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

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

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

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

对设备指针的使用限制总结如下：

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

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

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

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

得出以下结论：

不能使用标准C的free()函数来释放cudaMalloc()分配的内存，要释放，需要调用cudaFree()，这个函数的行为和free()的行为非常相似。

我们已经学习了如何在设备上分配内存和释放内存，在主机上不能对这块内存做任何修改，访问设备内存的两种最常见的方法是：在设备代码中使用设备指针以及调用cudaMemcpy()。

设备指针的使用方式与标准C中指针的使用方法完全一样。

主机指针上有类似的限制，虽然可以将主机指针传递给设备代码，但是如果想通过主机指针来访问设备代码中的内存，就会出现同样的问题。总的来说，主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。

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

cudaMemcpyHostToDevice将告诉运行时相反的含义，即源指针位于主机上，而目标指针是位于设备上。此外还可以通过参数传递cudaMemcpyDeviceToDevice来告诉运行时这两个指针都是位于设备上，如果两个指针都在主机上，则可以直接调用标准C的memcpy()。

3.3 查询设备

由于我们希望在设备上分配内存和执行代码，因此如果在程序中能够知道设备拥有多少内存以及具备哪些功能，将十分有用。而且，在一台计算机上拥有多个支持CUDA的设备也是很常见的情形。在这些情况中，希望通过某种方式来确定使用的是哪一个处理器。

在深入研究如何编写设备代码之前，需要通过某种机制来判断计算机中当前有哪些设备，以及每个设备都支持哪些功能。可以通过一个很简单的借口来获得这种信息。当我们希望知道系统中有多少个设备是支持CUDA架构的，并且这些设备能够运行于基于CUDA C编写的核函数。要获得CUDA设备的数量，可以调用cudaGetDeviceCount()，这个函数的作用从它的名字就可以看出来。

int count；

HANDLE_ERROR(cudaGetDeviceCount(&count));

在调用cudaGetDeviceCount()后，可以对每个设备进行迭代，并查询每个设备的相关信息。CUDA运行时将返回一个cudaDeviceProp类型的结构，其中包含了设备的相关属性。