BNU-FZH

Python版CUDA教程

发表于 2024-07-18 更新于 2024-09-26

本文主要介绍CUDA编程的基本流程和核心概念，并使用Python Numba编写GPU并行程序。

GPU硬件知识和基础概念：包括CPU与GPU的区别、GPU架构、CUDA软件栈简介。
GPU编程入门：主要介绍CUDA核函数，Thread、Block和Grid概念，并使用Python Numba进行简单的并行计算。
GPU编程进阶：主要介绍一些优化方法。
GPU编程实践：使用Python Numba解决复杂问题。

针对Python的CUDA教程

Python是当前最流行的编程语言，被广泛应用在深度学习、金融建模、科学和工程计算上。作为一门解释型语言，它运行速度慢也常常被用户诟病。著名Python发行商Anaconda公司开发的Numba库为程序员提供了Python版CPU和GPU编程工具，速度比原生Python快数十倍甚至更多。使用Numba进行GPU编程，你可以享受：

Python简单易用的语法；
极快的开发速度；
成倍的硬件加速。

为了既保证Python语言的易用性和开发速度，又达到并行加速的目的，本系列主要从Python的角度给大家分享GPU编程方法。关于Numba的入门可以参考我的另一篇文章。更加令人兴奋的是，Numba提供了一个GPU模拟器，即使你手头暂时没有GPU机器，也可以先使用这个模拟器来学习GPU编程！

初识GPU编程

兵马未动，粮草先行。在开始GPU编程前，需要明确一些概念，并准备好相关工具。

CUDA是英伟达提供给开发者的一个GPU编程框架，程序员可以使用这个框架轻松地编写并行程序。本系列第一篇文章提到，CPU和主存被称为主机（Host），GPU和显存（显卡内存）被称为设备（Device），CPU无法直接读取显存数据，GPU无法直接读取主存数据，主机与设备必须通过总线（Bus）相互通信。

在进行GPU编程前，需要先确认是否安装了CUDA工具箱，可以使用echo $CUDA_HOME检查CUDA环境变量，返回值不为空说明已经安装好CUDA。也可以直接用Anaconda里的conda命令安装CUDA：

1	$ conda install cudatoolkit

然后可以使用nvidia-smi命令查看显卡情况，比如这台机器上几张显卡，CUDA版本，显卡上运行的进程等。我这里是一台32GB显存版的Telsa V100机器。

安装Numba库：

1	$ conda install numba

检查一下CUDA和Numba是否安装成功：

1 2	from numba import cuda print(cuda.gpus)

如果上述步骤没有问题，可以得到结果：<Managed Device 0>...。如果机器上没有GPU或没安装好上述包，会有报错。CUDA程序执行时会独霸一张卡，如果你的机器上有多张GPU卡，CUDA默认会选用0号卡。如果你与其他人共用这台机器，最好协商好谁在用哪张卡。一般使用CUDA_VISIBLE_DEVICES这个环境变量来选择某张卡。如选择5号GPU卡运行你的程序。

1	CUDA_VISIBLE_DEVICES='5' python example.py

如果手头暂时没有GPU设备，Numba提供了一个模拟器，供用户学习和调试，只需要在命令行里添加一个环境变量。

Mac/Linux:

1	export NUMBA_ENABLE_CUDASIM=1

Windows:

1	SET NUMBA_ENABLE_CUDASIM=1

需要注意的是，模拟器只是一个调试的工具，在模拟器中使用Numba并不能加速程序，有可能速度更慢，而且在模拟器能够运行的程序，并不能保证一定能在真正的GPU上运行，最终还是要以GPU为准。

有了以上的准备工作，我们就可以开始我们的GPU编程之旅了！

GPU程序与CPU程序的区别

一个传统的CPU程序的执行顺序如下图所示：

CPU程序是顺序执行的，一般需要：

初始化。
CPU计算。
得到计算结果。

在CUDA编程中，CPU和主存被称为主机（Host），GPU被称为设备（Device）。

当引入GPU后，计算流程变为：

初始化，并将必要的数据拷贝到GPU设备的显存上。
CPU调用GPU函数，启动GPU多个核心同时进行计算。
CPU与GPU异步计算。
将GPU计算结果拷贝回主机端，得到计算结果。

一个名为gpu_print.py的GPU程序如下所示：

from numba import cuda

def cpu_print():
    print("print by cpu.")

@cuda.jit
def gpu_print():
    # GPU核函数
    print("print by gpu.")

def main():
    gpu_print[1, 2]()
    cuda.synchronize()
    cpu_print()

if __name__ == "__main__":
    main()

使用CUDA_VISIBLE_DEVICES='0' python gpu_print.py执行这段代码，得到的结果为：

1
2
3

print by gpu.
print by gpu.
print by cpu.

与传统的Python CPU代码不同的是：

使用from numba import cuda引入cuda库
在GPU函数上添加@cuda.jit装饰符，表示该函数是一个在GPU设备上运行的函数，GPU函数又被称为核函数。
主函数调用GPU核函数时，需要添加如[1, 2]这样的执行配置，这个配置是在告知GPU以多大的并行粒度同时进行计算。gpu_print[1, 2]()表示同时开启2个线程并行地执行gpu_print函数，函数将被并行地执行2次。下文会深入探讨如何设置执行配置。
GPU核函数的启动方式是异步的：启动GPU函数后，CPU不会等待GPU函数执行完毕才执行下一行代码。必要时，需要调用cuda.synchronize()，告知CPU等待GPU执行完核函数后，再进行CPU端后续计算。这个过程被称为同步，也就是GPU执行流程图中的红线部分。如果不调用cuda.synchronize()函数，执行结果也将改变，"print by cpu.将先被打印。虽然GPU函数在前，但是程序并没有等待GPU函数执行完，而是继续执行后面的cpu_print函数，由于CPU调用GPU有一定的延迟，反而后面的cpu_print先被执行，因此cpu_print的结果先被打印了出来。

