【NVBit: A Dynamic Binary Instrumentation Framework for NVIDIA GPUs】

NVBit 框架概述

NVBit (NVIDIA Binary Instrumentation Tool) 是一个专门用于 CUDA 编程环境的工具，用于在程序运行时动态插入和修改 CUDA 二进制代码（SASS 机器代码）。这对性能分析、错误检测和调试非常有用。

CUDA 编程和编译流程

1. 编写 CUDA 程序:

   • 用户使用 CUDA 语言编写并行程序，源代码文件通常是以 .cu 结尾。
   • 示例程序：simple_add.cu。

2. 前端编译器 (NVCC):

   • 使用 NVIDIA 的 NVCC 编译器将 CUDA 源代码编译成 PTX (Parallel Thread Execution) 中间代码。PTX 是一种虚拟指令集架构，它提供了一个稳定的编程模型和指令集，适用于通用并行编程。
   • PTX 是设备无关的，这意味着相同的 PTX 代码可以在不同的 GPU 上执行，而不需要修改。

3. 后端编译器:

   • PTX 汇编器 (ptxas): 在程序运行之前（ahead-of-time），PTX 代码被编译成 SASS (Streaming Assembly) 机器代码，SASS 是针对具体 GPU 硬件的低级别指令集。
   • 即时编译器 (JIT): 在程序运行时，CUDA 驱动程序中的 JIT 编译器将 PTX 代码编译成 SASS 机器代码。这种方式在需要灵活性的情况下使用。

NVBit 的工作原理

1. 直接与 CUDA 驱动程序交互:

   • NVBit 直接与 CUDA 驱动程序交互，处理已经编译成 SASS 机器代码的程序。它通过 LD_PRELOAD 机制在运行时注入库，使其能够在加载任何其他库之前加载指定的共享库。

2. 应用程序二进制接口 (ABI):

   • ABI 定义了调用者和被调用者之间的接口属性，例如寄存器的使用、参数传递方式等。NVBit 使用动态汇编器生成符合 ABI 的代码，以便能够将自定义 CUDA 设备函数注入到现有的应用程序中。

3. CUDA API 回调:

   • NVBit 为所有 CUDA 驱动程序 API 提供回调机制。这使得 NVBit 可以在 CUDA 程序调用任何 CUDA API 时截获这些调用，并插入自定义代码或进行分析。
   • NVBit 还提供在应用程序启动和终止时的特定回调功能，可以在程序的整个生命周期中进行监控和干预。

NVBit 工具开发流程

1. 开发 .cu 文件:

   • 使用 NVBit API 编写 CUDA 设备函数和回调函数。例如，在 .cu 文件中定义一个设备函数 incr_counter，用于在每次指令执行时计数。

2. 编译 .cu 文件:

   • 使用 NVCC 编译器将 .cu 文件编译成目标文件。

3. 链接生成共享库:

   • 将编译好的目标文件与静态库 libnvbit.a 链接，生成一个共享库（通常是 .so 文件）。
   • 示例生成的共享库：libmy_nvbit_tool.so。

NVBit 工具的使用

1. 注入共享库:
   • 在运行时通过 LD_PRELOAD 机制将共享库注入到目标应用程序中。
   • 示例命令：LD_PRELOAD=./libmy_nvbit_tool.so ./my_cuda_app。

示例代码详细解释

以下是一个完整的 NVBit 工具示例，用于计算每个线程级指令的执行次数，并在应用程序结束时打印计数器的值。

1 /* NVBit include, any tool must have it */
2 #include "nvbit.h"
3
4 /* Counter variable used to count instructions */
5 __managed__ long counter = 0;
6
7 /* Used to keep track of kernels already instrumented */
8 std::set<CUfunction> instrumented_kernels;
9
10 /* Implementation of instrumentation function */
11 extern "C" __device__ __noinline__ void incr_counter() {
12 atomicAdd(&counter, 1);
13 } NVBIT_EXPORT_DEV_FUNC(incr_counter);
14
15 /* Callback triggered on CUDA driver call */
16 void nvbit_at_cuda_driver_call(CUcontext ctx,
17 int is_exit, cbid_t cbid, const char *name,
18 void *params, CUresult *pStatus) {
19
20 /* Return if not at the entry of a kernel launch */
21 if (cbid != API_CUDA_cuLaunchKernel || is_exit)
22 return;
23
24 /* Get parameters of the kernel launch */
25 cuLaunchKernel_params *p = (cuLaunchKernel_params *) params;
26
27 /* Return if kernel is already instrumented */
28 if(!instrumented_kernels.insert(p->func).second)
29 return;
30
31 /* Instrument all instructions in the kernel */
32 for (auto &i: nvbit_get_instrs(ctx, p->func)) {
33 nvbit_insert_call(i, "incr_counter", IPOINT_BEFORE);
34 }
35 }
36
37 /* Callback triggered on application termination */
38 void nvbit_at_term() {
39 cout << "Total thread instructions " << counter << "\n";
40 }
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40

代码解释

CUDA 驱动程序调用回调:

