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 本章知识体系
为了让你更直观地理解本章的核心逻辑,我画了一张图:
这张图很清楚地展示了:CPU负责调用和同步,GPU负责并行执行。核函数就是中间的桥梁。
4.7 总结
好了,第一个CUDA程序就这么跑起来了。说实话,它虽然简单,但意义重大。你亲手验证了:
- 核函数怎么写
- 怎么启动多个线程
- 怎么编译和运行
- 怎么用Makefile管理
这些基础打牢了,后面学复杂的东西才不会慌。我当年就是靠这个Hello World,一步步走进了GPU并行计算的大门。你也一样。
核心要点回顾:
- 核函数用
__global__修饰,无返回值 - 启动配置
<<<grid, block>>>决定线程数量 - 别忘了
cudaDeviceSynchronize() - 用
nvcc编译,用Makefile管理项目
公众号:蓝海资料掘金营,微信deep3321