Thread层次结构

前面的程序中，核函数被GPU并行地执行了2次。在进行GPU并行编程时需要定义执行配置来告知以怎样的方式去并行计算，比如上面打印的例子中，是并行地执行2次，还是8次，还是并行地执行20万次，或者2000万次。2000万的数字��大，远远多于GPU的核心数，如何将2000万次计算合理分配到所有GPU核心上。解决这些问题就需要弄明白CUDA的Thread层次结构。

CUDA将核函数所定义的运算称为线程（Thread），多个线程组成一个块（Block），多个块组成网格（Grid）。这样一个grid可以定义成千上万个线程，也就解决了并行执行上万次操作的问题。例如，把前面的程序改为并行执行8次：可以用2个block，每个block中有4个thread。原来的代码可以改为gpu_print[2, 4]()，其中方括号中第一个数字表示整个grid有多少个block，方括号中第二个数字表示一个block有多少个thread。

实际上，线程（thread）是一个编程上的软件概念。从硬件来看，thread运行在一个CUDA核心上，多个thread组成的block运行在Streaming Multiprocessor（SM的概念详见本系列第一篇文章），多个block组成的grid运行在一个GPU显卡上。

CUDA提供了一系列内置变量，以记录thread和block的大小及索引下标。以[2, 4]这样的配置为例：blockDim.x变量表示block的大小是4，即每个block有4个thread，threadIdx.x变量是一个从0到blockDim.x - 1（4-1=3）的索引下标，记录这是第几个thread；gridDim.x变量表示grid的大小是2，即每个grid有2个block，blockIdx.x变量是一个从0到gridDim.x - 1（2-1=1）的索引下标，记录这是第几个block。

某个thread在整个grid中的位置编号为：threadIdx.x + blockIdx.x * blockDim.x。

利用上述变量，我们可以把之前的代码丰富一下：

``from numba import cuda

def cpu_print(N):
    for i in range(0, N):
        print(i)

@cuda.jit
def gpu_print(N):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 
    if (idx < N):
        print(idx)

def main():
    print("gpu print:")
    gpu_print[2, 4](8)
    cuda.synchronize()
    print("cpu print:")
    cpu_print(8)

if __name__ == "__main__":
    main()`

执行结果为：

```py
``gpu print:
0
1
2
3
4
5
6
7
cpu print:
0
1
2
3
4
5
6
7`

这里的GPU函数在每个CUDA thread中打印了当前thread的编号，起到了CPU函数`for`循环同样的作用。因为`for`循环中的计算内容互相不依赖，也就是说，某次循环只是专心做自己的事情，循环第i次不影响循环第j次的计算，所以这样互相不依赖的`for`循环非常适合放到CUDA thread里做并行计算。在实际使用中，我们一般将CPU代码中互相不依赖的的`for`循环适当替换成CUDA代码。

这份代码打印了8个数字，核函数有一个参数`N`，`N = 8`，假如我们只想打印5个数字呢？当前的执行配置共2 \* 4 = 8个线程，线程数8与要执行的次数5不匹配，不过我们已经在代码里写好了`if (idx < N)`的判断语句，判断会帮我们过滤不需要的计算。我们只需要把`N = 5`传递给`gpu_print`函数中就好，CUDA仍然会启动8个thread，但是大于等于`N`的thread不进行计算。==注意，当线程数与计算次数不一致时，一定要使用这样的判断语句，以保证某个线程的计算不会影响其他线程的数据。==

![线程数与计算次数不匹配](https://p1-jj.byteimg.com/tos-cn-i-t2oaga2asx/gold-user-assets/2019/11/20/16e86a62f739c803~tplv-t2oaga2asx-jj-mark:3024:0:0:0:q75.png)

## Block大小设置

不同的执行配置会影响GPU程序的速度，一般需要多次调试才能找到较好的执行配置，在实际编程中，执行配置`[gridDim, blockDim]`应参考下面的方法：

-   block运行在SM上，不同硬件架构（Turing、Volta、Pascal...）的CUDA核心数不同，一般需要根据当前硬件来设置block的大小`blockDim`（执行配置中第二个参数）。一个block中的thread数最好是32、128、256的倍数。==注意，限于当前硬件的设计，block大小不能超过1024。==
-   grid的大小`gridDim`（执行配置中第一个参数），即一个grid中block的个数可以由总次数`N`除以`blockDim`，并向上取整。

例如，我们想并行启动1000个thread，可以将blockDim设置为128，`1000 ÷ 128 = 7.8`，向上取整为8。使用时，执行配置可以写成`gpuWork[8, 128]()`，CUDA共启动`8 * 128 = 1024`个thread，实际计算时只使用前1000个thread，多余的24个thread不进行计算。

注意，这几个变量比较容易混淆，再次明确一下：`blockDim`是block中thread的个数，一个block中的`threadIdx`最大不超过`blockDim`；`gridDim`是grid中block的个数，一个grid中的`blockIdx`最大不超过`gridDim`。

以上讨论中，block和grid大小均是一维，实际编程使用的执行配置常常更复杂，block和grid的大小可以设置为二维甚至三维，如下图所示。这部分内容将在下篇文章中讨论。

![Thread Block Grid](https://p1-jj.byteimg.com/tos-cn-i-t2oaga2asx/gold-user-assets/2019/11/20/16e86a62ef74074b~tplv-t2oaga2asx-jj-mark:3024:0:0:0:q75.png)

## 内存分配

前文提到，GPU计算时直接从显存中读取数据，因此每当计算时要将数据从主存拷贝到显存上，用CUDA的术语来说就是要把数据从主机端拷贝到设备端。CUDA强大之处在于它能自动将数据从主机和设备间相互拷贝，不需要程序员在代码中写明。这种方法对编程者来说非常方便，不必对原有的CPU代码做大量改动。

我们以一个向量加法为例，编写一个向量加法的核函数如下：

```py
py@cuda.jit
def gpu_add(a, b, result, n):
    # a, b为输入向量，result为输出向量
    # 所有向量都是n维
    # 得到当前thread的索引
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n:
        result[idx] = a[idx] + b[idx]