• 定义一个回调函数 nvbit_at_cuda_driver_call，每次 CUDA 驱动程序调用时触发。
• 在函数入口和退出时分别触发回调，使用 is_exit 标识是否在退出时触发。
• 检查是否在内核启动时触发，如果不是则立即返回。
• 获取内核启动的参数，转换为 cuLaunchKernel_params 类型。
• 检查内核是否已经插桩，如果已经插桩则返回。
• 遍历内核的所有指令，使用 nvbit_get_instrs 获取指令列表，并在每条指令之前插入 incr_counter 调用。

应用程序终止回调:

• 在应用程序终止时触发，打印 counter 变量的值。

NVBit 用户级 API 概述

NVBit 框架提供了五类主要的用户级 API：回调（Callback）、检查（Inspection）、插桩（Instrumentation）、控制（Control）和设备（Device）。

回调 API（Callback API）

回调 API 在目标应用程序遇到特定事件时由 NVBit 核心触发。这些事件包括应用程序的启动或终止，以及任意 CUDA 驱动 API 调用的入口/出口。以下是回调 API 的主要函数：

/* 在应用程序启动/结束时触发 */
void nvbit_at_init();
void nvbit_at_term();

/* 在 CUDA 驱动调用 "name" (例如 cuMemAlloc) 的入口 (is_exit=0) 或出口 (is_exit=1) 触发。cbid 标识 CUDA 驱动调用（与 CUPTI 使用相同的枚举）。params 是指向驱动调用使用的参数结构的指针，需要转换为特定 "cbid" 的正确结构。pStatus 指向 CUDA 驱动调用的返回状态值（仅在出口有效）。 */
void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char* name, void* params, CUresult* pStatus);
1
2
3
4
5
6

这些 API 允许用户在特定事件发生时插入自定义代码，例如在内核启动时，nvbit_at_cuda_driver_call 回调会触发，并提供 CUfunction（即内核）作为参数。NVBit 的回调接口使用与 CUPTI 相同的事件枚举，使得 NVBit 易于使用。

检查 API（Inspection API）

检查 API 允许用户检索和检查组成 CUfunction 的指令。提供了两种视图方式：

1. 平面视图: 将 CUfunction 的指令按程序顺序表示为一个向量。
2. 基本块视图: 将指令表示为向量的向量，每个子向量代表一个基本块。基本块是连续执行的指令序列，没有中断控制流。

以下是检查 API 的主要函数：

/* 获取 CUfunction 的指令 */
const std::vector<Instr*>& nvbit_get_instrs(CUcontext c, CUfunction f);

/* 获取 CUfunction 的基本块 */
const std::vector<std::vector<Instr*>>& nvbit_get_basic_blocks(CUcontext c, CUfunction f);

/* 获取 CUfunction 的相关 CUfunction */
std::vector<CUfunction> nvbit_get_related_funcs(CUcontext c, CUfunction f);
1
2
3
4
5
6
7
8

此外，NVBit 提供了一个 Instr 类，用于抽象实际的机器级 SASS 指令，并通过更高级别的中间表示进行转换。以下是 Instr 类的一些主要方法：

class Instr {
public:
    /* 内存操作类型 */
    enum memOpType { NONE, LOCAL, GENERIC, GLOBAL, SHARED, TEXTURE, CONSTANT };

    /* 操作数类型 */
    enum operandType {
        IMM, // 立即数
        REG, // 寄存器编号
        PRED, // 断言寄存器编号
        CBANK, // 常量库 ID
        SREG, // 特殊寄存器编号
        MREF // 内存引用
    };

    /* 操作数结构 */
    typedef struct {
        operandType type; /* 操作数类型 */
        bool is_neg; /* 是否为负 */
        bool is_abs; /* 是否为绝对值 */
        long val[2]; /* 值 */
    } operand_t;

    /* 返回 SASS 字符串 */
    const char* getSass();

    /* 返回指令在函数中的偏移量（字节） */
    uint32_t getOffset();

    /* 返回指令在函数中的 ID */
    uint32_t getId();

    /* 检查指令是否使用了断言 */
    bool hasPred();

    /* 返回断言编号，仅在 hasPred() 为 true 时有效 */
    int getPredNum();

    /* 检查断言是否为否定（例如 @!P0），仅在 hasPred() 为 true 时有效 */
    bool isPredNeg();

    /* 返回完整的操作码（例如 IMAD.WIDE） */
    const char* getOpcode();

    /* 返回内存操作类型 */
    memOpType getMemOpType();

    /* 检查内存操作是否为加载 */
    bool isLoad();

    /* 检查内存操作是否为存储 */
    bool isStore();

    /* 返回内存操作的字节数 */
    int getMemOpBytes();

    /* 返回操作数的数量 */
    int getNumOperands();

    /* 获取特定操作数 */
    const operand_t* getOperand(int num_operand);

    /* 获取行信息，二进制必须使用生成行信息选项编译 (--generate-line-info/-lineinfo) */
    void getLineInfo(char** file, uint32_t* line);
};
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65

插桩 API（Instrumentation API）

