【整理贴】关于CUDA底层知识和编程

技术分享
#1

近期选了一门GPU编程，正好对GPU编程方向在学习，以及被迫营业的项目需求（。这里对我自己学习GPU编程的内容进行一些资料博客的整理。

学习CUDA编程神项目： PacktPublishing/Learn-CUDA-Programming: Learn CUDA Programming, published by Packt (github.com)
手把手带着从CUDA版的“Hello World“开始搓！
原书可以从zlibrary 嫖Z-Library – the world’s largest e-book library. Your gateway to knowledge and culture.（大陆没有这个源文件）

【资料整理】

Ⅰ. CUDA/cudnn/CUDA Toolkit/NVCC

参考：
CUDA Runtime API :: CUDA Toolkit Documentation (nvidia.com)

Pytorch 使用不同版本的 cuda - yhjoker - 博客园 (cnblogs.com)

显卡，显卡驱动,nvcc, cuda driver,cudatoolkit,cudnn到底是什么？ - 知乎 (zhihu.com)

安装多版本 cuda ，多版本之间切换_cuda 切换版本-CSDN博客

Linux中PATH、 LIBRARY_PATH、 LD_LIBRARY_PATH的区别_慕课手记 (imooc.com)

Nvidia-smi简介及常用指令及其参数说明_nvidia-smi的作用-CSDN博客

一张图了解GPU、CUDA、CUDA toolkit和pytorch的关系_cuda和pytorch的关系-CSDN博客

简介：

CUDA 为“GPU通用计算”构建的运算平台。

CUDA是显卡厂商NVIDIA推出的运算平台。CUDA™是一种由NVIDIA推出的通用并行计算架构，是一种并行计算平台和编程模型，该架构使GPU能够解决复杂的计算问题。CUDA英文全称是Compute Unified Device Architecture。

有人说：CUDA是一门编程语言，像C,C++,python 一样，也有人说CUDA是API。
官方说：CUDA是一个并行计算平台和编程模型，能够使得使用GPU进行通用计算变得简单和优雅。
运行CUDA应用程序要求系统至少具有一个具有CUDA功能的GPU和与CUDA Toolkit兼容的驱动程序。

cudnn为深度学习计算设计的软件库。

CUDNN是基于CUDA的深度学习GPU加速库，有了它才能在GPU上完成深度学习的计算；
来自知乎的解释：CUDA看作是一个工作台，上面配有很多工具，如锤子、螺丝刀等。cuDNN是基于CUDA的深度学习GPU加速库，有了它才能在GPU上完成深度学习的计算。它就相当于工作的工具，比如它就是个扳手。但是CUDA这个工作台买来的时候，并没有送扳手。想要在CUDA上运行深度神经网络，就要安装cuDNN，就像你想要拧个螺帽就要把扳手买回来。这样才能使GPU进行深度神经网络的工作，工作速度相较CPU快很多。
基本上所有的深度学习框架都支持cuDNN这一加速工具，例如：Caffe、Caffe2、TensorFlow、Torch、Pytorch、Theano等。

CUDA Toolkit (nvidia)： CUDA完整的工具安装包，其中提供了 Nvidia 驱动程序、开发 CUDA 程序相关的开发工具包等可供安装的选项。包括 CUDA 程序的编译器、IDE、调试器等，CUDA 程序所对应的各式库文件以及它们的头文件。

CUDA工具包的主要包含了CUDA-C和CUDA-C++编译器、一些科学库和实用程序库、CUDA和library API的代码示例、和一些CUDA开发工具。（通常在安装CUDA Toolkit的时候会默认安装CUDA Driver；但是我们经常只安装CUDA Driver，没有安装CUDA Toolkit，因为有时不一定用到CUDA Toolkit；比如我们的笔记本电脑，安装个CUDA Driver就可正常看视频、办公和玩游戏了）

CUDA Toolkit (Pytorch)： CUDA不完整的工具安装包，其主要包含在使用 CUDA 相关的功能时所依赖的动态链接库。不会安装驱动程序。

NVCC ：CUDA的编译器，只是 CUDA Toolkit 中的一部分

NVCC就是CUDA的编译器,可以从CUDA Toolkit的/bin目录中获取,类似于gcc就是c语言的编译器。由于程序是要经过编译器编程成可执行的二进制文件，而cuda程序有两种代码，一种是运行在cpu上的host代码，一种是运行在gpu上的device代码，所以nvcc编译器要保证两部分代码能够编译成二进制文件在不同的机器上执行。

image-20240227180055438
image-202402271800554381124×1094 107 KB

1. cuda 与 cudatoolkit 的区别

在使用 Anaconda 安装 Pytorch 深度学习框架时，可以发现 Anaconda 会自动为我们安装 cudatoolkit.

上述安装的 cudatoolkit 与通过 Nvidia 官方提供的 CUDA Toolkit 是不一样的。具体而言，Nvidia 官方提供的 CUDA Toolkit 是一个完整的工具安装包，其中提供了 Nvidia 驱动程序、开发 CUDA 程序相关的开发工具包等可供安装的选项。使用 Nvidia 官网提供的 CUDA Toolkit 可以安装开发 CUDA 程序所需的工具，包括 CUDA 程序的编译器、IDE、调试器等，CUDA 程序所对应的各式库文件以及它们的头文件。上述 CUDA Toolkit 的具体组成可参考 CUDA Toolkit Major Components.