初始化两个2千万维的向量，作为参数传递给核函数：

n = 20000000
x = np.arange(n).astype(np.int32)
y = 2 * x
gpu_result = np.zeros(n)

# CUDA执行配置
threads_per_block = 1024
blocks_per_grid = math.ceil(n / threads_per_block)

gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)

把上述代码整合起来，与CPU代码做对比，并验证结果正确性：

from numba import cuda
import numpy as np
import math
from time import time

@cuda.jit
def gpu_add(a, b, result, n):
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n:
        result[idx] = a[idx] + b[idx]

def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x

    gpu_result = np.zeros(n)
    cpu_result = np.zeros(n)

    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))

    if (np.array_equal(cpu_result, gpu_result)):
        print("result correct")

if __name__ == "__main__":
    main()

运行结果，GPU代码竟然比CPU代码慢10+倍！

1
2
3

gpu vector add time 13.589356184005737
cpu vector add time 1.2823548316955566
result correct

说好的GPU比CPU快几十倍上百倍呢？这里GPU比CPU慢很多原因主要在于：

向量加法的这个计算比较简单，CPU的numpy已经优化到了极致，无法突出GPU的优势，我们要解决实际问题往往比这个复杂得多，当解决复杂问题时，优化后的GPU代码将远快于CPU代码。
这份代码使用CUDA默认的统一内存管理机制，没有对数据的拷贝做优化。CUDA的统一内存系统是当GPU运行到某块数据发现不在设备端时，再去主机端中将数据拷贝过来，当执行完核函数后，又将所有的内存拷贝回主存。在上面的代码中，输入的两个向量是只读的，没必要再拷贝回主存。
这份代码没有做流水线优化。CUDA并非同时计算2千万个数据，一般分批流水线工作：一边对2000万中的某批数据进行计算，一边将下一批数据从主存拷贝过来。计算占用的是CUDA核心，数据拷贝占用的是总线，所需资源不同，互相不存在竞争关系。这种机制被称为流水线。这部分内容将在下篇文章中讨论。

原因2中本该程序员动脑思考的问题交给了CUDA解决，增加了时间开销，所以CUDA非常方便的统一内存模型缺点是计算速度慢。针对原因2，我们可以继续优化这个程序，告知GPU哪些数据需要拷贝到设备，哪些需要拷贝回主机。

from numba import cuda
import numpy as np
import math
from time import time

@cuda.jit
def gpu_add(a, b, result, n):
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n :
        result[idx] = a[idx] + b[idx]

def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x

    # 拷贝数据到设备端
    x_device = cuda.to_device(x)
    y_device = cuda.to_device(y)
    # 在显卡设备上初始化一块用于存放GPU计算结果的空间
    gpu_result = cuda.device_array(n)
    cpu_result = np.empty(n)

    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x_device, y_device, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))

    if (np.array_equal(cpu_result, gpu_result.copy_to_host())):
        print("result correct!")

if __name__ == "__main__":
    main()

这段代码的运行结果为：

1
2
3

gpu vector add time 0.19940638542175293
cpu vector add time 1.132070541381836
result correct!

至此，可以看到GPU速度终于比CPU快了很多。

Numba对Numpy的比较友好，编程中一定要使用Numpy的数据类型。用到的比较多的内存分配函数有：

cuda.device_array()：在设备上分配一个空向量，类似于numpy.empty()
cuda.to_device()：将主机的数据拷贝到设备

1 2	ary = np.arange(10) device_ary = cuda.to_device(ary)

cuda.copy_to_host()：将设备的数据拷贝回主机

1	host_ary = device_ary.copy_to_host()

总结

Python Numba库可以调用CUDA进行GPU编程，CPU端被称为主机，GPU端被称为设备，运行在GPU上的函数被称为核函数，调用核函数时需要有执行配置，以告知CUDA以多大的并行粒度来计算。使用GPU编程时要合理地将数据在主机和设备间互相拷贝。

CUDA编程的基本流程为：

初始化，并将必要的数据拷贝到GPU设备的显存上。
使用某个执行配置，以一定的并行粒度调用CUDA核函数。
CPU和GPU异步计算。
将GPU计算结果拷贝回主机。

参考声明

GPU硬件知识和基础概念

发表于 2024-07-18 更新于 2024-09-26

金融建模、自动驾驶、智能机器人、新材料发现、脑神经科学、医学影像分析...人工智能时代的科学研究极度依赖计算力的支持。提供算力的各家硬件芯片厂商中，最抢镜的当属英伟达Nvidia了。这家做显卡起家的芯片公司在深度学习兴起后可红得发紫，如果不聊几句GPU和英伟达，都不好意思跟别人说自己是做人工智能的。那么，英伟达的GPU是如何加速计算呢？本系列将介绍GPU计算加速的一些基础知识：

GPU硬件知识和基础概念：包括CPU与GPU的区别、GPU架构、CUDA软件栈简介。
GPU编程入门：主要介绍CUDA核函数，Thread、Block和Grid概念，并使用Python Numba进行简单的并行计算。
GPU编程进阶：主要介绍多核配置和存储管理。
GPU编程实践：使用Python Numba解决复杂问题。

什么是GPU

GPU全名为Graphics Processing Unit，又称视觉处理器、图形显示卡。GPU负责渲染出2D、3D、VR效果，主要专注于计算机图形图像领域。后来人们发现，GPU非常适合并行计算，可以加速现代科学计算，GPU也因此不再局限于游戏和视频领域。