插桩 API 允许用户在 CUfunction 的任意指令之前或之后注入多个设备函数。使用 nvbit_insert_call 插入函数，并指定位置（例如指令之前或之后）和要注入的函数名称。可以通过 nvbit_add_call_arg 添加参数，例如寄存器值、断言值和立即值。

以下是插桩 API 的主要函数：

/* 枚举用于指定插入设备函数的位置（在指令之前或之后） */
typedef enum { IPOINT_BEFORE, IPOINT_AFTER } ipoint_t;

/* 插入名为 "dev_func_name" 的设备函数调用，在指令 "Instr" 之前或之后。设备函数通过名称识别，需要使用宏 NVBIT_EXPORT_DEV_FUNC() 导出 */
void nvbit_insert_call(const Instr* instr, const char* dev_func_name, ipoint_t point);

/* 参数类型 */
typedef enum {
    PRED_VAL, // 指令的断言值
    PRED_REG, // 线程的断言寄存器
    IMM32, // 32 位立即值
    IMM64, // 64 位立即值
    REG_VAL, // 寄存器值
    CBANK_VAL // 常量库值
} arg_t;

/* 向最后插入的调用添加参数 */
void nvbit_add_call_arg(arg_t arg, long val0, long val1);

/* 移除原始指令 */
void nvbit_remove_orig(const Instr* instr);
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21

控制 API（Control API）

控制 API 允许用户在应用程序运行时控制插桩，例如动态选择执行插桩或非插桩版本的函数。用户可以随时重置应用的插桩，以便应用新的插桩选择。

以下是控制 API 的主要函数：

/* 基于标志值运行插桩或原始代码 */
void nvbit_enable_instrumented(CUcontext ctx, CUfunction func, bool flag);

/* 重置函数的插桩，允许重新应用插桩 */
void nvbit_reset_instrumentation(CUcontext ctx, CUfunction func);
1
2
3
4
5

设备 API（Device API）

设备 API 可在插桩（即注入）函数中使用。最重要的是，可以使用此 API 读取和写入应用程序内核或设备函数使用的任意寄存器。尽管任意写入寄存器值可能导致灾难性的应用程序级错误，但修改 GPU 状态的能力对于故障注入或指令仿真等用例是必要的。

以下是设备 API 的主要函数：

/* 读取寄存器值 */
__device__ int nvbit_read_reg32(int reg);
__device__ long nvbit_read_reg64(int reg);

/* 写入寄存器值 */
__device__ void nvbit_write_reg32(int reg, int val);
__device__ void nvbit_write_reg64(int reg, long val);
1
2
3
4
5
6
7

NVBit 核心组件和用户级 API 详细解释

图 3: NVBit 核心组件高层次示意图

图 3 显示了 NVBit 核心的高层次组件，包括驱动拦截器、工具函数加载器、硬件抽象层、指令提升器和代码生成器。下面我们详细描述这些组件的功能和作用。

驱动拦截器（Driver Interposer）

驱动拦截器位于 NVBit 核心层的底部，使用 LD_PRELOAD 提供的函数重载机制拦截 CUDA 驱动 API。当 CUDA 驱动加载应用程序函数（CUfunction）时，驱动拦截器记录其属性，包括：

• 最大寄存器使用量
• 最大堆栈使用量
• 依赖函数（即当前函数可以调用的函数）
• 指令加载的位置

这些属性供 NVBit 核心库的其他组件使用。例如，计算跳转到插桩函数之前需要保存的寄存器数量时会用到最大寄存器消耗。驱动拦截器还负责将 CUDA 驱动回调 API 传播到 NVBit 用户级回调 API。

工具函数加载器（Tool Functions Loader）

工具函数加载器负责加载 NVBit 工具动态库中的所有设备函数。这个过程不会在应用程序启动时自动发生，因为 CUDA 驱动不了解 NVBit 工具库中包含的设备和全局函数。

一些加载的设备函数（使用宏 NVBIT_EXPORT_DEV_FUNCTION 导出的）记录在一个映射中，函数名称与包含函数属性的结构体相关联，例如寄存器使用数量、请求的堆栈大小和代码在 GPU 内存中的位置。代码生成器在创建跳转到插桩函数所需的代码时会使用这些信息。

工具函数加载器还负责加载其他预构建的设备函数（嵌入在 libnvbit.a 中），例如在跳转到用户注入的函数之前用于保存和恢复寄存器的函数。NVBit 实现了一组固定的保存和恢复函数，每个函数针对特定数量的一般用途寄存器。

硬件抽象层（HAL）

硬件抽象层在 CUcontext 在特定设备上启动时初始化。在 HAL 初始化期间，记录设备特定的信息，例如：

• 每条指令的字节大小
• 对齐要求
• 每个线程可用的寄存器数量
• ABI 版本