实际上，Nvidia 官方提供安装的 CUDA Toolkit 包含了进行 CUDA 相关程序开发的编译、调试等过程相关的所有组件。但对于 Pytorch 之类的深度学习框架而言，其在大多数需要使用 GPU 的情况中只需要使用 CUDA 的动态链接库支持程序的运行( Pytorch 本身与 CUDA 相关的部分是提前编译好的 )，就像常见的可执行程序一样，不需要重新进行编译过程，只需要其所依赖的动态链接库存在即可正常运行。故而，Anaconda 在安装 Pytorch 等会使用到 CUDA 的框架时，会自动为用户安装 cudatoolkit，其主要包含应用程序在使用 CUDA 相关的功能时所依赖的动态链接库。在安装了 cudatoolkit 后，只要系统上存在与当前的 cudatoolkit 所兼容的 Nvidia 驱动，则已经编译好的 CUDA 相关的程序就可以直接运行，而不需要安装完整的 Nvidia 官方提供的 CUDA Toolkit .

通过 Anaconda 安装的应用程序包位于安装目录下的 /pkg 文件夹中，如笔者的目录即为 /home/xxx/anaconda3/pkgs/ ，用户可以在其中查看 conda 安装的 cudatoolkit 的内容，如下图所示。可以看到 conda 安装的 cudatoolkit 中主要包含的是支持已经编译好的 CUDA 程序运行的相关的动态链接库。( Ubuntu 环境下 )

在大多数情况下，上述 cudatoolkit 是可以满足 Pytorch 等框架的使用需求的。但对于一些特殊需求，如需要为 Pytorch 框架添加 CUDA 相关的拓展时( Custom C++ and CUDA Extensions )，需要对编写的 CUDA 相关的程序进行编译等操作，则需安装完整的 Nvidia 官方提供的 CUDA Toolkit.

Pytorch 确定所使用的 cuda 版本

实际使用过程中，Pytorch 检测运行时使用的 cuda 版本的代码位于 torch/utils/cpp_extension.py 的_find_cuda_home 函数 ( Pytorch 1.1.0, Line 24 )中.这里主要介绍 Linux 环境下的 cuda 版本的确认过程，关于 Windows 环境下多版本 cuda 的使用可以参考上述文件中的具体实现.

确定 cuda 路径

若在运行时需要使用 cuda 进行程序的编译或其他 cuda 相关的操作，Pytorch 会首先定位一个 cuda 安装目录( 来获取所需的特定版本 cuda 提供的可执行程序、库文件和头文件等文件 )。具体而言，Pytorch 首先尝试获取环境变量 CUDA_HOME/CUDA_PATH 的值作为运行时使用的 cuda 目录。若直接设置了 CUDA_HOME/CUDA_PATH 变量，则 Pytorch 使用 CUDA_HOME/CUDA_PATH 指定的路径作为运行时使用的 cuda 版本的目录。

若上述环境变量不存在，则 Pytorch 会检查系统是否存在固定路径 /usr/local/cuda 。默认情况下，系统并不存在对环境变量 CUDA_HOME 设置，故而 Pytorch 运行时默认检查的是 Linux 环境中固定路径 /usr/local/cuda 所指向的 cuda 目录。 /usr/local/cuda 实际上是一个软连接文件，当其存在时一般被设置为指向系统中某一个版本的 cuda 文件夹。使用一个固定路径的软链接的好处在于，当系统中存在多个安装的 cuda 版本时，只需要修改上述软连接实际指向的 cuda 目录，而不需要修改任何其他的路径接口，即可方便的通过唯一的路径使用不同版本的 cuda. 如笔者使用的服务器中，上述固定的 /usr/local/cuda 路径即指向一个较老的 cuda-8.0 版本的目录。

需要注意的是， /usr/local/cuda 并不是一个 Linux 系统上默认存在的路径，其一般在安装 cuda 时创建( 为可选项，不强制创建 )。故而 Pytorch 检测上述路径时也可能会失败。

若 CUDA_HOME 变量指定的路径和默认路径 /usr/local/cuda 均不存在安装好的 cuda 目录，则 Pytorch 通过运行命令 which nvcc 来找到一个包含有 nvcc 命令的 cuda 安装目录，并将其作为运行时使用的 cuda 版本。具体而言，系统会根据环境变量 PATH 中的目录去依次搜索可用的 nvcc 可执行文件，若环境变量 PATH 中包含多个安装好的 cuda 版本的可执行文件目录( 形如/home/test/cuda-10.1/bin )，则排在 PATH 中的第一个 cuda 的可执行文件目录中的 nvcc 命令会被选中，其所对应的路径被选为 Pytorch 使用的 cuda 路径。同样的，若 PATH 中不存在安装好的 cuda 版本的可执行目录，则上述过程会失败，Pytorch 最终会由于找不到可用的 cuda 目录而无法使用 cuda.比较推荐的做法是保持 PATH 路径中存在唯一一个对应所需使用的 cuda 版本的可执行目录的路径。

