CUDA编程入门历险记

说来惭愧，做CUDA相关工作已经有两年多，但是对CUDA编程的研究并不深入，最近想夯实一下基础，于是找了一个教程“cuda crash course“。

这个教程在B站和youtube上都有：

B站：https://www.bilibili.com/video/BV127411G76m?p=1&vd_source=5d09aad9eacf6d90c0a17cf81ee41eef

youtube：https://www.youtube.com/watch?v=2NgpYFdsduY&list=PLxNPSjHT5qvtYRVdNN1yDcdSl39uHV_sU

教学视频对应的源代码都在这里：https://github.com/CoffeeBeforeArch/cuda_programming/tree/master

这算是一个不错的CUDA入门教程，我看完了前12集，包括sum reduction的前五集，准备实战一下。虽说平时也有CUDA编程经验，但都是在一个成熟的项目中修修补补的工作，像这样start from scratch还是头一回。于是就遇到了诸多困难，在此记录，以便总结提高。

1. docker run

首先尝试使用了一台比较旧的T4服务器，cuda driver版本440。想用一个比较纯净的cuda环境，于是就选择了cuda11.8的devel版cuda镜像（而没有使用triton镜像、pytorch镜像等等），拉取镜像后遇到的第一个问题是镜像启动之后自动exit，手动docker start之后还是会自动exit，网上搜索到一个命令可以查看docker启动失败的原因：

docker logs [OPTIONS] CONTAINER

docker logs具体手册：https://docs.docker.com/reference/cli/docker/container/logs/

于是执行命令后出现了这样一条日志：

/bin/bash: /bin/bash: cannot execute binary file

看起来是用命令/bin/bash执行了/bin/bash这个二进制文件，而理论上/bin/bash后面跟的应该是shell脚本（*.sh）。

在stackoverflow上找到一个类似的问题：https://stackoverflow.com/questions/61055324/docker-cannot-execute-binary-file

于是把docker run命令最后的/bin/bash删掉了，我不知道为什么在另一台A800机器上没有遇到这个问题，也没有深究。

插曲1:

另外关于cuda环境还有两点和本地环境相关的问题，一个是docker run启动脚本里挂进去了宿主机的/usr/local/cuda目录，这样会导致覆盖掉docker容器里面的cuda版本，这是个历史问题，删掉这个挂载就好了。另外一个是需要export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/lib64，这个也是由于我本地的环境没有使用nvidia-docker而使用docker导致的。

插曲2:

nvidia提供的cuda镜像里不含cuda sample，git clone最新的cuda sample工程（https://github.com/NVIDIA/cuda-samples）又暂时没法装git，按照git的时候apt源还有问题，详见下文。于是就用了宿主机上/usr/local/cuda/samples里面的sample，但是编译这里面的sample时遇到报错compute_30架构不支持，看Makefile发现里面写的编译的架构比较旧，从30到最新的，于是删掉75以下的架构支持就没问题了。

插曲3:

T4机器的驱动版本太低，不能运行cuda11.8，换了机器。

做好这些之后nvcc --version可以看到正确的cuda版本，可以编译运行cuda sample。

总结：遇到问题不要慌，查清报错再搜索。

2. apt源

没想到这个问题竟然折磨我许久，因为cuda镜像里面没有git vim这些基本的工具，所以急需apt install，但是apt Install就出现连不上网的情况，即使加了代理也是这样，看了一下报错的网址都是ubuntu的原生地址，想起来国内访问这些apt源好像会有问题，于是查怎么解决。

首先想到是配置厂内源，但是这里出了一个差错就是我以为我用的是ubuntu22.04而其实是ubuntu20.04（debug的时候一定不能想当然！！！），厂内没有22.04的源，于是我开始尝试使用厂外普遍使用的阿里源，还是不行，于是我又仔细看了一下我的ubuntu版本，最终成功使用了厂内的ubuntu22.04源。但是还是会连一个cuda的网站连不上，于是我就查看/etc/apt/source.list路径附近的文件夹，发现有个sourced.list的文件夹，里面有个cuda.txt文件，写了一些nvidia的网址，大概这就是链接失败的原因，我就把这个文件备份后删除了，没有深究，厂内源终于可以成功使用，速度飞起。

关于what is the function of /etc/apt/sources.list.d?：https://stackoverflow.com/questions/26020917/what-is-the-function-of-etc-apt-sources-list-d

3. nvcc编译命令的参数

然后开始编译第一个vector add的例子，我使用的命令就是凭直觉的nvcc vector_add.cu（事实证明千万不能靠直觉！！！又想当然了！！！），而且后来我在网上搜到了几乎所有入门教程里都写的是这样编译。这样编译可以成功，但是运行的时候都会在代码最后cpu与gpu计算结果比较时assert报错，并且core dump，一开始还想用gdb看core dump的原因，后来才知道assert出core是正常的，这个地方是一个计算机基础的缺失（https://stackoverflow.com/questions/21275643/core-dumped-in-assert），这个有点不应该，但是计算机的know-how实在是太多。