在一个 GPU 系列中，指令大小是唯一且固定的。Kepler、Maxwell 和 Pascal 具有 64 位宽的编码，而 Volta 具有 128 位宽的编码。ABI 版本指定在进入和退出插桩函数之前必须保存和恢复的寄存器和特殊寄存器（例如在 Volta 中保存收敛障碍状态的寄存器）。HAL 还初始化设备特定的汇编/反汇编函数。这些函数用于在代码生成器中汇编代码或在指令提升器中反汇编代码。使用 HAL 提高了 NVBit 在不同 GPU 代际间的可移植性，因为 SASS ISA 不是固定不变的。

指令提升器（Instruction Lifter）

指令提升器负责检索每个应用程序级 CUfunction 的“原始” SASS 指令缓冲区。当用户请求检查 CUfunction 的指令时（使用 nvbit_get_instrs 或 nvbit_get_basic_blocks），指令提升器将每条指令转换为 Instr 类的对象。Instr 类是机器独立的，表示单个 SASS 指令。反汇编的指令可以排列成一个向量或细分成向量（表示基本块），具体取决于用户的 API 使用情况。

代码生成器（Code Generator）

在 CUDA 驱动回调退出时，如果应用了插桩，代码生成器开始工作。图 4 展示了 NVBit 插桩代码生成的过程。

1. 复制原始代码: 将原始代码复制到系统内存（我们称之为插桩代码）。
2. 生成新代码区域: 在 GPU 内存中分配一个新代码区域，命名为 trampoline。
3. 修改插桩代码: 将插桩代码中高亮指令（如 STS [R15], R8）替换为跳转到 trampoline 的指令（如 JMP L1）。插入 trampoline 优雅地保留了指令布局，而原地扩展会复杂得多，可能需要额外的运行时数据结构。

生成的 trampoline 通常包含以下指令：

1. 保存线程状态: 调用例程保存线程状态，NVBit 仅保存最少数量的一般用途寄存器，并通过分析原始代码和注入函数的寄存器需求来选择适当的保存例程。一般用途寄存器、条件码和断言由这个例程保存到堆栈。
2. 传递参数: 执行一系列指令（本例中只有一条 MOV 指令）传递用户指定的参数。参数传递约定由目标设备的特定 ABI 定义，由 HAL 初始化和处理。
3. 跳转到插桩函数: 跳转到插桩函数 foo 的程序计数器，该函数通过访问工具函数加载器填充的注入函数映射检索。
4. 恢复线程状态: 调用例程从堆栈恢复线程状态。
5. 执行原始指令: 执行“重新定位”的原始指令（STS [R15], R8）。如果这个重新定位的指令是相对控制流指令，必须调整偏移量以考虑新位置和原始目标位置。
6. 返回插桩代码: 跳转回插桩代码的下一个程序计数器（JMP NPC）。

每条插桩指令都有一个 trampoline，但出于效率考虑，这些 trampoline 的空间分配由自定义内存分配器批量处理。trampoline 的内容可能会有所不同，具体取决于在相同 GPU 指令之前或之后插入了多少注入函数，以及注入发生在之前、之后或两者之间。如果使用 nvbit_remove_orig（见前面的插桩 API），“重新定位”的原始指令也必须转换为 NOP。

代码加载器/卸载器（Code Loader/Unloader）

在运行时，用户可以决定是否为特定 CUfunction 启用或禁用插桩。代码加载器/卸载器根据传递给控制 API nvbit_enable_instrumented 的值按需交换原始代码和插桩代码。这个操作的成本与从主机到设备的 cudaMemcpy 操作相同，字节数等于原始代码的大小。为了允许交换，原始代码和插桩代码必须具有相同的字节数，并占用 GPU 内存中的相同位置。只有这样，NVBit 才能保证针对 CUfunction 的绝对跳转在无论运行哪种版本（插桩或非插桩）时继续工作。由于 trampoline 仅在设备内存中创建，因此除非使用控制 API nvbit_reset_instrumented 或卸载特定 CUfunction 的 CUmodule，否则不需要移除它们。代码加载器/卸载器还根据将要执行的代码版本计算内核启动的堆栈和寄存器需求。

当然可以。为了更好地理解 NVBit 的工作原理和各个组件的功能，我们来看一些具体的例子。

例子 1：简单的指令计数器

假设我们有一个简单的 CUDA 内核 simple_add，它只是将两个数组的元素逐一相加。我们想要使用 NVBit 工具来计算每个线程执行的指令数量。

原始 CUDA 内核

__global__ void simple_add(int *a, int *b, int *c) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {
        c[tid] = a[tid] + b[tid];
    }
}
1
2
3
4
5
6

NVBit 插桩工具

我们将编写一个 NVBit 工具来插桩这个内核，以便计算每个线程执行的指令数量。

1. 编写插桩函数：

   我们需要一个设备函数来计数。每次执行指令时，这个函数都会增加一个计数器。

   __managed__ long counter = 0;
   
   extern "C" __device__ __noinline__ void incr_counter() {
       atomicAdd(&counter, 1);
   }
   NVBIT_EXPORT_DEV_FUNC(incr_counter);
   1
   2
   3
   4
   5
   6