在确定好使用的 cuda 路径后，基于 cuda 的 Pytorch 拓展即会使用确定好的 cuda 目录中的可执行文件( /bin )、头文件( /include )和库文件( /lib64 )完成所需的编译过程。

Pytorch 使用特定的 cuda 版本

image-20240227180207442
image-202402271802074421114×756 135 KB

从 Pytorch 确定使用的 cuda 版本的流程来看，想要指定 Pytorch 使用的 cuda 版本，主要有两种方法，第一种是修改软链接 /usr/local/cuda 所指向的 cuda 安装目录( 若不存在则新建 )，第二种是通过设置环境变量 CUDA_HOME 指向所需使用的 cuda 版本的安装目录。除此之外，还建议将对应 cuda 目录中的可执行文件目录( 形如/home/test/cuda-10.1/bin )加入环境变量 PATH 中。

对于第一种方法，由于 /usr/ 和 /usr/local/ 目录下的文件均为 root 用户所管理，故而普通用户无法对其进行修改。对于具备了 root 权限的用户而言，在安装有多版本 cuda 的 Linux 系统上，只需切换 /usr/local/cuda 所指向的 cuda 目录，让其指向所需的 cuda 版本的安装位置，即可让 Pytorch 在运行时使用指定版本的 cuda 运行程序。修改软链接的方法如下命令所示，命令删除原有的软链接，并新建指向新路径的软链接。

  sudo rm -rf /usr/local/cuda           //删除软链接，注意是 /usr/local/cuda 而不是 /usr/local/cuda/,前者仅删除软链接，而后者会删除软链接所指向的目录的所有内容，操作请小心
  sudo ln -s cuda_path /usr/local/cuda    //创建名为 /usr/local/cuda 的软链接，其指向 cuda_path 所指定的 cuda 安装目录

或者直接强制修改原始的软链接

    sudo ln -sf cuda_path /usr/local/cuda    //修改或创建软链接 /usr/local/cuda 使其指向指定版本的 cuda 目录

对于非 root 用户而言，主要通过第二种方法进行设置。若想要指定 Pytorch 使用的 cuda 版本，则首先需要设置 CUDA_HOME 环境变量，之后在 PATH 中加入指定 cuda 版本的可执行目录，也就时 cuda_path/bin/ 目录。完成设置后，运行 Pytorch 时所使用的即为对应的 cuda 版本。

实例

以笔者的服务器账户为例，笔者在 /home/test/cuda-10.1 目录中安装了 cuda-10.1 ，而服务器上的 /usr/local/cuda 目录指向的是之前安装的老版本的 cuda-8.0，直接运行 Pytorch 时，其会基于上面的确认流程直接使用老版本的 cuda .若想要临时设置 Pytorch 使用新安装的 cuda ，则可以通过 export 命令修改全局变量。这种设置方式在当前终端退出后即失效。

    export CUDA_HOME=/home/test/cuda-10.1/           //设置全局变量 CUDA_HOME
    export PATH=$PATH:/home/test/cuda-10.1/bin/        //在 PATH 变量中加入需要使用的 cuda 版本的路径,使得系统可以使用 cuda 提供的可执行文件，包括 nvcc

想要永久设置上述 cuda 设置，用户可以直接在自己的 bash 设置文件 ~/.bashrc 文件尾部加入上述命令，保存后再通过 source ~/.bashrc 执行文件，即可完成当前终端的环境变量修改。如果需要使用新的 cuda 来编译文件，还可以通过 LD_LIBRARY_PATH 变量指定进行链接的 cuda 库文件的路径。

位于 ~/.bashrc 文件中的指令在每次终端启动时均会自动运行，后续本用户所打开的终端中的环境变量均会首先执行上述文件中的命令，从而获得对应的 cuda 变量。

获取 Pytorch 使用的 cuda 版本

目前，网络上比较多的资源会讨论如何获得 Pytorch 使用的 cuda 的版本的方法。比较主流的一种方法是使用 Pytorch 提供的方法 torch.version.cuda .

    >>>import torch
    >>>torch.version.cuda    #输出一个 cuda 版本

事实上，上述输出的 cuda 的版本并不一定是 Pytorch 在实际系统上运行时使用的 cuda 版本，而是编译该 Pytorch release 版本时使用的 cuda 版本。

torch.version.cuda 是位于 torch/version.py 中的一个变量， Pytorch 在基于源码进行编译时，通过 tools/setup_helpers/cuda.py 来确定编译 Pytorch 所使用的 cuda 的安装目录和版本号，确定的具体流程与 Pytorch 运行时确定运行时所使用的 cuda 版本的流程较为相似，具体可以见其源码( Pytorch 1.1.0, Line 66 ).在进行 Pytorch 源码编译时，根目录下的 setup.py 会调用上述代码，确定编译 Pytorch 所使用的 cuda 目录和版本号，并使用获得的信息修改 torch/version.py 中的 cuda 信息( Pytorch, Line 286 )。上述 torch.version.cuda 输出的信息即为编译该发行版 Pytorch 时所使用的 cuda 信息。若系统上的 Pytorch 通过 conda 安装，用户也可以直接通过 conda list | grep pytorch 命令查看安装的 Pytorch 的部分信息。

    conda list | grep pytorch    //查看安装的 Pytorch 的信息