CPU和GPU

现代CPU处理数据的速度在纳秒级别，为何还要使用GPU来加速？CPU能被GPU替代吗？

对于计算机体系不了解的朋友可以先阅读我之前的文章，有助于你理解下面的一些概念。

无论是CPU还是GPU，在进行计算时，都需要用核心（Core）来做算术逻辑运算，比如加减乘与或非等。核心中有ALU（逻辑运算单元）和寄存器等电路。在进行计算时，一个核心只能顺序执行某项任务。不可能“吃着火锅唱着歌”，因为吃饭唱歌都占着嘴呢。所以为了同时并行地处理更多任务，芯片公司开发出了多核架构，只要相互之间没有依赖，每个核心做自己的事情，多核之间互不干扰，就可以达到并行计算的效果，极大缩短计算时间。

个人桌面电脑CPU只有2到8个CPU核心，数据中心的服务器上也只有20到40个左右CPU核心，GPU却有上千个核心。与CPU的核心不同，GPU的核心只能专注于某些特定的任务。知乎上有人把CPU比作大学教授，把GPU比作一个学校几千个小学生：同样是做加减法，几千个小学生所能做的计算，远比几十个大学教授要多得多。俗话说，三个臭皮匠，顶一个诸葛亮。大学教授的知识结构和个人能力远强于小学生，能独立解决复杂问题，小学生的知识有限，只能进行简单的计算。目前来看GPU在处理简单计算任务上有更大的优势，但是主要还是靠人海战术，并不能像CPU那样可以独当一面，短时间内也无法替换掉CPU。如下图所示，在整个计算机系统中，CPU起到协调管理的作用，管理计算机的主存、硬盘、网络以及GPU等各类元件。

如果只关注CPU和GPU，那么计算结构将如下图所示。CPU主要从主存（Main Memory）中读写数据，并通过总线（Bus）与GPU交互。GPU除了有超多计算核心外，也有自己独立的存储，被称之为显存。一台服务器上可以安装多块GPU卡，但GPU卡的发热量极大，普通的空调系统难以给大量GPU卡降温，所以大型数据中心通常使用水冷散热，并且选址在温度较低的地方。

GPU核心在做计算时，只能直接从显存中读写数据，程序员需要在代码中指明哪些数据需要从内存和显存之间相互拷贝。这些数据传输都是在总线上，因此总线的传输速度和带宽成了部分计算任务的瓶颈。也因为这个瓶颈，很多计算任务并不适合放在GPU上，比如笔者这两年关注的推荐系统虽然也在使用深度学习，但因为输入是大规模稀疏特征，GPU加速获得的收益小于数据互相拷贝的时间损失。当前最新的总线技术是NVLink，IBM的Power CPU和英伟达的高端显卡可以通过NVLink直接通信。同时，单台机器上的多张英伟达显卡也可以使用NVLink相互通信，适合多GPU卡并行计算的场景。

Intel的CPU目前不支持NVLink，只能使用PCI-E技术，如下图所示。NVLink和PCI-E都是总线技术的一种。

由于CPU和GPU是分开的，在英伟达的设计理念里，CPU和主存被称为Host，GPU被称为Device。Host和Device概念会贯穿整个英伟达GPU编程。

以上结构也被称为异构计算：使用CPU+GPU组合来加速计算。世界上顶尖的数据中心和超级计算机均采用了异构计算架构。例如超越天河2号成为世界第一的超级计算机Summit使用了9216个IBM POWER9 CPU和27648个英伟达Tesla GPU。

GPU架构

英伟达不同时代产品的芯片设计不同，每代产品背后有一个架构代号，架构均以著名的物理学家为名，以向先贤致敬。当前比较火热的架构有：

Turing 图灵
- 2018年发布
- 消费显卡：GeForce 2080 Ti
Volta 伏特
- 2017年末发布
- 专业显卡：Telsa V100 (16或32GB显存 5120个核心)
Pascal 帕斯卡
- 2016年发布
- 专业显卡：Telsa P100(12或16GB显存 3584个核心)

在英伟达的设计里，多个小核心组成一个Streaming Multiprocessor（SM），一张GPU卡有多个SM。从“multiprocessor”这个名字上也可以看出SM包含了多个处理器。实际上，英伟达主要以SM为运算和调度的基本单元。上图为当前计算力最强的显卡Tesla V100，密密麻麻的绿色小格子就是GPU小核心，多个小核心一起组成了一个SM。

单个SM的结构如图所示。可以看到一个SM中包含了：

针对不同计算的小核心（绿色小格子），包括优化深度学习的TENSOR CORE，32个64位浮点核心（FP64），64个整型核心(INT)，64个32位浮点核心(FP32)。
计算核心直接从寄存器（Register）中读写数据。
调度和分发器（Scheduler和Dispatch Unit）。
L0和L1级缓存。

前面提到的以物理学家命名的是英伟达各代GPU的架构代号。对于消费者而言，英伟达主要有两条产品线：

消费级产品 GeForce系列：GeForce 2080 Ti...
高性能计算产品 Telsa系列：Telsa V100、Telsa P100、Telsa P40...

软件生态

英伟达能够在人工智能时代成功，除了他们在长期深耕显卡芯片领域，更重要的是他们率先提供了可编程的软件架构。2007年，英伟达发布了CUDA编程模型，软件开发人员从此可以使用CUDA在英伟达的GPU上进行并行编程。在此之前，GPU编程并不友好。CUDA简单到什么程度？有经验的程序员经过半天的培训，掌握一些基础概念后，能在==半小时==内将一份CPU程序修改成为GPU并行程序。