2. 编写回调函数：

   我们需要在 CUDA 内核启动时插入这个计数函数。

   void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char* name, void* params, CUresult* pStatus) {
       if (cbid != API_CUDA_cuLaunchKernel || is_exit)
           return;
   
       cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
       if(!instrumented_kernels.insert(p->func).second)
           return;
   
       for (auto &i : nvbit_get_instrs(ctx, p->func)) {
           nvbit_insert_call(i, "incr_counter", IPOINT_BEFORE);
       }
   }
   1
   2
   3
   4
   5
   6
   7
   8
   9
   10
   11
   12

3. 编写初始化和终止回调：

   我们需要在应用程序启动和终止时初始化和打印计数器的值。

   void nvbit_at_init() {
       counter = 0;
   }
   
   void nvbit_at_term() {
       printf("Total thread instructions: %ld\n", counter);
   }
   1
   2
   3
   4
   5
   6
   7

4. 编译和运行：

   将上述代码编译成 NVBit 工具库，并在运行时注入到 CUDA 程序中。

   nvcc -o simple_add simple_add.cu
   LD_PRELOAD=./libnvbit_tool.so ./simple_add
   1
   2

   当程序运行时，simple_add 内核将被插桩，计数器将记录每个线程执行的指令数量，并在程序结束时打印出来。

例子 2：具体解释图 4 中的插桩过程

图 4 显示了 NVBit 插桩代码生成的过程。我们通过一个具体的例子来解释这个过程。

假设我们有以下原始代码片段：

...
SHL R8, R0, 0x1
STS [R15], R8
LDG [R15 + 0x8], R12
...
1
2
3
4
5

我们希望在 STS [R15], R8 指令之前插入一个计数函数 foo。下面是具体的插桩过程：

1. 复制原始代码：

   将原始代码复制到系统内存中，作为插桩代码。

2. 生成 trampoline：

   在 GPU 内存中分配一个新的代码区域，命名为 trampoline。trampoline 是用于保存和恢复状态，并跳转到插桩函数的代码段。

3. 修改插桩代码：

   将 STS [R15], R8 指令替换为跳转到 trampoline 的指令，例如 JMP L1。

4. trampoline 的内容：

   生成的 trampoline 包含以下指令：

   1. 保存线程状态，例如保存寄存器、条件码和断言到堆栈。
   2. 传递参数，例如将参数传递给插桩函数 foo。
   3. 跳转到插桩函数 foo 的程序计数器。
   4. 恢复线程状态，从堆栈恢复保存的寄存器等。
   5. 执行原始的 STS [R15], R8 指令。如果这是相对控制流指令，需要调整偏移量以考虑新位置和原始目标位置。
   6. 跳转回插桩代码的下一个程序计数器。

具体的代码结构如下：

L1:
    JCAL save_thread_state  // 保存线程状态
    MOV32I R4, arg          // 传递参数
    JCAL "foo"              // 跳转到插桩函数
    JCAL restore_thread_state // 恢复线程状态
    STS [R15], R8           // 执行原始指令
    JMP NPC                 // 跳转回插桩代码
1
2
3
4
5
6
7

NVBit 工具使用示例：内存访问地址分歧分析

背景

理解内存访问模式对于优化应用程序或设计内存子系统非常重要。NVBit 允许通过对每个内存操作进行插桩来收集引用地址，然后可以直接在 GPU 上分析这些数据，或者将其发送到 CPU 进行进一步处理。整个缓存模拟器可以围绕这些机制构建。我们来看一个 NVBit 工具示例，该工具计算每个 warp 级全局内存指令请求的唯一缓存行数量。

示例代码解析

插桩函数

__managed__ float uniq_lines = 0;
__managed__ long mem_instrs = 0;

extern "C" __device__ __noinline__ void ifunc(int pred, int r1, int r2, int imm) {
    if (!pred) return;

    long addr = (((long)r1) | ((long)r2 << 32)) + imm;
    int mask = __ballot(1);

    if (get_lane_id() == __ffs(mask) - 1)
        atomicAdd(&mem_instrs, 1);

    long cache_addr = addr >> LOG2_CACHE_LINE_SIZE;
    int cnt = __popc(__match_any_sync(mask, cache_addr));

    atomicAdd(&uniq_lines, 1.0f / cnt);
}
NVBIT_EXPORT_DEV_FUNC(ifunc);
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

这个设备函数 ifunc 用于计数每个 warp 级全局内存指令请求的唯一缓存行数量。函数接受四个参数：一个断言值、两个寄存器值和一个立即数。

1. 断言检查：如果断言值为假，则返回。
2. 地址计算：将两个寄存器值和一个立即数组合成内存地址。
3. 活跃线程掩码计算：计算 warp 中所有活跃线程的掩码。
4. 内存指令计数：只有 warp 中的第一个活跃线程才增加全局内存引用计数器。
5. 缓存行地址计算：每个线程计算它访问的缓存行地址，并使用 CUDA 内置函数将其归约为单个值。
6. 缓存行计数：每个线程按比例增加缓存行计数器。

回调函数