笔者环境下上述命令的结果如图所示，可以看到显示的 cuda 信息与 torch.version.cuda 保持一致。

想要查看 Pytorch 实际使用的运行时的 cuda 目录，可以直接输出之前介绍的 cpp_extension.py 中的 CUDA_HOME 变量。

    >>> import torch
    >>> import torch.utils
    >>> import torch.utils.cpp_extension
    >>> torch.utils.cpp_extension.CUDA_HOME        #输出 Pytorch 运行时使用的 cuda

2.nvcc和nvidia-smi显示的CUDA版本不同？

在实验室的服务器上nvcc --version显示的结果如下：

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2018 NVIDIA Corporation
Built on Tue_Jun_12_23:07:04_CDT_2018
Cuda compilation tools, release 9.2, V9.2.148

而nvidia-smi显示结果如下：

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 410.104      Driver Version: 410.104      CUDA Version: 10.0     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla V100-PCIE...  On   | 00000000:01:00.0 Off |                  Off |
| N/A   28C    P0    26W / 250W |      0MiB / 16130MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   1  Tesla P100-PCIE...  On   | 00000000:02:00.0 Off |                  Off |
| N/A   24C    P0    30W / 250W |      0MiB / 16280MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+

+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|  No running processes found                                                 |
+-----------------------------------------------------------------------------+

可以看到nvcc的CUDA 版本是9.2，而nvidia-smi的CUDA版本是10.0。很奇怪的是有时候绝大多数情况代码也能整成跑起来，stackoverflow上的一个解释如下：

CUDA有两个主要的API：runtime(运行时) API和driver API。这两个API都有对应的CUDA版本（如9.2和10.0等）。

  • 用于支持driver API的必要文件(如libcuda.so)是由GPU driver installer安装的。nvidia-smi就属于这一类API。
  • 用于支持runtime API的必要文件(如libcudart.so以及nvcc)是由CUDA Toolkit installer安装的。（CUDA Toolkit Installer有时可能会集成了GPU driver Installer）。nvcc是与CUDA Toolkit一起安装的CUDA compiler-driver tool，它只知道它自身构建时的CUDA runtime版本。它不知道安装了什么版本的GPU driver，甚至不知道是否安装了GPU driver。

综上，如果driver API和runtime API的CUDA版本不一致可能是因为你使用的是单独的GPU driver installer(显卡驱动里的GPU driver installer)，而不是CUDA Toolkit installer里的GPU driver installer。runtime和driver API在很多情况非常相似，也就是说用起来的效果是等价的，但是你不能混合使用这两个API，因为二者是互斥的。也就是说在开发过程中，你只能选择其中一种API。简单理解二者的区别就是：runtime是更高级的封装，开发人员用起来更方便，而driver API更接近底层，速度可能会更快。两者的区别：

  • 复杂性

    • runtime API通过提供隐式初始化、上下文管理和模块管理来简化设备代码管理。这使得代码更简单，但也缺乏驱动程序API所具有的控制级别。
    • 相比之下，driver API提供了更细粒度的控制，特别是在上下文和模块加载方面。实现内核启动要复杂得多，因为执行配置和内核参数必须用显式函数调用指定。

  • 控制

    • 对于runtime API，其在运行时，所有内核都在初始化期间自动加载，并在程序运行期间保持加载状态。
    • 而使用driver API，可以只加载当前需要的模块，甚至动态地重新加载模块。driver API也是语言独立的，因为它只处理cubin对象。

  • 上下文管理

    • 上下文管理可以通过driver API完成，但是在runtime API中不公开。相反，runtime API自己决定为线程使用哪个上下文。
    • 如果一个上下文通过driver API成为调用线程的当前上下文，runtime将使用它。
    • 如果没有这样的上下文，它将使用“主上下文(primary context)”。

image-20240227172346968
image-202402271723469681043×1126 94.8 KB

image-20240227174044447
image-202402271740444471139×681 99.6 KB

3. PATH

PATH是可执行文件路径，是三个中我们最常接触到的，因为我们命令行中的每句能运行的命令，如ls、top、ps等，都是系统通过PATH找到了这个命令执行文件的所在位置，再run这个命令（可执行文件）。 比如说，在用户的目录~/mycode/下有一个bin文件夹，里面放了有可执行的二进制文件、shell脚本等。如果想要在任意目录下都能运行上述bin文件夹的可执行文件，那么只需要把这个bin的路径添加到PATH即可，方法如下：

# vim ~/.bashrc
PATH=$PATH:~/mycode/bin

4. LIBRARY_PATH和LD_LIBRARY_PATH

这两个路径可以放在一起讨论，

  • LIBRARY_PATH是程序编译期间查找动态链接库时指定查找共享库的路径
  • LD_LIBRARY_PATH是程序加载运行期间查找动态链接库时指定除了系统默认路径之外的其他路径

