在本节中,我们将通过使用CUDA C编写一个非常基础的程序来学习CUDA编程。我们将从编写一个“Hello,CUDA!”开始,在CUDA C中编程并执行它。在详细介绍代码之前,有一件事你应该记得,主机代码是由标准C编译器编译的,设备代码是由NVIDIA GPU编译器来执行。NVIDIA工具将主机代码提供给标准的C编译器,例如 Windows的Visual Studio和 Ubuntu 的GCC编译器,并使用macOS执行。同样需要注意的是,GPU编译器可以在没有任何设备代码的情况下运行CUDA代码。所有CUDA代码必须保存为*.cu扩展名。
下面就是Hello,CUDA!的代码
#include <iostream>
#include <stdio.h>
__global__ void myfirstkernel(void)
{
}
int main()
{
myfirstkernel << <1, 1 >> >();
printf("Hello, CUDA!\n");
return 0;
}
如果你仔细查看代码,它看起来将非常类似于简单地用C语言编写的Hello,CUDA !用于CPU执行的程序。这段代码的功能也类似。它只在终端或命令行上打印“Hello,CUDA!”。因此,你应该想到两个问题:这段代码有何不同?CUDA C在这段代码中扮演何种角色?这些问题的答案可以通过仔细查看代码来给出。它与用简单的C编写的代码相比,有两个主要区别:
- 一个名为myfirstkernel 的空函数,前缀为_global_
- 使用<<1,1>>>调用myfirstkernel函数
_global_是CUDA C在标准C中添加的一个限定符,它告诉编译器在这个限定符后面的函数定义应该在设备上而不是在主机上运行。在前面的代码中,myfirstkernel将运行在设备上而不是主机上,但是,在这段代码中,它是空的。
那么,main 函数将在哪里运行? NVCC编译器将把这个函数提供给C编译器,因为它没有被global关键字修饰,因此main函数将在主机上运行。
代码中的第二个不同之处在于对空的myfirstkernel函数的调用带有一些尖括号和数值。这是一个CUDA C技巧:从主机代码调用设备代码。它被称为内核调用。内核调用的细节将在后面的章节中解释。尖括号内的值表示我们希望在运行时从主机传递给设备的参数。基本上,它表示块的数量和将在设备上并行运行的线程数。因此,在这段代码中,<<1,1> >>表示myfirstkernel将运行在设备上的一个块和一个线程或块上。虽然这不是对设备资源的最佳使用,但是理解在主机上执行的代码和在设备上执行的代码之间的区别是一个很好的起点。