void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char* name, void* params, CUresult* pStatus) {
    if (cbid != API_CUDA_cuLaunchKernel || is_exit)
        return;

    cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
    for (auto &i : nvbit_get_instrs(ctx, p->func)) {
        if (i->getMemOpType() != Instr::GLOBAL) continue;

        for (int n = 0; n < i->getNumOperands(); n++) {
            operand_t *op = i->getOperand(n);
            if (op->type != Instr::MREF) continue;

            nvbit_insert_call(i, "ifunc", IPOINT_BEFORE);
            nvbit_add_call_arg(PRED_VAL);
            nvbit_add_call_arg(REG_VAL, op->val[0]);
            nvbit_add_call_arg(REG_VAL, op->val[0] + 1);
            nvbit_add_call_arg(IMM32, op->val[1]);
        }
    }
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

这个回调函数在 CUDA 内核启动时插入 ifunc 函数：

1. 检查 CUDA 内核启动：如果不是在内核启动时或退出时触发，直接返回。
2. 获取内核指令：使用 nvbit_get_instrs 获取 CUfunction 的指令。
3. 检查全局内存操作：如果指令不是全局内存操作，跳过。
4. 遍历操作数：遍历指令的操作数，找到内存引用（MREF）。
5. 插入 ifunc 函数：在每个全局内存操作指令之前插入 ifunc 函数，并传递四个参数（断言值和三个操作数）。

应用程序终止回调

void nvbit_at_term() {
    printf("Average cache lines requests per memory instruction: %f\n", uniq_lines / mem_instrs);
}
1
2
3

这个回调函数在程序终止时触发，打印每个内存指令请求的平均缓存行数量。

内存访问地址分歧分析

测试结果

图 6 显示了对各种机器学习工作负载（如 AlexNet、ENet、GoogLeNet、ResNet 和 VGG）进行内存访问地址分歧分析的结果。这些工作负载使用了 NVIDIA 开发的预编译库（如 cuBLAS 和 cuDNN）。

• 绿色柱状图：插桩了预编译库。
• 橙色柱状图：未插桩预编译库。

结果表明，未插桩预编译库会导致内存分歧分析的不准确，并显著高估应用程序的内存分歧。因为这些预编译库包含大量不同的内核，而编译器方法无法捕获这些库内的内存引用，导致分析不完整。

通过 NVBit，我们可以在运行时对任何使用这些库的应用程序二进制文件进行插桩，而无需访问库的源代码。这大大简化了分析过程，并提供了更准确的内存访问模式信息。

NVBit 使用示例：内核采样与指令直方图

背景

在进行应用程序优化时，理解指令执行的分布和内核的执行情况是非常重要的。然而，频繁的插桩会导致显著的性能开销。为了减少插桩带来的执行开销，NVBit 允许使用采样技术，仅在特定条件下运行插桩版本的内核。这种方法通过减少插桩回调的频率来降低开销，同时保持数据收集的准确性。

示例代码解析

我们来看一个具体的示例，如何使用 NVBit 实现采样，并构建执行指令的直方图。

插桩函数

我们将实现一个工具，收集所有执行的指令，以构建前五大执行指令的直方图。以下是插桩函数的实现：

__managed__ long instruction_counts[128] = {0}; // 假设共有 128 种不同指令
__managed__ long total_instructions = 0;

extern "C" __device__ __noinline__ void count_instructions(int opcode) {
    atomicAdd(&instruction_counts[opcode], 1);
    atomicAdd(&total_instructions, 1);
}
NVBIT_EXPORT_DEV_FUNC(count_instructions);
1
2
3
4
5
6
7
8

这个设备函数 count_instructions 用于计数每个指令的执行次数。函数接受一个操作码参数 opcode，并将其计数增加。

回调函数

我们在 CUDA 内核启动时插入 count_instructions 函数：

void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char* name, void* params, CUresult* pStatus) {
    if (cbid != API_CUDA_cuLaunchKernel || is_exit)
        return;

    cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
    for (auto &i : nvbit_get_instrs(ctx, p->func)) {
        nvbit_insert_call(i, "count_instructions", IPOINT_BEFORE);
        nvbit_add_call_arg(IMM32, i->getOpcode());
    }
}
1
2
3
4
5
6
7
8
9
10

这个回调函数在 CUDA 内核启动时插入 count_instructions 函数，并传递操作码作为参数。

采样选择逻辑

我们希望仅在每组唯一的网格维度值下运行一次插桩版本的内核。我们可以使用 NVBit 的 nvbit_enable_instrumented API 来实现这个选择逻辑：

std::set<std::tuple<int, int, int>> unique_grids;