两者的共同点是库，库是这两个路径和PATH路径的区别，PATH是可执行文件。

两者的差异点是使用时间不一样。一个是编译期，对应的是开发阶段，如gcc编译；一个是加载运行期，对应的是程序已交付的使用阶段。

配置方法也是类似：

export  LD_LIBRARY_PATH=LD_LIBRARY_PATH:XXXX

Ⅱ. CUDA Toolkit具体组成

参考：CUDA 12.3 Update 2 Release Notes (nvidia.com)

  • CUDA Samples: 演示如何使用各种CUDA和library API的代码示例。可在Linux和Mac上的samples/目录中获得，Windows上的路径是C：\ProgramData\NVIDIA Corporation\CUDA Samples中。在Linux和Mac上，samples/目录是只读的，如果要对它们进行修改，则必须将这些示例复制到另一个位置。
  • CUDA Driver: 运行CUDA应用程序需要系统至少有一个具有CUDA功能的GPU和与CUDA工具包兼容的驱动程序。每个版本的CUDA工具包都对应一个最低版本的CUDA Driver，也就是说如果你安装的CUDA Driver版本比官方推荐的还低，那么很可能会无法正常运行。CUDA Driver是向后兼容的，这意味着根据CUDA的特定版本编译的应用程序将继续在后续发布的Driver上也能继续工作。通常为了方便，在安装CUDA Toolkit的时候会默认安装CUDA Driver。在开发阶段可以选择默认安装Driver，但是对于像Tesla GPU这样的商用情况时，建议在官方安装最新版本的Driver。 目前(2024年2月)的CUDA Toolkit和CUDA Driver版本的对应情况如下：

image
image1639×1051 75.4 KB

Ⅲ. NVCC简介

  • nvcc其实就是CUDA的编译器，cuda程序有两种代码， 在cpu上的host代码和在gpu上的device代码。
  • .cu后缀：cuda源文件，包括host和device代码
  • nvcc编译例子
nvcc –cuda x.cu –keep

# x.cudafe1.gpu

# x.cudafe2.gpu

# x.cudafe1.cpp
1 个赞
#2

program 1

#include<stdio.h> 
__global__ void cuda_hello() 
{ 
    printf("Hello World from GPU!\n"); 
}
 
int main() 
{ 
    printf("Hello World from CPU!\n");
    cuda_hello<<<2,3>>>();      
    cudaDeviceReset(); 
    return 0; 
}

首先打印一遍Hello World from CPU!

接着调用cuda_hello()函数，开启2个block，3个thread，共6个线程，打印6遍Hello World from GPU!

image-20240313194025892
image-20240313194025892973×282 18.9 KB

program2

#include<stdio.h>
__global__ void hello_from_gpu()
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    printf("Hello World from block %d and thread %d!\n", bid, tid);
}
int main(void)
{
    hello_from_gpu<<<2, 3>>>();
    cudaDeviceReset();
    return 0;
}

调用hello_from_gpu<<<2,3>>>,开启2个block，3个thread，共6个线程

其中每次使用一个线程，打印blockIdx.x,和threadIdx.x，且blockIdx为逆序（1到0），threadIdx为顺序(0-1-2)

image-20240313194115206
image-20240313194115206958×246 21.6 KB

执行顺序原因：

CUDA并行执行kernel时，其执行顺序是不确定的。虽然我们在代码中按照某种顺序启动了线程，但实际上这些线程在GPU上的执行顺序并不一定按照我们启动它们的顺序。这是因为GPU的调度器会根据各种因素（如资源可用性等）来决定具体执行哪个线程。

在我的例子中，虽然block 0的线程可能先启动，但block 1的线程可能先得到执行。这就是为什么你看到的输出结果可能先从block 1开始的原因。

这种情况在并行编程中非常常见，我们通常不能假定或依赖于特定的线程或block执行顺序。如果我们的程序需要某种特定的执行顺序，那么我们需要使用一些同步机制（如CUDA中的__syncthreads()函数）来确保正确的顺序。

program3

#include<stdio.h>
__global__ void hello_from_gpu()
{
    int gDim = gridDim.x;
    int bDim = blockDim.x;
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    printf("Hello World from block %d/%d and thread %d/%d!\n", bid,  gDim , tid, bDim);
}
int main(void)
{
    hello_from_gpu<<<2, 3>>>();
    cudaDeviceReset();
    return 0;
}

调用hello_from_gpu<<<2,3>>>,开启2个block，3个thread，共6个线程

其中每次使用一个线程，打印bid/gridDim， tid/bDim， 其中gridDim = 2， bDim = 3；

image-20240313194208630
image-20240313194208630969×247 23.7 KB

program4