继CUDA之后，英伟达不断丰富其软件技术栈，提供了科学计算所必须的cuBLAS线性代数库，cuFFT快速傅里叶变换库等，当深度学习大潮到来时，英伟达提供了cuDNN深度神经网络加速库，目前常用的TensorFlow、PyTorch深度学习框架的底层大多基于cuDNN库。英伟达能在人工智能时代击败Intel、AMD等强大对手，很大一部分是因为它丰富的软件体系。这些软件工具库使研发人员专注于自己的研发领域，不用再去花大量时间学习GPU底层知识。CUDA对于GPU就像个人电脑上的Windows、手机上的安卓系统，一旦建立好生态，吸引了开发者，用户非常依赖这套软件生态体系。

GPU编程可以直接使用CUDA的C/C++版本进行编程，也可以使用其他语言包装好的库，比如Python可使用Numba库调用CUDA。CUDA的编程思想在不同语言上都很相似。

CUDA及其软件栈的优势是方便易用，缺点也显而易见：

软件环境复杂，库以及版本很多，顶层应用又严重依赖底层工具库，入门者很难快速配置好一整套环境；多环境配置困难。
用户只能使用英伟达的显卡，成本高，个人用户几乎负担不起。

因此，如果没有专业的运维人员维护GPU机器，最好还是在公有云上按需购买GPU虚拟机。入门者可以考虑云厂商的Telsa P4虚拟机，大约10+元/小时，云厂商会配置好CUDA及工具库。如自己购买物理机，可以考虑消费级的GeForce 2080Ti，这张卡足以应对绝大多数自然语言处理任务。

Archlinux安装USB无线网卡驱动

发表于 2024-07-17 更新于 2024-09-26

今天获得了一个微形的USB无线网卡，其在Windows下是可以自动安装驱动程序的，因为此USB网卡中包含了Windows版的驱动程序。但是我现在想在Archlinux下使用这个USB无线网卡作为Wifi的发射无线信息供手机上网使用，而windows驱动就无能为力了。通过将USB无线网卡插入到windows电脑上配置好wifi就可以查询到此网卡的具体型号为realtek 8188gu wireless lan 802.11n usb nic linux, 而使用 pacman -Ss无法找到realtek相关驱动，于是使用paru查找:

paru -Ss realtek |grep 8188gu

    Driver for Realtek RTL8188CUS (8188C, 8192C) chipset wireless cards
aur/8188eu-dkms v4.3.0.8_13968-17 [+25 ~0.14]
    Driver for Realtek RTL8188EUS (RTL8188EUS, RTL8188ETV) WLAN
aur/8188eu-aircrack-dkms-git 5.3.9.r196.3fae723-5 [+7 ~0.00]
    Realtek RTL8188EUS and RTL8188ETV Wi-Fi driver with monitor mode & frame injection support
aur/8188eu-dkms-git 5.2.2.4.r794.53ed527-3 [+5 ~0.00]
    Standalone driver for Realtek RTL8188EUS (RTL8188EUS, RTL8188ETV) WLAN
aur/8188fu-kelebek333-dkms-git r115.751882b-1 [+4 ~0.00] [无人维护]
    Kernel driver module for Realtek RTL8188FU and RTL8188FTV Wireless 802.11n USB 2.0 adapters (device id 0bda:f179), for kernel versions < 6
aur/8188fu-dkms-git r13.586d795-1 [+3 ~0.00] [无人维护]
    Kernel driver module for Realtek RTL8188FU and RTL8188FTV Wireless 802.11n USB 2.0 adapters (device id 0bda:f179), for kernel versions less than 5.8
    Kernel module for Realtek RTL8188SU/RTL8191SU/RTL8192SU devices
aur/8188fu-supremegamers-dkms-git r120.40d4a49-1 [+2 ~0.00]
    Kernel driver module for Realtek RTL8188FU and RTL8188FTV Wireless 802.11n USB 2.0 adapters (device id 0bda:f179), for kernel versions 6.0 and 6.1
aur/rtl8188gu-dkms-git r10.699d0cc-1 [+2 ~0.00] [已安装：r12.9dec829-1]
    Linux driver for Realtek RTL8188GU

`rtl8xxxu`支持的`wifi`芯片

RTL8188CU/CUS/CTV
RTL8188EU/EUS/ETV
RTL8188FU/FTV
RTL8188GU | RTL8188RU
RTL8191CU | RTL8192CU 
RTL8192EU | RTL8192FU
RTL8723AU | RTL8723BU

支持Linux 内核版本: 5.5.x ~ 6.11.x
测试工作良好的Linux发行版：
- Arch Linux (kernel version: 6.6.39-1-lts)
- Debian 11.10 (kernel version: 5.10.0-30-amd64 / 6.1.0-0.deb11.21-amd64)
- Linux Mint 20.3 (kernel version: 5.15.0-113-generic)

安装驱动`rtl8xxxu`

克隆下rtl8xxxu的源文件：

1	git clone https://github.com/a5a5aa555oo/rtl8xxxu.git

如果您使用的USB Wifi发射器基于RTL8188GU或RTL8192FU芯片，您需要首先使用命令usb_modeswitch或eject切换到Wifi模式。
安装构建模块需要的软件包：gcc,make,linux-headers,dkms和其他的依赖软件。
构建和安装模块
- 使用传统方式
  
  make clean modules && sudo make install
- 使用 DKMS 方式(推荐)
  
  sudo dkms install $PWD
- 基于ArchLinux的发行版用户
  
  在AUR中安装软件包： rtl8xxxu-dkms-git
  1
  paru -S rtl8xxxu-dkms-git
为芯片 RTL8188EU/RTL8188FU/RTL8188GU/RTL8192EU/RTL8192FU 安装固件(关键)

sudo make install_fw
载入模块

sudo modprobe rtl8xxxu_git

安装成功后，可以在设置→Wi-Fi中找到网卡 Realtek RTL8188GU,顺利完成无线网卡驱动的安装。