void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char* name, void* params, CUresult* pStatus) {
    if (cbid != API_CUDA_cuLaunchKernel || is_exit)
        return;

    cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
    auto grid_dim = std::make_tuple(p->gridDimX, p->gridDimY, p->gridDimZ);

    if (unique_grids.find(grid_dim) == unique_grids.end()) {
        unique_grids.insert(grid_dim);
        nvbit_enable_instrumented(ctx, p->func, true);
    } else {
        nvbit_enable_instrumented(ctx, p->func, false);
    }

    for (auto &i : nvbit_get_instrs(ctx, p->func)) {
        nvbit_insert_call(i, "count_instructions", IPOINT_BEFORE);
        nvbit_add_call_arg(IMM32, i->getOpcode());
    }
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21

这个回调函数在每次 CUDA 内核启动时检查网格维度：

1. 如果这是一个新的网格维度组合，启用插桩并将该组合添加到集合中。
2. 如果这是一个已知的网格维度组合，禁用插桩。

应用程序终止回调

void nvbit_at_term() {
    printf("Top-5 Instructions:\n");
    std::vector<std::pair<long, int>> instruction_counts_vec;
    for (int i = 0; i < 128; i++) {
        instruction_counts_vec.push_back({instruction_counts[i], i});
    }

    std::sort(instruction_counts_vec.rbegin(), instruction_counts_vec.rend());

    for (int i = 0; i < 5; i++) {
        printf("Opcode %d: %ld times\n", instruction_counts_vec[i].second, instruction_counts_vec[i].first);
    }
    printf("Total Instructions: %ld\n", total_instructions);
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14

这个回调函数在程序终止时触发，打印前五大执行指令的统计信息。

测试结果

我们在一系列 OpenACC SpeccAccel 基准测试上运行这个工具，以分析所有执行的指令并构建直方图。

全插桩与采样比较

1. 全插桩方法：插桩所有内核，收集所有数据。这种方法会导致显著的性能开销。
2. 采样方法：仅在每组唯一的网格维度下运行一次插桩版本的内核，其余时间运行未插桩版本的内核。

图 8 显示了全插桩方法和采样方法相对于原生执行的减速：

• 平均而言，全插桩方法比原生执行慢 36.4 倍。
• 采样方法的减速仅为 2.3 倍。

准确性分析

尽管采样方法可以显著减少性能开销，但可能会导致准确性的下降。图 9 显示了采样方法的误差，每个基准测试的误差以单个数字报告，平均误差小于 0.6%。

这种采样技术的误差取决于内核执行的控制流特性。如果内核的控制流仅是网格维度的函数而不依赖于计算值，那么采样误差为 0%。

NVBit 工具使用示例：模拟 Warp 级 FFT 指令

背景

在进行体系结构探索和预硅编译器测试时，指令模拟是一种常见的技术。NVBit 提供了修改可见状态的设备 API，使我们可以模拟不存在的指令。本例中，我们演示如何使用 NVBit 模拟一个假想的 warp 级（32 点）FFT 指令 WFFT32。

示例代码解析

插桩工具（Listing 9）

首先，我们定义一个设备函数 wfft32_emu，用于模拟 WFFT32 指令的功能。然后，我们编写回调函数，在内核启动时将代理指令替换为 wfft32_emu 函数。

/* Compute a 32-point warp-wide FFT (across lanes) */
extern "C" __device__ __noinline__ void wfft32_emu(int reg_dst_num, int reg_src_num) {
    /* Read input register */
    long in = nvbit_read_reg64(reg_src_num);
    /* Implementation of the warp-wide FFT function */
    shuffle_fft_warp(in, out);
    /* Write value in destination registers */
    nvbit_write_reg64(reg_dst_num, out);
}
NVBIT_EXPORT_DEV_FUNC(wfft32_emu);

void nvbit_at_cuda_driver_call(CUcontext ctx, int is_exit, cbid_t cbid, const char *name, void *params, CUresult *pStatus) {
    if (cbid != API_CUDA_cuLaunchKernel || is_exit)
        return;

    cuLaunchKernel_params *p = (cuLaunchKernel_params *)params;
    for (auto &i : nvbit_get_instrs(ctx, p->func)) {
        operand_t *ops = i->get_operands();

        /* Identify "proxy" instruction */
        asm("or.b32 %0, %1, 0xfefefefe;" , ops[2].val[0] == "0xfefefefe" );

        if (i->getOpcode() == "LOP32I.OR" && ops[2].val[0] == "0xfefefefe") {
            nvbit_insert_call(i, "wfft32_emu", IPOINT_BEFORE);
            nvbit_add_call_arg(REG_VAL, ops[0]->val[0]);
            nvbit_add_call_arg(REG_VAL, ops[1]->val[0]);
            /* remove the "proxy" instruction */
            nvbit_remove_orig(i);
        }
    }
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31

详细解释：

1. 设备函数 wfft32_emu：

   • 读取输入寄存器：使用 nvbit_read_reg64 读取输入寄存器的值。
   • FFT 功能实现：调用 shuffle_fft_warp 实现 warp 级 FFT 功能。
   • 写入输出寄存器：使用 nvbit_write_reg64 将结果写入目标寄存器。

2. CUDA 驱动程序回调函数：

   • 检查 CUDA 内核启动：如果不是在内核启动时或出口时触发，直接返回。
   • 获取内核指令：使用 nvbit_get_instrs 获取 CUfunction 的指令。
   • 识别代理指令：如果指令是 LOP32I.OR 并且立即数操作数为 0xfefefefe，则认为是代理指令。
   • 插入模拟函数：在代理指令之前插入 wfft32_emu 函数，并传递源和目标寄存器编号作为参数。
   • 移除代理指令：插入模拟函数后移除代理指令。

内核代码（Listing 10）

在 CUDA 内核中使用代理指令表示假想的 WFFT32 指令。

__global__ void fft32_kernel(float2 *in, float2 *out) {
    /* 获取线程标识符 */
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    /* 插入表示 WFFT32 的代理指令 */
    asm("or.b64 %0, %1, 0xfefefefe;" : "=l"(in[tid]) : "l"(out[tid]));
}
1
2
3
4
5
6

详细解释：

1. 获取线程标识符：

   • int tid = blockIdx.x * blockDim.x + threadIdx.x; 计算当前线程在网格中的全局索引。

2. 插入代理指令：

   • 使用内联汇编插入 or.b64 指令，该指令用作代理指令，表示假想的 WFFT32。0xfefefefe 是用于区分的魔数。

执行分析

当使用 NVBit 工具对上述内核进行插桩时，内联汇编 PTX 指令将被替换为 wfft32_emu 函数。这使得我们可以结合指令模拟和指令跟踪来跟踪不存在的指令集，从而启用基于跟踪的 GPU 模拟器。

示例运行结果

• 执行 fft32_kernel 内核，该内核使用 WFFT32 计算每个 warp 的 32 点 FFT。
• 使用 WFFT32，每个 warp 执行 21 条指令。
• 使用 CUDA 代码执行 warp 级 FFT，每个 warp 执行 150 条指令。

NVBit 的限制与讨论

尽管 NVBit 设计上具有广泛的适应性，允许任意注入任何 CUDA 设备函数，但它也有一些限制。以下是对这些限制的详细解释：

共享内存和常量内存的使用

• 限制原因：注入函数不能使用共享内存和常量内存，因为这些内存可能被应用程序本身使用。在实践中，程序通常会使用所有的共享内存容量，留给插桩库的空间不足。
• 影响：如果插桩工具尝试使用共享内存或常量内存，可能会导致被插桩的程序失败。

插桩函数中的库使用

• 限制原因：插桩函数不能使用目标应用程序使用的加速库。这样做可能会导致插桩递归，即插桩函数本身也被插桩。
• 影响：这会导致无限递归和程序崩溃。

非确定性应用程序

• 限制原因：尽管 NVBit 的设计是尽量减少入侵性，但用户级插桩和 NVBit 框架本身的额外指令、寄存器压力和缓存效应可能会改变包含竞争条件或依赖于特定调度或时间假设的应用程序的行为。
• 影响：如果应用程序已经容易出现非确定性行为，那么使用 NVBit 插桩可能会加剧这种非确定性。例如，在一个依赖内存同步通过自旋循环的应用程序上收集内存地址，可能会导致多次运行应用程序时执行的内存指令数量非常不同。
• 说明：这种限制是所有插桩方法（静态或动态）的常见问题，而不仅仅是 NVBit 特有的问题。

执行开销

• 限制原因：在 NVBit 中，每个活跃线程进入并离开插桩函数，因此所有线程都必须支付这些开销。在插桩函数内，可以基于线程标识符实现线程特化，例如让一部分线程立即返回。
• 计划改进：我们计划在跳转到插桩函数之前实现某种形式的谓词匹配，以允许更精细的线程选择。此外，在跳转到插桩函数之前，每个线程的特定寄存器会被保存到内存中。虽然这是完全并行化的，但需要很多周期，并可能破坏缓存局部性。我们考虑过划出插桩寄存器以限制需要保存的状态，但这会带来新的挑战，包括降低占用率和需要新的特制 ABI。

NVBit 可访问的信息

• 现状：通过 NVBit 检查 API 获取的信息与 NVIDIA 的工具（如 nvdisasm 和 cuda-gdb）可以观察到的信息相当。例如，开发人员可以使用 nvdisasm 观察任何 GPU 二进制文件的 SASS 代码（如果 SASS 存在），并使用 cuda-gdb 观察任何嵌入的 PTX 代码到 SASS 的翻译和映射。此外，用户可以使用 cuda-gdb 从内存和寄存器中读取值，允许手动检查整个 ISA 可见的机器状态（如同 NVBit）。
• 优势：NVBit 相比这些工具的主要优势在于，可以用高性能 C/C++ 代码在运行时分析这些信息，比用户用 nvdisasm 和 cuda-gdb 交互操作快了几个数量级。此外，NVBit 的动态插桩特性，允许用户在运行的应用程序中注入通用 CUDA 函数，这是目前其他任何工具都无法做到的。

动态插桩

• 现状：NVBit 允许在内核启动之前进行插桩。然而，一旦内核开始执行，代码就不能再进一步修改（直到下一次启动）。这与 CPU 上的现有方法形成对比，CPU 方法可以随时中断并动态修改代码。
• 原因：这是由于 GPU 不能自编译代码，必须依赖 CPU 来驱动执行。这并不是 NVBit 的限制，而是 GPU 本身的限制。