#include<stdio.h>
__global__ void hello_from_gpu(){
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int bz = blockIdx.z;

    int gdx = gridDim.x;
    int gdy = gridDim.y;
    int gdz = gridDim.z;

    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int tz = threadIdx.z;

    int bdx = blockDim.x;
    int bdy = blockDim.y;
    int bdz = blockDim.z;

    printf("Hello World from block-[%d/%d, %d/%d, %d/%d] and thread-[%d/%d, %d/%d, %d/%d]!\n", 
    bx, gdx, by, gdy, bz, gdz, tx, bdx, ty, bdy,  tz, bdz);
}
int main(void){
    dim3 block_size(2, 3, 4);
    dim3 grid_size(2, 1, 1);
    hello_from_gpu<<<grid_size, block_size>>>();
    cudaDeviceSynchronize();
    printf("\nblock_size.x:%d, block_size.y:%d, block_size.z:%d", block_size.x, block_size.y, block_size.z);
    cudaDeviceReset();
    return 0;
}

通过dim3数据类型定义定义了三维的block和grid，调用hello_from_gpu函数打印三维的block和grid的执行进程数，最后打印block的三维信息

image-20240313194330253
image-202403131943302531047×705 77.9 KB

program5

运行以下代码，观察线程束，你觉得线程束(warp)的大小是多少？总共有多少线程束？

#include<stdio.h>
#define N 136
__global__ void hello_threads()
{
    int tid = threadIdx.x;
    printf("%d  ", tid);
}
 
__global__ void hello_blocks()
{
    int bid = blockIdx.x;
    printf("%d  ", bid);
}
 
int main() {
    printf("Order of threads: \n");
    hello_threads<<<1, N>>>();
    cudaDeviceSynchronize();
    printf("\n\n");
 
    printf("Order of blocks: \n");
    hello_blocks<<<N, 1>>>();
    cudaDeviceSynchronize();
    printf("\n");
    
    cudaDeviceReset();
    return 0;
}
  • 对于hello_tnread, 这个kernel 在一个block里启动了136个线程，由于一个warp包含32个线程，所以136个线程会被分为5个warp，前四个对应32个线程，最后一个warp包含8个线程
  • 对于hello_blocks,这个kernel 每个block只含一个线程，每个block都会形成一个warp，每个warp包含一个线程，故会形成136个warp

image-20240313195042973
image-202403131950429731207×271 24.1 KB

program6

输入以下代码，去除下述注释行，观察代码差异以及是否能编译运行。如果不能运行，为什么？

1.1打开注释// r1 = func1(tid);

1.2撤销上一步的操作，打开注释// r1 = func2(tid);

1.3撤销上一步的操作，打开注释// r1 = func3(tid);

2.1打开注释// // r2 = func1(nBlk);

2.2打开注释// r2 = func2(nBlk);

3.1打开注释// cube_gpu2<<<nGrid, nBlk>>>();

3.2 在保证前面代码正确运行的情况下（报错的行可以暂时先注释），打开注释// cudaDeviceSynchronize();分析该函数的作用。

#include<stdio.h> 
 
int func1(int x){
    return 2*x;
}
 
__device__ int func2(int x){
    return 2*x;
}
 
__host__ __device__ int func3(int x){
    return 2*x;
}
 
void __global__ cube_gpu1(){
    int tid = threadIdx.x;
    int r1;
    // r1 = func1(tid);
    // r1 = func2(tid);
    // r1 = func3(tid);
    printf("tid:%d, cube:%d\n", tid, r1);
}
 
__global__ void cube_gpu2(){
    int tid = threadIdx.x;
    int r1;
    r1= func3(tid);
    printf("tid:%d, cube:%d\n", tid, r1);
}
 
int main() 
{ 
    printf("Host and device functions!\n");
    printf("\nResults from device:\n");
    int nBlk = 3;
    int nGrid = 2;
    cube_gpu1<<<nGrid, nBlk>>>();      
    // cube_gpu2<<<nGrid, nBlk>>>();   
    // cudaDeviceSynchronize();
 
    int r2;
    // r2 = func1(nBlk);
    // r2 = func2(nBlk);
    r2 = func3(nBlk);
    printf("\nResults from host:%d\n", r2);
    r2 = func3(nGrid);
    printf("\nResults from host:%d\n", r2);
 
    cudaDeviceReset(); 
    return 0; 
}

1.1 打开注释// r1 = func1(tid);，代码将无法编译。

因为func1是一个只能在CPU上执行的函数，而这里试图在一个GPU kernel中调用它。(int)

1.2 打开注释// r1 = func2(tid);，代码将可以编译并运行。

因为func2是一个只能在GPU上执行的函数，而在一个GPU kernel中调用它是正确的。(_ device _)

1.3 打开注释// r1 = func3(tid);，代码将可以编译并运行。

因为func3是一个既可以在CPU上执行，也可以在GPU上执行的函数，在一个GPU kernel中调用它是正确的。(_ _ host _ _ _ _ device_ _)

2.1 打开注释// r2 = func1(nBlk);，代码将可以编译并运行。

因为func1是一个只能在CPU上执行的函数，在CPU的代码中调用它是正确的。

2.2 打开注释// r2 = func2(nBlk);，代码将无法编译。

因为func2是一个只能在GPU上执行的函数，而这里试图在CPU的代码中调用它。

3.1 打开注释// cube_gpu2<<<nGrid, nBlk>>>();，