2024年07月17日, 在办公室的台式机90M2CTO1WW (QiTianM428-N000)测试失败，或许是USB版的网卡不稳定？由于时间的关系，不再研究无线wifi在linux上的启用。
2024年08月12日, 在办公室的台式90M2CTO1WW(QiTianM428-N000)再次测试，重新安装rtl8188gu-dkms-git之后仍然将网卡识别为光盘，但是执行sudo modprobe rtl8188之后再重启，同时更换的插入的USB接口，发现无线网卡启动成功。
2024年08月13日, 通过rtl8xxxu-dkms官网https://kkgithub.com/a5a5aa555oo/rtl8xxxu安装固件.

安装驱动`rtl8xxxu-dkms-git`

根据文章RTL8188GU驱动安装, rtl8188gu-dkms-git有已知性能问题，在用户态可以看到一个RTW_CMD_THREAD。我使用时延迟很大并且使用一会就会断开，只能通过拔插恢复。同时也说明rtl8xxxu-dkms-git由Linux主线移植的驱动，性能优异，于是决定卸载rtl8188gu-dkms-git后安装rtl8xxxu-dkms-git。

rtl8xxxu驱动安装

1 2	paru -S rtl8xxxu-dkms-git sudo modprobe rtl8xxxu

~~如果已经安装了前面的McMCCRU版本的驱动，可以在/etc/modprobe.d/blacklist中添加一行blacklist 8188gu屏蔽旧驱动。~~

安装 dnsmasq

当安装好rtl8xxxu-dkms-git后，发现使用Gnome自带的开启wifi热成功，但是使用手机连接热点总是失败，同时wifi热点也随之关闭。原因是没有安装dnsmasq，它的主要作用是提供 DNS 缓存和 DHCP 服务功能。作为域名解析服务器(DNS)，dnsmasq可以通过缓存 DNS 请求来提高对访问过的网址的连接速度。作为DHCP 服务器，dnsmasq 可以用于为局域网电脑分配内网ip地址和提供路由。DNS和DHCP两个功能可以同时或分别单独实现。dnsmasq轻量且易配置，适用于个人用户或少于50台主机的网络。此外它还自带了一个 PXE 服务器。所以如果没有这个服务，那wifi热点就无法为手机分配 dns, 也就无法启动wifi热点。安装此服务即可：

1	sudo pacman -S dnsmasq

根据文章Arch Linux上简单地创建wifi热点（只用一块无线网卡）, 共享wifi需要安装linux-wifi-hotspot, 但是按其配置好我的电脑并不能通过linux-wifi-hotspot共享网络，或许在其文章底部的~~当然不是用来共享校园网~~说明了问题。于是果断删除此软件,在安装好dnsmasq之后，使用Gnome自带的创建wifi共享网络成功。

为何有些M-2固态硬盘会有两个缺口

发表于 2024-07-16 更新于 2024-09-26

提到M.2，大家都知道它是固态硬盘接口，不过并不是所有M.2接口都能安装各种类型的固态硬盘。为什么有些M.2固态硬盘接口上会有两个缺口，而有些则只有一个？其中的诸多的讲究，让我们一起了解下。

我们平时看到的M.2模块大多是M.2 2242或M.2 2280规格，除此之外M.2还提供了多种标准规格。一个基本的原则是小体积可以安装到大体积插槽内，M.2 2242就比M.2 2280更灵活。

并不是所有的M.2接口都是给固态硬盘使用的，比如下图中的M.2 E Key插槽就是留给无线网卡专用。仔细观察你会发现E Key的缺口位置和普通固态硬盘上的B Key有所不同，M.2插槽就是通过缺口的位置来表达自己能够支持的M.2设备类型。

常见的M.2固态硬盘分为一个缺口（M Key）和两个缺口（B Key）两种类型。下图中从上到下分别是东芝饥饿鲨RD400 NVMe固态硬盘、东芝Q200 240G M.2固态硬盘和东芝RC100 NVMe固态硬盘，分别对应PCIE NVMe、SATA AHCI、PCIE NVMe接口。也就是说，两个缺口的M.2固态硬盘既有可能是NVMe协议，也有可能为SATA协议，而只有一个缺口（M Key）的通常只有NVMe协议一种可能。

与固态硬盘相对应的是，M.2插槽上的卡口位置也有一些不同。下图比较了Socket 2和Socket 3两种最常见M.2插槽。Socket 2插槽支持PCIe x2或SATA通道，而Socket 3插槽则支持PCIe x4或SATA通道。

现在结论就比较明朗了，由于插槽卡口的限制，两个缺口（B+M Key）的固态硬盘可以安装到Socket 2或Socket 3两种形式的M.2插槽中，而一个缺口（M Key）的固态硬盘只能安装到Socket 3插槽。

两个缺口的固态硬盘只能支持PCIe x2通道，在持续读写带宽上会有所限制，但决定日常使用性能的4K随机读写效能却不受影响。

虽说x2通道对性能会有所限制，但也不乏优秀产品。譬如东芝的迷你NVMe SSD——RC100 就使用了B+M Key规格，在保持优良的兼容性的同时，4K随机读取速度达到58.55MB/s，基本和高端的旗舰级NVMe固态硬盘保持一致。

受益于B+M Key的双卡口特点，RC100甚至能够安装到为4GB上网卡（B Key接口）预留的M.2 2242插槽中并正常工作，已有玩家在Thinkpad T470p等机型中验证了这个特殊用法。

总结来说:

从极致性能出发，一个缺口（M Key）的M.2 NVMe固态硬盘最强（不是100%一定，具体还要看主控规格）；
从兼容性出发，两个缺口的M.2 SATA固态硬盘更好（如Q200 M.2版）；
要兼顾性能和兼容性，两个缺口的M.2 NVMe固态硬盘（如RC100）比较理想。

Riccati方程

发表于 2024-07-14 更新于 2024-09-26

