0%

我发现 optimizer.step() 这一步会改变 torch.nn.BatchNorm2d 层的 weight 和 bias 的梯度(看上去是每个元素按照相同的比例进行了缩放)。如果想要比较梯度,应该在 optimizer.step() 之前来对比,想要对比更新后的权重,要在 optimizer.step() 之后对比。

这很不应该啊?

版本 2.4.1 + CUDA 11.7(自行构建)上和版本 2.2.2 + CUDA 11.8 上都验证过会出现这种情况。

使用 debugpy 调试

在被 C++ 调用的 python 文件中加上这样的内容(端口可以随便选):

import debugpy
debugpy.listen(5678)
debugpy.wait_for_client()

这样在第一次加载这个 python 模块的时候,这行代码就会暂停等待调试器连接。我们可以创建这样的 Python 调试配置,随后连接正在监听端口的 python 程序。

{
  "version": "0.2.0",
  "configurations": [
    {
      "name": "Python attach",
      "type": "debugpy",
      "request": "attach",
      "connect": {
        "host": "127.0.0.1",
        "port": 5678
      }
    }
  ]
}

如果 C++ 程序的调试器还开着,工作区就会同时存在两个调试器。通过切换调试会话可以对 C++ 和 Python 同时调试。

注册钩子

通过注册钩子,收集模型运行时的输出,可以对每一层的输出值进行调试。

假设模型是 model,我们可以把每一层的输入的梯度和输出的梯度保存在字典中:

module_names = {v: k for k, v in model.named_modules()}
grad_inputs = {}
grad_outputs = {}

def hook(m, grad_input, grad_output) -> Tuple[torch.Tensor] | None:
    nonlocal module_names, grad_inputs, grad_outputs
    name = module_names[m]
    grad_inputs[name] = grad_input
    grad_outputs[name] = grad_output

for m in model.modules(): # 或者在这里遍历 named_modules() 并记录名称和模型的对应关系
    m.register_full_backward_hook(hook)

注意 grad_input 的含义是输入的梯度,grad_output 的含义是输出的梯度。假设一个 module 的 forward() 函数负责计算 y = 2 * x,那么 grad_input 相当于 retain_grad() 之后的 x.gradgrad_output 相当于 retain_grad() 之后的 y.grad

另外还要注意:

说明

文章是按照我解决问题的过程来写的,不是一个一步式的教程,所以显得有点凌乱。如果要操作请务必先看完全文,以免跟着中间过程走了同样的弯路。如果不想看前面的内容可以直接跳到 conda 打包这一节

编译和安装 PyTorch(egg 格式)

PyTorch 官方只为每个 PyTorch 版本准备了几个可选的 CUDA 版本,如果需要对 PyTorch 使用不同的 CUDA 就需要自己从源码中编译。我们想要用 PyTorch 2.4.1 源码构建支持 CUDA 11.7 的包。

参考官方的说明 https://github.com/pytorch/pytorch#from-source ,过程非常简单,如果下载顺畅,在 32 核服务器上需要近一个小时编译完成。我下载的版本是 v2.4.1。总体流程是先准备好一个 conda 环境,只安装好 python 即可。然后根据说明安装各种东西,最后用 python setup.py install 或者 python setup.py develop 来安装 PyTorch。

Note

如何选择 installdevelop

  • 根据 https://stackoverflow.com/a/26588871/develop 会在 site-packages 中创建一个 .egg-link 文件,将包的路径指向工作区,这样就可以通过在工作区修改代码来影响系统中安装的 Python 包。
  • 根据 setuptools/command/install.pyinstall 会调用命令 bdist_egg,会生成一个 egg 格式的发布包并安装,在 site-packages 对应的 egg 文件夹中确实有源码的一份副本。

注意 .egg 和 .egg-link 有区别!

然后 conda 环境中就会出现 PyTorch,而且用的 CUDA 版本也是系统里面的版本(如果用的镜像基于 NVIDIA 的 CUDA 镜像就肯定没问题),可以用 python -c "import torch; print(torch.version.cuda)" 来验证一下。不过,这时如果下载其他的 pip 包,pip 就会认为 torch 没有被安装,并且帮你去下载一个新的,下载完成之后 CUDA 的版本就变了!

