4、第一个CUDA程序:编写Hello World核函数、编译与运行、nvcc编译器使用、Makefile编写

好,咱们正式开始动手写代码了。

前面聊了那么多GPU的架构、线程模型,说实话,都是纸上谈兵。真正要理解CUDA,你得亲手写一个能跑在GPU上的程序。我当年刚接触CUDA时,也是从“Hello World”开始的。别小看这个例子,它背后藏着CUDA编程最核心的几个概念:核函数、线程层次、内存模型。

今天我们就一步步来,把第一个CUDA程序跑起来。

4.1 什么是核函数?

核函数,说白了就是跑在GPU上的函数。它跟普通的C函数最大的区别在于:

  • 前面要加一个 __global__ 修饰符
  • 调用时要用特殊的语法 <<<grid, block>>>
  • 它不能有返回值,必须是 void

嗯,这里要注意:__global__ 是告诉编译器,这个函数要被编译成GPU设备代码,而不是CPU主机代码。调用它的地方在CPU上,但执行却在GPU上。

核心概念:核函数 = 在GPU上并行执行的函数。每个线程执行相同的代码,但处理不同的数据。

4.2 编写Hello World核函数

先看代码,我习惯把完整的例子贴出来,然后一行行拆解。

// hello.cu
#include <stdio.h>

// 定义核函数
__global__ void helloFromGPU()
{
    printf("Hello World from GPU! threadId: %d\n", threadIdx.x);
}

int main()
{
    // 调用核函数
    helloFromGPU<<<1, 10>>>();

    // 同步等待GPU执行完成
    cudaDeviceSynchronize();

    printf("Hello World from CPU!\n");
    return 0;
}

这段代码很简单,但有几个关键点我得强调一下:

  • __global__ void helloFromGPU() — 这就是核函数。注意它没有返回值。
  • <<<1, 10>>> — 启动配置。1表示1个线程块,10表示每个块有10个线程。所以总共10个线程在跑。
  • threadIdx.x — 每个线程都有一个内置变量,告诉你“我是谁”。这里x从0到9。
  • cudaDeviceSynchronize() — 这个很重要!CPU调用核函数是异步的,它不会等GPU跑完就继续往下执行。加上这个,CPU会阻塞直到GPU完成。

个人经验:我刚开始写CUDA时,经常忘记加 cudaDeviceSynchronize(),结果CPU先打印了“Hello World from CPU”,GPU的打印才慢悠悠出来。输出顺序乱得一塌糊涂。所以,记得同步!

4.3 编译与运行:nvcc编译器

CUDA的编译器叫 nvcc,它是NVIDIA提供的。你想想看,它要处理两种代码:CPU代码(用gcc/g++编译)和GPU代码(用NVCC自己的编译器编译)。所以nvcc其实是个“编译器驱动”,它会自动分离代码,调用不同的后端。

编译上面的程序很简单:

nvcc -o hello hello.cu

然后运行:

./hello

输出应该是:

Hello World from GPU! threadId: 0
Hello World from GPU! threadId: 1
...
Hello World from GPU! threadId: 9
Hello World from CPU!

注意,GPU线程的执行顺序是不确定的。你每次运行,打印的顺序可能都不一样。这是因为GPU的线程是并行调度的,谁先跑完谁先打印。

避坑指南:我曾经在服务器上编译时,忘了加 -arch=sm_70 之类的架构参数,结果报错说“不支持的GPU架构”。nvcc默认会生成一个兼容性很差的代码,建议你根据自己显卡的算力指定架构。比如:nvcc -arch=sm_75 -o hello hello.cu

4.4 nvcc常用编译选项

这里我整理了一个表格,方便你查阅:

选项 说明 示例
-o 指定输出文件名 -o myprogram
-arch 指定GPU架构(算力) -arch=sm_75
-G 生成调试信息(用于cuda-gdb) -G -o debug_prog
-lineinfo 生成行号信息(性能分析用) -lineinfo
-Xcompiler 传递给主机编译器的选项 -Xcompiler -Wall

我个人习惯在开发阶段用 -G -lineinfo,方便调试。发布时去掉这些,用 -O2 优化。

4.5 编写Makefile

手动敲 nvcc 命令一次两次还行,项目大了就受不了。所以Makefile是必备的。我写一个通用的模板给你:

# Makefile for CUDA Hello World

# 编译器
NVCC = nvcc

# 编译选项
NVCC_FLAGS = -arch=sm_75 -O2

# 目标文件
TARGET = hello

# 源文件
SRC = hello.cu

# 默认目标
all: $(TARGET)

$(TARGET): $(SRC)
	$(NVCC) $(NVCC_FLAGS) -o $@ $<

# 清理
clean:
	rm -f $(TARGET)

# 运行
run: $(TARGET)
	./$(TARGET)

.PHONY: all clean run

使用方法:

  • make — 编译
  • make run — 编译并运行
  • make clean — 清理

小技巧:如果你不确定自己的显卡算力,可以用 nvidia-smi 查看GPU型号,然后去NVIDIA官网查对应的算力。或者直接用 deviceQuery 示例程序,它会打印所有信息。

4.6 本章知识体系

为了让你更直观地理解本章的核心逻辑,我画了一张图:

第一个CUDA程序:知识体系 主机端 (CPU) 设备端 (GPU) 1. 编写 main() 函数 2. 调用核函数 helloFromGPU<<<1,10>>>(); 3. 同步等待 cudaDeviceSynchronize(); 4. 编译:nvcc -o hello hello.cu 1. 定义核函数 __global__ void helloFromGPU() 2. 每个线程执行相同代码 3. 使用 threadIdx.x 区分线程 4. 10个线程并行打印 调用 核函数是连接CPU和GPU的桥梁

这张图很清楚地展示了:CPU负责调用和同步,GPU负责并行执行。核函数就是中间的桥梁。

4.7 总结

好了,第一个CUDA程序就这么跑起来了。说实话,它虽然简单,但意义重大。你亲手验证了:

  • 核函数怎么写
  • 怎么启动多个线程
  • 怎么编译和运行
  • 怎么用Makefile管理

这些基础打牢了,后面学复杂的东西才不会慌。我当年就是靠这个Hello World,一步步走进了GPU并行计算的大门。你也一样。

核心要点回顾

  • 核函数用 __global__ 修饰,无返回值
  • 启动配置 <<<grid, block>>> 决定线程数量
  • 别忘了 cudaDeviceSynchronize()
  • nvcc 编译,用Makefile管理项目

公众号:蓝海资料掘金营,微信deep3321