十七世纪，意大利数学家黎卡提提出如下方程：

\[\begin{equation}\label{eq:riccati} y'=P(x)y^2+Q(x)y+R(x) \end{equation}\]

式$\eqref{eq:riccati}$称为Riccati方程。1841年法国数学家刘维尔证明了黎卡提(Riccati)方程一般没有初等解法，但是很多实际问题与理论问题又迫切需要求得这个方程的解，这也使得这一方程成为世界著名难题。

黎卡提方程自从十七世纪黎卡提提出以来，历经三百多年一直未有一般解法，虽然有众多特例解法，但是都未能从根本上解决这个方程。但是在物理上它的地位相当重要，所以解决这一方程仍是一个重要的任务。经过简单化简，我们可以给出一个更加简洁对称的形式，即：

\[\begin{equation}\label{eq:riccati1} y'+y^2+R(x)=0 \end{equation}\]

只要能够求解出式$\eqref{eq:riccati1}$则式$\eqref{eq:riccati}$的解也就可以得到了。由于其异常困难，我将此方程单独放在这里，希望获得式$\eqref{eq:riccati1}$的解。2024年07月14日

Shell脚本中识别方向键

发表于 2024-07-13 更新于 2024-09-26

在vim中使用H、J、K、L移动光标，但是有时候人们习惯于使用方向键，这在有一些脚本中实现移动光标选择选项时比较重要。在网络上找到两个方案：

方案一

#!/bin/bash
akey=(0 0 0)

cESC=`echo -ne "\033"`

while :
do
#这里要注意，这里有-n选项，后跟一个数字，指定输入的字符长度最大值，所以不管key变量的值有多少个字符，都是一个值一个值做循环， 有多少个值就循环多少次
#输入键盘的上下左右键时，都有三个字符(ESC键算一个)，所以都是做三次循环，每做一次循环，akey的值都会发生改变
      read -s -n 1 key

      akey[0]=${akey[1]}
      akey[1]=${akey[2]}
      akey[2]=${key}

      if [[ ${key} == ${cESC} && ${akey[1]} == ${cESC} ]]
      then
         echo "ESC键"
      elif [[ ${akey[0]} == ${cESC} && ${akey[1]} == "[" ]]
      then
           if [[ ${key} == "A" ]];then echo "上键"
           elif [[ ${key} == "B" ]];then echo "向下"
           elif [[ ${key} == "D" ]];then echo "向左"
           elif [[ ${key} == "C" ]];then echo "向右"
           fi
      fi
done

方案二

因为上下左右键当中分别含有ABDC，可以利用这一点来做while循环控制输出上下左右键的信号

#!/bin/bash
key=() #----->自定义一个数组key，为空值
while :
do
     read -s -n 1 key  #----->这里跟第一种方法类似，就是利用while循环，不过这里有点不一样的地方，就是这个不考虑完整方向键的字符
     #的方向键字符组成只要循环中包含A就会输出上键，包含B就会输出下键，以此类推，所以个人觉得这种方法没有第一种方法严谨，第一种方法更适合用来写比较大的脚本

     case ${KEY[0]} in
     "B")
           echo "下键"
           ;;
     "A")
           echo "上键"
           ;;
     "D")
           echo "左键"
           ;;
     "C")
           echo "右键"
           ;;
       *)
           continue
     esac
done

方案三

通过分析上述两种方案，我们给出第三种方案，即识别[A、[B、[C、[D 也可以区分方向键。代码如下:

NEO_ESC=`echo -ne "\033"`
echo -n "请输入按键"; read -s -n 1 key
if [[ $key == ${NEO_ESC} ]]; then
   read -sn 2 SubListDo
   case "$SubListDo" in
       "[A")
           echo "向上"
           ;;
       "[B")
           echo "向下"
           ;;
       "[C")
           echo "向右"
           ;;
       "[D")
           echo "向左"
           ;;
       *)
           exit
           ;;
   esac
fi