今天写 atan 反向传播的 CUDA kernel 发现和 torch 算出来的不一样。核心代码如下:

__global__ void elementwise_atan_backward(float* in, float* din, float* dout, int N) {
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (N); i += blockDim.x * gridDim.x) {
    // din[i] = dout[i] / (in[i] * in[i] + 1.f);
    //
    // pytorch/tools/autograd/derivatives.yaml
    // - name: atan(Tensor self) -> Tensor
    //   self: grad / (self * self + 1).conj()
    //   result: auto_element_wise
    //
    float square = in[i] * in[i] + 0.f; // disable fma so we can align with torch
    din[i] = dout[i] / (square + 1.f);
  }
}

PyTorch 是用规则文件生成自动梯度求导的,torch.atan 的求导规则为 grad / (self * self + 1).conj(),其中 gradself 都是张量,conj() 在实数张量的场景下是可以不管的。我按照同样的方法写 kernel,发现和 torch 算出来的结果有极小的差异。

上面的代码 我也放在 Compiler Explorer 了,从汇编来看,in[i] * in[i] + 1.f 被优化成了一条指令,这样不仅更快,而且精度更高(从网上的资料来看,nvcc 默认会进行不少优化,fma 就是其中一种)。PyTorch 生成梯度求导规则的方法将乘法和加法放在了两个不同的 kernel 里面,导致编译器无法使用 fma 优化。

我尝试把乘法和加法拆开,但是编译器太聪明了,还是使用了 fma 优化:

通过 sudo usermod -aG docker xx 把 xx 加到 docker 组之后,需要先退出当前登录才能生效。

通过 mamba/conda 安装

mamba install -c mempo opencv

通过 pip 安装

pip install opencv-python opencv-python-headless

可能需要安装 opencv-contrib-python-headless,但是我这里不安装也成功了。

注意

两种安装方式不能混合使用。尤其是在通过 pip 安装 opencv-python 之后,发现不生效,是不能通过 conda 安装 opencv 来补救的。只能继续安装 opencv-python-headless。

安装 pybind11:

pip install 'pybind11[global]'
# add your executable
set(PYTHON_EXECUTABLE /opt/miniforge3/bin/python)
find_package(pybind11 REQUIRED)
target_link_libraries(your_executable PRIVATE pybind11::pybind11 pybind11::embed)

设置 PYTHON_EXECUTABLE 可以确保 pybind11 使用正确的 python 版本,这样才能在里面找到你安装好的其他的包。

释放 NAT 服务占用的端口

查看被占用的 TCP 端口,检查自己所用的服务端口是否被包含在结果中:

netsh interface ipv4 show excludedportrange protocol=tcp

重启 winnat 服务来释放大多数端口(需要管理员权限),验证 TCP 端口是否被释放:

net stop winnnat && net start winnat

设置动态端口号范围

看看应用程序的动态端口范围:

分类

我们平时说的超线程可能指同时多线程(SMT),也可能指 Intel 的超线程(HT)。

flowchart TD
    SMT("Simultaneous multithreading (SMT)")
    HT("Hyper-threading (HT)")
    TMT("Temporal multithreading, or super-threading")
    SMT -->|Intel's implementation| HT
    TMT <-->|SMT 同一条流水线能执行来自多个线程的指令,另外一个不能| SMT
	HT -.-> HT_note("两个逻辑线程共享包括 TLB 和 cache 等资源,但 CR3 等资源独立")

也存在一个物理核心对应超过 2 个逻辑线程的 SMT 实现,比如 IBM Power8: 12 cores, 8T each, (32 FX + 32 FP) registers per thread,FX 是定点数,FP 是浮点数。

更完整的分类:

We can consider that Super-Threading the same as Temporal multithreading (coarse-grain). I.e. types of multithreading is vary as follows - from shared anything to shared nothing in CPU-core between threads: None (fully software multithreading), TM (fine-grained) in barrel processors, TM (coarse-grained) - Super Threading, SMT - Hyper Threading, CMT (Chip Multithreading) - multicore with shared expensive resource (single FPU on some cores), CMP (Chip Multiprocessors) - multicore with shared Last Level Cache & interconnect.