分析NVIDIA开源GPU驱动

按照是否与操作系统直接相关来分类代码

NVIDIA GPU驱动的大多数内核模块可分为两个组件（部分）：

• 一个“与操作系统无关”的组件：这是独立于操作系统的每个内核模块的组件。

• 一个“内核接口层”：这是特定于Linux内核版本和配置的每个内核模块的组件。

当把编译好的驱动模块（也可以叫子驱动）打包在NVIDIA.run安装包（比如NVIDIA-Linux-x86_64-560.35.03.run）中时，与操作系统无关的组件以二进制形式提供，因为这些代码编译起来非常耗时，因此NVIDIA提供了预构建版本，这样用户就不必在每次安装驱动程序时都进行编译。

比如对于nvidia.ko内核模块，这部分组件被命名为“nv-kernel.o_binary”。对于nvidia-modeset.ko内核模块，这部分组件被命名为“nv-modeset-kernel.o_binary”。nvidia-drm.ko和nvidia-uvm.ko都没有与操作系统无关的组件。

代码目录结构

驱动的代码大都位于根目录下的kernel-open和src目录中。

kernel-open目录中的代码为内核接口层。

kernel-open/nvidia/

nvidia.ko的内核接口层

kernel open/nvidia-drm/

nvidia-drm.ko的内核接口层

kernel-open/nvidia-modeset/

nvidia-modeset.ko的内核接口层

kernel-open/nvidia-uvm/

nvidia-uvm.ko的内核接口层

src目录中为与操作系统无关的代码。

src/nvidia/

nvidia.ko的操作系统无关代码

src/nvidia-modeset/

nvidia-modeset.ko的操作系统无关代码

src/common/

由nvidia.ko和nvidia-modeset.ko中的一个或多个使用的实用程序代码

nouveau/

与nouveau设备驱动程序集成的工具

驱动构成和依赖关系

整个NVIDIA GPU驱动包含了几个子驱动，编译好后的模块文件名分别是nvidia.ko、nvidia-modeset.ko、nvidia-drm.ko、nvidia-uvm.ko和nvidia-peermem.ko。

nvidia.ko是其它所有子驱动的基础，负责创建/proc/driver/nvidia目录、为nvidia-caps和nvidia-nvlink创建cdev、为nvswitch和显卡本身注册PCI驱动等。显卡PCI驱动的probe函数也是在此驱动中实现的，它负责寻找有效的BAR、解析ACPI表获取GPU内存的物理地址、映射寄存器、获取GPU的PCIe配置等。

nvidia-modeset.ko的主要任务是为应用层提供各种配置GPU的ioctl命令，比如SET_MODE、SET_DISP_ATTRIBUTE等。它还提供了一些和内存管理有关的回调函数，比如AllocateVideoMemory，主要是给nvidia-drm驱动使用的。

nvidia-drm.ko调用系统函数为每个GPU注册了一个DRM设备，应该主要和GPU的本职工作“显示”有关。由于我在分析过程中只拿GPU作为计算设备使用，所以没有（至少我没发现）用到这个驱动提供的功能。

nvidia-peermem.ko从名字可以看出来涉及到GPU之间互相访问内存。由于我在分析过程中只用了一个GPU，此驱动根本就没有加载（手动加载也失败了，还没研究原因），所以在此不对它做分析。

nvidia-uvm.ko负责配合cuda库做一些内存管理有关的工作。它和cuda库之间的主要交互方式为ioctl、mmap、poll等传统方式。后文会有一些更详细的介绍。

从驱动注释中得到的一些概念和信息

表面布局（Surface layouts）

BlockLinear是NVIDIA GPU的原生平铺格式，将像素排列成块或平铺，以便在常见的GPU操作中更好地定位。

Pitch是一种简单的“线性”表面布局，像素在内存中逐行顺序排列，为了对齐，可以在每行末尾添加一些填充。

驱动{BANNED}最佳多支持32个GPU

#define NV_MAX_GPUS 32

在cuda进行内存分配时，nvidia-uvm驱动做了哪些工作

我写了一个简单的cuda程序，其中调用了两次cudaMallocManaged（分配统一内存，即Unified Memory，这种内存可以被CPU和GPU共同访问）分配了两个buffer。代码如下。

• 
  

 __global__void add(int n, float *x, float *y) {   for (int i = 0; i < n; i++)   
 y[i] = x[i] + y[i]; } int main(void) {   int N = 1024;   float *x, *y;   
// 内存分配，在GPU或者CPU上统一分配内存  
cudaMallocManaged(&x, N*sizeof(float));   //4KB 
 cudaMallocManaged(&y, N*sizeof(float));   
// initialize x and y arrays on the host   for (int i = 0; i < N; i++) 
{    
x[i] = 1.0f;    y[i] = 2.0f;  
} //执行计算  
add<<<1, 1>>>(N, x, y); // CPU需要等待cuda上的代码运行完毕，才能对数据进行读取 
 cudaDeviceSynchronize(); // Check for errors (all values should be 3.0f)   float maxError = 0.0f;   for (int i = 0; i < N; i++)    
maxError = fmax(maxError, fabs(y[i]-3.0f));   
std::cout << "Max error: " << maxError << std::endl;// Free memory  cudaFree(y); 
 cudaFree(x);   return 0; }