【AI】来Get下CUDA

最近在GPU上搞模型部署，绕不开的就是CUDA，索性花点时间把它摸一摸。

看效果

先看个效果——两个向量相加，直接用GPU算出来。没错，就是下面这段代码干的事，结果是并行计算出来的。

话原理

概念

CUDA全称Compute Unified Device Architecture，是NVIDIA搞出来的并行计算框架，让GPU不光能渲染画面，还能干通用计算的活。对比CPU那点可怜的计算单元，GPU的计算单元多到让人眼红，天生适合批量并行任务。

NVIDIA GPU上做并行计算，核心流程就两步：

1. CPU调用一个叫核函数的函数，这个函数由GPU来执行。
2. GPU根据给定的并行量，把这些函数实例撒出去并行跑。

CUDA里，执行核函数的基本单位叫线程（thread），若干个线程凑成一个线程块（block），一次调用中所有线程块组成一个网格（grid）。

CPU调用核函数的时候，得指定线程块数量和每个块里的线程数。换句话说，核函数里的内容会被并行执行：线程块数 × 每块线程数 这么多次。

核函数

• __global__ 是CUDA C/C++里的函数修饰符，告诉编译器这是个核函数
• 核函数在GPU上跑，但由主机（CPU）代码调用
• 返回类型必须写成 void
• 调用时函数名后面跟 <<>>：b是线程块数量，t是每块线程数

__global__ void myKernel() {
    printf("Hello world
");
}
int main(int argc, char const *argv[]) {
    myKernel<<<4,2>>>();
    return 0;
}

这个例子指定了4个线程块，每块2个线程，一共8个线程并行执行，所以会输出8个"Hello world"。

除了核函数本身，还得管好内存的分配和回收，以及主机（CPU）和设备（GPU）之间的数据拷贝、错误处理之类的活儿。下面这个向量相加的例子，基本把整套流程都串起来了：

1. 定义核函数 vectorAdd，用 __global__ void 修饰
2. 写主函数，里面依次做：定义参数 → 分配GPU内存 → 初始化数据 → 数据从主机拷到设备 → 调用核函数 → 结果从设备拷回主机 → 打印 → 释放GPU内存

来实践

纸上谈兵没意思，直接上手操作一波。

1. 搞一台带GPU的机器
   。申请云GPU也行，本地有更省事。这里用的是RTX 4090，按小时租大概¥2.7，给了120G内存、16核CPU。

2. 装CUDA Toolkit
   。用云镜像直接装了CUDA 11，省去自己折腾环境。

3. 配环境变量
   。编辑 /root/.bashrc，把CUDA路径加进去，然后 source 一下。

4. 写CUDA代码
   。新建文件 cuda_program.cu，敲入向量相加的程序（就是上面原理部分贴的那段）。

5. 用nvcc编译
   。执行 nvcc -o cuda_program cuda_program.cu，生成可执行文件。

6. 运行
   。直接 ./cuda_program，看到打印结果就算成了。

写在最后

这周还顺便琢磨了几点感慨。模型部署比想象中麻烦，云GPU镜像倒是装好了，结果模型从Hugging Face上下载慢得离谱，已经挂了两天还没完……不过Keep moving就对了。

另外记录几个学到的英文单词（看书时碰到的）：beads（珠子）、Paleolithic period（旧石器时代）、slabs（厚板）、bare walls（裸露的墙壁）、depicted in ca ve art（在洞xue艺术中被描绘）、herds of game（成群的猎物）、mammoths（长毛象）、reindeer（驯鹿）等等。学无止境，不管是CUDA还是词汇。

参考资料

网上有不少好的CUDA入门教程，动手实践一遍比看十遍理论都管用。本文的代码示例和思路就是参考了这些材料整理出来的。