2022CUDA夏季训练营Day1实践

 

张小白曾经成功地运行了第一个CUDA程序：

https://bbs.csdn.net/topics/607184868

但是对其只是知其然不知其所以然。所以CUDA训练营是帮你知其所以然的。

我们把 CPU,内存这块区域叫做“主机（HOST）”，把GPU，显存这块区域叫做“设备（DEVICE）”。

CUDA的代码执行包含以下几步：

​

​

​

简述一下，就是 host_to_device-》在device上并行计算-》device_to_host。

cuda程序其实是一个对C的扩展程序。其后缀名为.cu，如果头文件则为.cuh。

这个.cu 程序除了C程序的语法外，还有一些cuda的特有部分，比如它在函数前面加了前缀，分为 __global__, __host__,__device__ 三种。

对于__global__，训练营是这么说的：

​

所谓的“执行配置”，下面会提到，比如说是 <<< >>>中间的内容。

这个标识符将一个C函数声明成一个 核函数。它只能在设备（device）上执行。

对于__host__，训练营是这么说的：

​

现在好像还不太懂。。等下次看示例吧。

对于__device__，训练营是这么说的：

​

个人理解，这几个前缀定义了这些代码运行的设备，这会让程序决定在哪个设备上运行。

对于一个简单的Hello World代码而言：

#include <stdio.h>

void hello_from_cpu()
{
    printf("Hello World from the CPU!\n");
}

int main(void)
{
    hello_from_cpu();
    return 0;
}

如果我们想让它在GPU上运行，仅需要做两步：

（1）将被调用的函数 hello_from_cpu 改为 hello_from_gpu ，前面加上 __global__ 将其定义为核函数。

（2）在main主函数调用的时候，加上执行配置<<< >>>部分，如加上<<<1,1>>>则为并行1次，如加上<<<2,4>>>则运行 2X4次。

我们看看实际代码修改后的效果：

#include <stdio.h>

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

int main(void)
{
    hello_from_gpu<<<1,1>>>();
    return 0;
}

cu代码必须使用nvcc编译，编译的时候要根据GPU架构的不同填不同的参数。

​

其中，arch参数如下：

​

code参数如下：

举个简单的例子，张小白这台笔记本的显卡是Quardo P1000，是Pascal架构，那么参数是compute_61和sm_61.

我们执行以下语句：

/usr/local/cuda/bin/nvcc -arch=compute_61 -code=sm_61 hello_cuda.cu -o hello_cuda

./hello_cuda

​

如果将执行配置改为2,4：

​

可以发现这个核函数被执行了8次。

下面我们来看看这段代码能不能使用jupyter lab执行。

先看下IP地址：

要想执行ifconfig得先 sudo apt install net-tools​

我们解压下 训练营提供的 jupyter练习包：CUDA_on_ARM.zip

解压：unzip，要想执行先得装unzip:

​

执行解压：unzip xx.zip

​

安装PIP:

sudo apt install python3-pip

​

​

​

看来python3和pip都准备OK了。

pip install jupyterlab -i https://pypi.tuna.tsinghua.edu.cn/simple​

​sudo apt install jupyter

​

​启动 jupyter lab --no-browser看看：

​

报错了，没关系，度娘一下：https://blog.csdn.net/weixin_45438997/article/details/124261720

pip show markupsafe

​python -m pip install markupsafe==2.0.1

​重新启动：jupyter lab --no-browser

​

当然启动前可以先设一下密码：

jupyter server password

输入两次123456

​

浏览器打开：http://127.0.0.1:8888/lab​

​

可见，在jupyter lab中也可以体验CUDA代码的。

番外1：解决jupyter lab中sudo时需要密码的问题

Jupyter Lab是一个很好的东西。它让编程可视化了。

然而，如果你执行这样的指令：

!sudo /usr/local/cuda/bin/nvprof --print-api-trace ./hello_cuda

那就有点糟糕了：

​因为它会问你要密码，但是却不给你输入密码的地方。

于是你只有等着左边的星号一时亮着。。。并不会像其他的cell那样执行后变成一个数字。

然而，并不是没有解决的方法的。你可以试着这样：

!echo nano|sudo -S /usr/local/cuda/bin/nvprof --print-api-trace ./hello_cuda

把密码作为管道后面的输入。

你会发现，居然他能运行了！

​

​

以上是CUDA查看程序性能的程序，目前貌似在Jetson Nano上运行良好。但是张小白曾经尝试在Quardo的P1000 GPU上运行，它会报错：

​

这个问题只有等待后续解决了。

番外2：一个错误导致张小白加深了对CUDA异步执行的理解

张小白在Jupyter Lab中执行cuda的helloworld（对不起，都第三遍helloworld了）发现一个奇怪的现象，有时候，编译执行后会有输出，有时候，编译执行后并没有输出。但是执行nprof的时候一直都有输出。

​

于是张小白就在群里问了何老师和欢老师。

并且贴出了代码：

​

两位老师一眼（对不起，是两眼。。）就看出，我的代码中少了个同步过程。即少了 cudaDeviceSynchronize();

因为CUDA的代码本身是异步执行的。程序指令发出就结束了。这个时候，如果没有 执行cudaDeviceSynchronize()，做同步等待，那么有可能结果在程序执行完毕前会出来，有可能结果会在程序执行完毕后才会出来。那个时候，调用方的程序早就退出了。这也导致了张小白发现的古怪现象——有时候有结果，有时候没结果。

加上同步后代码如下：

#include <stdio.h>

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

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

执行后效果如下：

​

这回，无论是nvcc编译时直接加上-run参数，还是编译后单独运行，或者是在nvprof中运行，都有打印的结果了。

这样也好，加深一下印象。以后代码中也不会出现忘记加同步的错误了吧。。。（没准还会有！）

（未完待续）