代码将可以编译并运行。因为cube_gpu2是一个GPU kernel，可以在CPU的代码中启动它。

3.2 cudaDeviceSynchronize();函数用来阻塞CPU的执行，直到GPU完成所有之前的操作。

如果在启动GPU kernel之后打开这个注释，那么CPU将会等待，直到cube_gpu2完成执行。这样，就可以确保在cudaDeviceSynchronize();之后的CPU代码能正确地访问GPU上的数据。

program7

#include<stdio.h>
int main() {
    int dCount;
    cudaGetDeviceCount(&dCount);
    for(int i=0; i<dCount; i++)
    {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        printf("CUDA Device#%d\n", i);
        printf("Device name:%s\n", prop.name);
        printf("multiProcessorCount:%d\n", prop.multiProcessorCount);
        printf("maxThreadsPerBlock:%d\n", prop.maxThreadsPerBlock);
        printf("warpSize:%d\n", prop.warpSize);
        printf("maxThreadsDim[3]:%d, %d, %d\n", 
        prop.maxThreadsDim[0], 
        prop.maxThreadsDim[1], 
        prop.maxThreadsDim[2]);
        printf("maxGridSize[3]:%d, %d, %d\n", 
        prop.maxGridSize[0], 
        prop.maxGridSize[1], 
        prop.maxGridSize[2]);
    }
    cudaDeviceReset();
    return 0;
}

dCount是一个整型变量，用于存储系统中CUDA设备（GPU）的数量。在调用cudaGetDeviceCount(&dCount)函数后，dCount会被赋值为系统中CUDA设备的数量。&dCount是dCount变量的地址，这是因为cudaGetDeviceCount函数需要一个指针参数，以便在函数内部修改该变量的值。

  • Device name：这是设备的名称，例如 “Tesla V100”。
  • multiProcessorCount：这是设备上多处理器的数量。在CUDA中，一个多处理器包含一组CUDA核心，这些核心可以并行执行多个线程。
  • maxThreadsPerBlock：这是一个block可以包含的最大线程数量。在CUDA中，线程被组织成一个三维的线程块（block），一个block中的所有线程可以并行执行。
  • warpSize：这是一个warp的大小，也就是一个warp包含的线程数量。在CUDA中，一个warp是并行执行的最小单位，一个warp中的所有线程会同时执行相同的指令。
  • maxThreadsDim[3]：这是一个block的最大尺寸，也就是一个block可以包含的最大线程数量在每个维度上的值。在CUDA中，线程被组织成一个三维的线程块（block）。
  • maxGridSize[3]：这是一个grid的最大尺寸，也就是一个grid可以包含的最大block数量在每个维度上的值。在CUDA中，block被组织成一个三维的grid。

program8

尝试利用内置(build-in)变量来设计出一种在任意给定numOfBlocks和numOfThreads参数值的情况下，都可以唯一标识每个线程的编号方法。

例如：当numOfBlocks=3，numOfThreads=4时，0表示tid为0并且bid为0的线程的唯一标识，1表示tid为1并且bid索引为0的线程的唯一标识，…，4表示tid为0并且bid索引为1的线程的唯一标识，依此类推，如下图所示。

image-20240313195550008
image-20240313195550008886×703 158 KB 

#include<stdio.h>
__global__ void cuda_hello(void)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    printf("bid:[%d], tid:[%d] Hello World from GPU!\n", bid, tid);
}
 
int main() 
{
    int numOfBlocks = 5;
    int numOfThreads = 3;
    cuda_hello<<<numOfBlocks , numOfThreads>>>();
    cudaDeviceReset();
    return 0;
}

通过组合blockIdx.x，blockDim.x和threadIdx.x这三个内置变量来为每个线程生成一个唯一的编号。这可以通过以下公式实现：

Unique Thread ID=blockIdx.x×blockDim.x+threadIdx.xUnique Thread ID=blockIdx.x×blockDim.x+threadIdx.x

假设你有3个线程块，每个线程块有4个线程。那么：

  • 对于第一个线程块（blockIdx.x = 0），线程ID将是0, 1, 2, 3（对应于threadIdx.x = 0, 1, 2, 3）。
  • 对于第二个线程块（blockIdx.x = 1），线程ID将是4, 5, 6, 7（因为blockIdx.x * blockDim.x + threadIdx.x将是4 + 0, 4 + 1, 4 + 2, 4 + 3）。
  • 对于第三个线程块（blockIdx.x = 2），线程ID将是8, 9, 10, 11（因为blockIdx.x * blockDim.x + threadIdx.x将是8 + 0, 8 + 1, 8 + 2, 8 + 3）。
//添加：
__global__ void cuda_hello(void)
{
    int bid = blockIdx.x;
    int tid = threadIdx.x;
    int idx = bid * blockDim.x + tid;
    printf("idx:[%d], bid:[%d], tid:[%d] Hello World from GPU!\n",idx, bid, tid);
}

image-20240313195757914
image-202403131957579141092×547 51.9 KB

program9

素数(prime number)定义：是指在大于1的自然数中，除了1和它本身以外不再有其他因数的自然数。