表面粗浅理解这个assert是cuda kernel计算错误导致的。（事实证明这个理解却是太粗浅了！！！又想当然了！！！）

于是想要debug这个问题，首先想到就是打印，开始发现CPU代码打印不起作用，想到python编程中类似的问题使用print("Hello", flush=True)解决，搜索C++类似的方法使用fflush(stdout)解决。然后发现kernel中的printf不起作用。查了很多关于cuda kernel printf不起作用的原因，加了cudaDeviceSynchronize()，加了cudaDeviceReset()都不起作用，一筹莫展，其实是走了弯路，这里不单单是printf不起作用的问题，整个cuda kernel的运行行为都不正确（因为逻辑上教程里的代码是正确的，教程一般也不会瞎说，但是本地运行会assert，这就是个很大的疑点！！！但是我要查清这个疑点就要打印，然而又打印不了，就把我困在其中了！！！），但是因为无法打印我也无从debug。

最终还是聚焦问题，先解决打印问题，重新写了一个cuda Hello world的最小kernel，发现也无法打印。在cuda samples里面找到了一个simplePrintf测试用例，欣喜，然后发现是可以打印的。缩小了问题的范围，运行环境是没有问题的。那么就是两个可能性，代码的问题，编译命令的问题。于是进一步测试，使用我的编译命令编译simplePrintf测试用例而不使用它的Makefile（删掉了一些不必要的依赖头文件），发现也是无法打印的，这样就确定了是编译命令的问题。

但是！这个小小的simplePrintf测试用例的Makefile竟然有332行！！！我到底是缺少了什么编译参数导致我无法编译这个问题变得无从下手，其实这里有个目标导向的策略是我从此就开始使用这个测试用例的Makefile即可，但是本次也是想自底向上的搞懂一些cuda的基础概念，于是没有这样做。

Debug到这时候，搞定cuda的初心和信心已经消失殆尽，情绪开始不稳定，开始怀疑人生怀疑研发职业生涯，在老婆的鼓励下，我开始使用bing的chatGPT（copilot）帮我工作！！！

其实之前也问了它几个问题，但是效果不佳，于是改为自己搜索，但其实对于这种入门的问题chatGPT还是很有用的，我修改了问问题的方式（Prompt工程啊！！！让大模型为我所用的关键啊！！！）

就在刚刚！由于我昨晚把那个bing页面关了而且bing也没有登陆，我想不起那个prompt了！！！我问了之后给我的答案还是错的！我也忘了昨晚问的什么了！！！刚才差点崩溃地笑出声，但是！我又重新设计了一个prompt就可以了！所以prompt还是要简洁命令关注关键问题！

我有一段cuda代码，在kernel中打印hello world，应该使用什么样的编译命令

你好！要编译一个CUDA程序，你需要使用NVIDIA的CUDA编译器 nvcc。下面是一个简单的CUDA版“Hello World”小程序，它在CPU和GPU上输出“Hello World”。你可以使用以下命令编译这个程序：

nvcc -arch sm_60 hello_world_gpu.cu -o hello

以下是完整的CUDA代码示例：

#include <stdio.h>

__global__ void helloFromGPU() {
    printf("Hello World from GPU!\\n");
}

int main(int argc, char** argv) {
    printf("Hello World from CPU!\\n");
    helloFromGPU<<<1, 2>>>();
    cudaDeviceReset();
    return 0;
}

这段代码有以下几个关键知识点：

1. __global__：这是一个修饰符，用于标识GPU内核函数。内核函数与C++函数类似，但必须使用此修饰符，并且返回类型必须为void。
2. <<<1, 2>>>：在调用内核函数时，需要指定在设备中分配多少个线程。这里的数字表示线程块数和每个线程块中的线程数。
3. cudaDeviceReset()：显式释放和清空与当前设备相关的所有资源。

请注意，这只是一个简单的示例，实际的CUDA程序可能涉及更复杂的操作。¹²

一个错误的prompt：

请帮我写一段cuda代码，在核函数中打印Hello world，并给出编译命令

当然，我可以帮你写一段简单的 CUDA 代码来在核函数中打印 “Hello, world!”。以下是一个示例：

#include <iostream>

// CUDA 核函数
__global__ void helloWorld() {
    printf("Hello, world from thread %d!\n", threadIdx.x);
}

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

    return 0;
}

编译命令如下（假设你的文件名为 hello.cu）：

nvcc -o hello hello.cu

这将生成一个可执行文件 hello，你可以运行它来查看输出。请注意，你需要安装 NVIDIA CUDA 工具包并配置好编译环境才能成功编译和运行此代码。

希望这对你有帮助！如果你还有其他问题，请随时告诉我。

所以区别就是在于要用-arch=sm_80指定架构！！！

我根据C++的编程经验，应该有类似于readelf或者objdump可以从编译后的二进制文件中看出如果不设置-arch参数默认指定的是什么架构的（程序员的自我修养这本书没白看），查到nvidia提供了一个类似工具cuobjdump，查看到默认使用的是sm_52架构，所以这就是使用nvcc默认编译出来的二进制文件无法正常计算也无法打印的原因！！！！！！