GPU（四）CUDA hello world

nuczzz

已于 2024-03-31 22:16:22 修改

阅读量916

点赞数 19

分类专栏： GPU 文章标签： GPU AI

于 2024-03-09 20:10:06 首次发布

本文链接：https://blog.csdn.net/NUCEMLS/article/details/136589916

版权

GPU 专栏收录该内容

7 篇文章 4 订阅

订阅专栏

通过前面的文章，我们对GPU、CUDA有了初步的了解，从本文开始，我们来学习c语言语法的CUDA编程。

1、c语言的hello world

如果学过C语言得到读者，可能还记的初学编程时的第一个hell world程序：

// main.c
#include<stdio.h>

int main()
{
        printf("hello world\n");
}

用gcc编译器编译，并执行：

// 编译main.c为可执行程序c_hello
$ gcc main.c -o c_hello

// 查看当前目录下的文件
$ ll
total 24
drwxr-xr-x 2 nuczzz nuczzz 4096 Mar  9 10:39 ./
drwxr-xr-x 3 nuczzz nuczzz 4096 Mar  9 10:19 ../
-rwxr-xr-x 1 nuczzz nuczzz 8296 Mar  9 10:39 c_hello*
-rw-r--r-- 1 nuczzz nuczzz   60 Mar  9 10:39 main.c

// 执行可执行程序c_hello
$ ./c_hello
hello world

2、nvcc编译纯c代码

前面我们安装CUDA环境的时候一起安装了CUDA程序的编译器nvcc，nvcc编译器支持纯c/c++代码编译，我们看看使用nvcc来编译上述c代码。

在文件名后缀上和c/c++文件类似，nvcc对应编译的文件后缀也有自己的约定，用.cu。我们先复制一下上述c文件：

$ cp main.c main.cu
$ ll main.*
-rw-r--r-- 1 nuczzz nuczzz 60 Mar  9 10:39 main.c
-rw-r--r-- 1 nuczzz nuczzz 60 Mar  9 10:57 main.cu

注意，此时仅仅是复制了一下c代码文件，并没有对文件内容做任何修改。

使用nvcc编译main.cu并执行：

$ nvcc main.cu -o nvcc_hello
$ ./nvcc_hello
hello world

可以看到，编译正常，输出结果也和gcc编译输出结果一致，这就是nvcc对c语言语法（其实c++也一样）的友好性。

3、CUDA的hello world

我们再准备一个文件名为cuda.cu、文件内容如下的程序：

#include<stdio.h>

__global__ void hello_world()
{
        printf("hello world\n");
}

int main()
{
        hello_world<<<2, 2>>>();
        cudaDeviceSynchronize();
        return 0;
}

使用nvcc编译并执行：

$ nvcc cuda.cu -o cuda_hello
$ ll cuda*
-rw-r--r-- 1 nuczzz nuczzz    159 Mar  9 11:10 cuda.cu
-rwxr-xr-x 1 nuczzz nuczzz 989224 Mar  9 11:15 cuda_hello*
$ ./cuda_hello
hello world
hello world
hello world
hello world

可以看到编译正常并且正常输出，恭喜你，你已经成功的完成了c语法的CUDA hello world编程！

4、从CUDA hello world开始

很多小伙伴可能还是一脸问号：我怎么就完成CUDA hello world了？上面的cuda.cu代码和纯c语法有些不同，我都看不懂。别着急，我们先了解一些背景资料。

前面介绍过， GPU是无法单独完成计算功能的，它需要与CPU组成异构架构才能正常工作 ，我们可以把GPU当做CPU的协处理器，也就是把GPU当做CPU的外设，因此GPU如何执行、需要处理哪些数据，都需要经过CPU的任务下达与指令控制。

在这里插入图片描述

所以 在编写CUDA程序时，既要编写在CPU（host）上运行的主机代码，也要编写在GPU（device）上运行的设备代码，主机对设备的调用是通过核函数进行的。

现在我们再来详细讲解上述cuda.cu程序。为了方便讲述，我们先给cuda.cu程序加上行号：

  1 #include<stdio.h>
  2
  3 __global__ void hello_world()
  4 {
  5         printf("hello world\n");
  6 }
  7
  8 int main()
  9 {
 10         hello_world<<<2, 2>>>();
 11         cudaDeviceSynchronize();
 12         return 0;
 13 }

1、第1行：引入头文件stdio.h，这和c语言是一样的

2、第3~6行：如果把__global__去掉，其实就是一个普通的c语言函数，这个函数使用printf输出hello world，没有入参，返回值是void。

__global__是一种CUDA编程特定修饰前缀，除了__global__，CUDA有2中特定的修饰前缀，分别是__host__和__device__，这三种修饰前缀的含义分别为：

__host__：与c/c++中函数相同，由CPU调用CPU执行的函数
__global__：表示一个核函数，是一组由GPU执行的 并行 计算任务，具体并行度在函数调用处指定。核函数必须是CPU调用。
__device__：表示一个由GPU中一个线程调用的函数。由于Tesla架构的GPU允许线程调用函数，因此实际上是将__device__ 函数以__inline形式展开后直接编译到二进制代码中实现的，并不是真正的函数。

核函数除了需要加__global__前缀修饰，还有一些注意事项：

核函数要求函数的返回值必须是void
核函数只能访问GPU内存，也就是显存
核函数不能使用变长参数
核函数不能使用静态变量
核函数不能使用函数指针
核函数具有异步性
核函数不支持c++的iostream

3、第8行，c语言的main函数入口

4、第10行，核函数调用。与普通c/c++函数调用不同，该核函数调用时多了三对尖括号，并且尖括号里有两个逗号分隔的数字。尖括号中的数字是用来指定核函数执行的线程数，其中第一个数字指的是现线程块的个数，第二个数值指的是每个线程块中线程的数量。

很显然<<<2, 2>>>的写法就是表示指定2个线程块、每个线程块2个线程来执行这个核函数，因此总共有2*2=4个线程会执行这个核函数，所以输出内容有4行hello world。

<<<2, 2>>>是一维指定线程块、线程数量的写法，还有二维、三维的写法，本文不展开讨论，相关内容会在后续的GPU线程模型中再做详细说明。

5、第11行，同步host与device，促使缓冲区刷新，从而在终端打印hello world。

因核函数具有异步性，CPU并不会等待GPU核函数执行完毕再继续执行后续代码，如果不调用cudaDeviceSynchronize函数，CPU逻辑都执行完程序退出了，就无法打印hello world。

6、第12行，main函数返回0，表示执行成功。

总结

本文通过c语言的hello world到CUDA的hello world，初步了解了如何编写CUDA程序。其实编写CUDA程序访问GPU有两种API：Driver API和Runtime API，本文使用的是对用户更加友好的Runtime API方式。也有一些基于CUDA Runtime API封装的SDK，这是后话了。

在这里插入图片描述