请实现GPU版本的素数筛选函数prime_number_gpu，完成代码中的//ToDo部分并提交运行结果截图以及代码到实验报告。

(注意：不能修改其他代码，可参考以下代码中的CPU版本函数prime_number_cpu)。

在实验报告中记录完整代码（使用编辑器中的“代码语言”->"C/C++"后输入代码）以及运行结果。

#include<stdio.h>
__global__ void prime_number_gpu(void)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x; //唯一标识
    int is_prime = 1;
    if(idx < 2)
        return;
    for(int i = 2; i < idx; i++)
    {
        if(idx %i==0)
            {
                is_prime = -1;
                break;
            }  
    }
    if(is_prime== 1)
    {
       printf("%d  ", idx);
    }
 
}
 
void prime_number_cpu(int x)
{
    for(int p=2; p<x; p++)
    {
        int is_prime = 1;
        for(int i=2; i<p; i++)
        {   
            if(p%i==0)
            {
                is_prime = -1;
                break;
            }      
        }
        if(is_prime==1)
        {
            printf("%d  ", p);
        }
    }
}
 
int main() 
{
    int numUpperBound = 50;
    printf("CPU version:\n");
    prime_number_cpu(numUpperBound);
    printf("\nGPU version:\n");
    prime_number_gpu<<<1, numUpperBound>>>();
       printf("\n");
    cudaDeviceReset();
    return 0;
}

image-20240314150604059
image-20240314150604059964×212 14.8 KB

program10

水仙花数（Narcissistic number）也被称为超完全数字不变数（pluperfect digital invariant, PPDI）、自恋数、自幂数、阿姆斯壮数或阿姆斯特朗数（Armstrong number），水仙花数是指一个 3 位数，它的每个位上的数字的 3次幂之和等于它本身（例如：1^3 + 5^3+ 3^3 = 153）。

四叶玫瑰数（Four-leaf rose number）是指四位数各位上的数字的四次方之和等于本身的数（例如：1^4 + 6^4+ 3^4+4^4 = 1634）。

请实现GPU版本的水仙花数和四叶玫瑰数筛选函数narcissistic_and_four_leaf_rose_number_gpu，完成代码中的//ToDo部分。

(注意：不能修改其他代码，可参考以下代码中的CPU版本函数narcissistic_and_four_leaf_rose_number_cpu)。

在实验报告中记录完整代码（使用编辑器中的“代码语言”->"C/C++"后输入代码）以及运行结果。

#include<stdio.h>
__global__ void narcissistic_and_four_leaf_rose_numbers_gpu(int x1, int x2)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int num = index + x1;
    if (num >= x1 && num < x2) {
        int tho, hun, ten, ind;
        tho = num / 1000;
        hun = num / 100;
        ten = (num-hun*100) / 10;
        ind = num % 10;
        if (tho > 0) {
            hun = (num-tho*1000) / 100;
            ten = (num-tho*1000-hun*100) / 10;
            ind = num % 10;
            if (num == tho*tho*tho*tho + hun*hun*hun*hun + ten*ten*ten*ten + ind*ind*ind*ind) {
                printf("%d  ", num);
            }
        } 
        else {
            if (num == hun*hun*hun + ten*ten*ten + ind*ind*ind) {
                printf("%d  ", num);
            }
        }
    }
}
void narcissistic_and_four_leaf_rose_numbers_cpu(int x1, int x2)
{
    int tho, hun, ten, ind;
    for(int i=x1; i<x2; i++)
    {
        tho = i/1000;
        hun = i/100;
        ten = (i-hun*100)/10;
        ind = i%10;
        if(tho>0)
        {
            hun=(i-tho*1000)/100;
            ten=(i-tho*1000-hun*100)/10;
            ind=i%10;
            if(i==tho*tho*tho*tho + hun*hun*hun*hun + ten*ten*ten*ten + ind*ind*ind*ind)
            {
                printf("%d  ", i);
            }
        }
        else
        {
            if(i==hun*hun*hun + ten*ten*ten + ind*ind*ind)
            {
                printf("%d  ", i);
            }
        }
    }
}
 
int main() 
{
    int numLowerBound = 100;
    int numUpperBound = 10000;
    printf("Narcissistic and four-leaf rose numbers from %d to %d (CPU version):\n", numLowerBound, numUpperBound);
    narcissistic_and_four_leaf_rose_numbers_cpu(numLowerBound, numUpperBound);
    printf("\nNarcissistic and four-leaf rose numbers from %d to %d (GPU version):\n", numLowerBound, numUpperBound);
    int numNumbers = numUpperBound - numLowerBound;
    int numThreadsPerBlock = 256;
    int numBlocks = (numNumbers + numThreadsPerBlock - 1) / numThreadsPerBlock;
    narcissistic_and_four_leaf_rose_numbers_gpu<<<numBlocks, numThreadsPerBlock>>>(numLowerBound, numUpperBound);
    cudaDeviceReset();
    return 0;
}

image-20240314153412487
image-202403141534124871819×193 29.4 KB

#3

很好的CUDA，让我的GPU旋转