2024年07月13日, 由于在脚本diary.sh和zugit.sh中菜单使用了单一的vim方案，在没有实现完美的方向键识别前暂不加入方向键移动光标选择功能。
2024年07月14日, 在上述方案中首先检测是否输入特殊字符，若为特殊字符前导符\033[, 则读取后面两个字符，再对后两个字符分类探测。使用此方法成功在diary.sh和zugit.sh中引入方向键,home,end,pgup,pgdn控制列表显示，使之成为更加友好的同时兼容vim快捷键的脚本程序。

计算机求导数的数学理论

发表于 2024-07-12 更新于 2024-09-26

对于解析函数，可以将其Taylor展开为： \[\begin{equation}\label{eq:taylor} f(x_{i+1})=f(x_i)+f'(x_i)h+\frac{f''(x_i)}{2}h^2+\cdots \end{equation}\] 现在据式$\eqref{eq:taylor}$来计算多阶导数，取近似到二阶项，则一阶导数为 \[\begin{equation}\label{eq:taylor1} f'(x_i)=\frac{f(x_{i+1})-f(x_i)}{h}+O(h) \end{equation}\] 将式$\eqref{eq:taylor1}$中的函数$f$替换为相应的导数$f'$得二阶导数,然后出现的一阶导数使用式$\eqref{eq:taylor1}$替换，从而得到二阶导数用零阶导数表达的算式 \[\begin{equation}\label{eq:taylor2} f''(x_i)=-\frac{f(x_{i+2})-2f(x_{i+1})+f(x_i)}{h^{2}}+O(h) \end{equation}\] 按前述方法可以求出任意阶导数使用零阶导数表达的算式，于是可以将式$\eqref{eq:taylor}$取更多的项以提高其导数$f'(x_i)$的精度，比如取到二阶项 \[\begin{equation}\label{eq:taylor3} f'(x_i)=\frac{f(x_{i+1})-f(x_i)}{h}-\frac{f(x_{i+2})-2f(x_{i+1})+f(x_i)}{2h^{2}}h+O(h^2) \end{equation}\] 化简式$\eqref{eq:taylor3}$得 \[\begin{equation}\label{eq:taylor4} f'(x_i)=\frac{-f(x_{i+2})+4f(x_{i+1})-3f(x_i)}{2h}+O(h^2) \end{equation}\]

tar命令详解

发表于 2024-07-12 更新于 2024-09-26

tar的选项参数

-c: 建立压缩档案
-x：解压
-t：查看内容
-r：向压缩归档文件末尾追加文件
-u：更新原压缩包中的文件

这五个是独立的命令，压缩解压都要用到其中一个，可以和别的命令连用但只能用其中一个。下面的参数是根据需要在压缩或解压档案时可选的。

-z：有gzip属性的
-j：有bz2属性的
-Z：有compress属性的
-v：显示所有过程
-O：将文件解开到标准输出

--delete从存档中删除注意是两个减号。

下面的参数-f是必须的
-f: 使用档案名字，切记，这个参数是最后一个参数，后面只能接档案名。

举例说明

压缩一个目录

1	eg: tar -cvzf test.tar.gz tar #压缩当前目录下的tar目录，压缩包的名字叫test.tar.gz

压缩其中某一类型的文件

1	eg: tar -cvzf test.tar.gz ./*.c #压缩当前目录下的.c为后缀的文件。但是不能压缩以.开头的文件

查看一个压缩文件的内容

1	eg: tar -tvf test.tar.gz #查看当前目录下test.tar.gz压缩文件中的文件

解压一个文件

1	eg. tar -xvzf test.tar.gz #解压test.tar.gz到当前目录下面

文件夹做一个归档文件、备份一个目录、

1	eg: tar -cvf test.tar tar #将tar文件夹下的文件生成一个test.tar归档文件

将某一类型的文件做成归档文件、备份某一类型的文件

1	eg:tar -cvf test.tar ./tar/*.c # 将tar文件夹下所有以.c结尾的文件放到test.tar 文件中，但是以.开头的文件不行

向归档文件中添加一个文件夹

1	eg: tar -rf test.tar ./test1 #将test1文件夹中的所有文件，添加到test.tar 归档文件中

删除归档文件中的文件

1	eg: tar -f test.tar --delete a.c #删除test.tar文件中的 a.c文件

删除归档文件中的某一类的文件

1	eg: tar -f test.tar --delete *.c #删除test.tar文件中的以.c结尾的文件

释放归档文件

1	eg: tar -xvf test.tar #释放test.tar归档文件中的内容到当前目录下

查看归档文件中的内容
1
eg: tar -tvf test.tar

总结

格式与解压命令

1、*.tar           用 tar –xvf 解压
2、*.gz            用 gzip -d或者gunzip 解压
3、*.tar.gz和*.tgz 用 tar –xzf 解压
4、*.bz2           用 bzip2 -d或者用bunzip2 解压
5、*.tar.bz2       用tar –xjf 解压
6、*.Z             用 uncompress 解压
7、*.tar.Z         用tar –xZf 解压
8、*.rar           用 unrar e解压
9、*.zip           用 unzip 解压

使用举例

tar cvf  etcbak.tar etc/      打包一个tar
tar xvf  etcbak.tar           解开一个tar
tar cvzf etcbak.tar.gz etc/   打包压缩一个 tar
tar zxvf etcbak.tar.gz        解压一个tar

关于LaTeX排版系统的说明

发表于 2024-07-09 更新于 2024-09-26

LaTeX排版系统主要有CTeX套装和TeXLive, 然而我们建议大家尽量使用TeXLive, 原因是：

CTeX套装

CTeX套装是科学院吴凌云研究员的个人作品。在 CTeX套装刚刚问世之时，因其解决了繁琐的中文字体安装工作，所以广受欢迎。但是，一方面 CTeX套装已经很久不更新，内里的宏包、工具陈旧；另一方面，随着 XeLaTeX 的发展，以及 xeCJK 等技术的成熟，上述这些繁琐的工作已经没有必要而失去意义；因此，现在不推荐使用CTeX套装.

CTeX宏集

虽然它的名字也是「CTeX」，但是 CTeX 宏集和 CTeX 套装是两个不同的东西。CTeX 宏集是集成了中文支持、操作系统判定、字体选择、版式预设为一体的一组宏包和文档类的合集。我们推荐在任何情况下，优先使用 CTeX 宏集处理中文。

TeXLive

TeXLive 是 TUG (TeX User Group) 维护和发布的 TeX 系统，可说是「官方」的 TeX 系统。我们推荐任何阶段的 TeX 用户，都尽可能使用 TeXLive，以保持在跨操作系统平台、跨用户的一致性。TeXLive 的官方站点是 https://tug.org/texlive/.

归纳总结

不要安装和使用 CTeX 套装！
请在任何情况下优先使用 CTeX 宏集在 LaTeX 中处理中文！
具体安装配置请参考：ArchLinux安装指南

LaTeX中文档类revtex4中禁止引用文献中显示脚注

发表于 2024-07-08 更新于 2024-09-26

简介

Includes styles for American Physical Society, American Institute of Physics, and Optical Society of America. The distribution consists of the RevTeX class itself, and several support packages.

文献引用显示和隐藏脚注

在使用revtex4编写论文时，发现所有的footnote都会在引用文献中出现，调用其说明文档revtex4.pdf后发现其使用规则并没有详细介绍。于是经过一翻查找得到设置方法：即使用选项footinbib实现在参考文献列表中出现footnote的内容，使用选项nofootinbib则不在参考文献中出现footnote. 即

1 2	\documentclass[footinbib]{revtex4} %默认启用此选项 \documentclass[nofootinbib]{revtex4} %关闭需手动设置