使用 NVBit 进行内存访问跟踪指南
使用 NVBit 进行内存访问跟踪指南
NVBit（NVIDIA Binary Instrumentation Tool） 是 NVIDIA 提供的一款轻量级、灵活的 GPU 二进制插桩框架，可以帮助开发者在不修改源代码的情况下，跟踪 CUDA 程序的内存访问。
具有以下特点：
- 轻量级和高效：对应用程序的性能影响较小。
- 灵活性：支持用户自定义插件，实现多种分析功能。
- 透明性：无需修改应用程序源码或重新编译。
- 适用性：支持 Maxwell（Compute Capability 5.0）及以上架构的 GPU。
NVBit 通过在程序运行时对 CUDA 内核的二进制代码（SASS）进行动态插桩，实现对内核执行的监控和分析。
NVBit 的工作原理
- 直接与 CUDA 驱动程序交互:
NVBit 直接与 CUDA 驱动程序交互，处理已经编译成 SASS 机器代码的程序。它通过 LD_PRELOAD 机制在运行时注入库，使其能够在加载任何其他库之前加载指定的共享库。 - 应用程序二进制接口 (ABI):
ABI 定义了调用者和被调用者之间的接口属性，例如寄存器的使用、参数传递方式等。NVBit 使用动态汇编器生成符合 ABI 的代码，以便能够将自定义 CUDA 设备函数注入到现有的应用程序中。 - CUDA API 回调:
NVBit 为所有 CUDA 驱动程序 API 提供回调机制。这使得 NVBit 可以在 CUDA 程序调用任何 CUDA API 时截获这些调用，并插入自定义代码或进行分析。
NVBit 还提供在应用程序启动和终止时的特定回调功能，可以在程序的整个生命周期中进行监控和干预。
NVBit 工具开发流程
- 开发 .cu 文件:
使用 NVBit API 编写 CUDA 设备函数和回调函数。例如，在 .cu 文件中定义一个设备函数 incr_counter，用于在每次指令执行时计数。 - 编译 .cu 文件:
使用 NVCC 编译器将 .cu 文件编译成目标文件。
3. 链接生成共享库:
将编译好的目标文件与静态库 libnvbit.a 链接，生成一个共享库（通常是 .so 文件）。
示例生成的共享库：libmy_nvbit_tool.so。
https://blog.csdn.net/m0_63471305/article/details/139804027
1.下载地址：
https://github.com/NVlabs/NVBit/releases
解压 .bz2：
tar -xvjf ./nvbit-Linux-x86_64-1.7.1.tar.bz2
如果直接git clone https://github.com/NVlabs/NVBit.git会只有两个txt、md两个文件
或者git clone https://github.com/CoffeeBeforeArch/nvbit_tools.git —>版本较旧
2.进入 NVBit 目录并设置环境变量：
cd nvbit-release
export NVBIT_ROOT=$(pwd)
export LD_LIBRARY_PATH=$NVBIT_ROOT:$LD_LIBRARY_PATH
3.进入tools里会有mem_trace目录：
mem_trace.cu是一个插槽功能的函数
编译：
make ARCH=sm_80 #这里Makefile里是`ARCH?=ALL` 手动指定架构
编译成功后，将生成 mem_trace.so 插件文件，成为共享库。
4.编译cuda程序：
假设是矩阵相乘 matrixmul.cu :
nvcc -o matrixmul matrixmul.cu
#include <stdio.h>
#include <cuda_runtime.h>
// 矩阵维度
#define N 1024 // 矩阵大小 N x N
// CUDA内核，执行矩阵乘法
__global__ void matrixMulKernel(float* A, float* B, float* C, int n) {
// 获取当前线程对应的行和列索引
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float value = 0.0;
// 矩阵乘法的核心计算：C(row, col) = A(row, :) * B(:, col)
if (row < n && col < n) {
for (int k = 0; k < n; ++k) {
value += A[row * n + k] * B[k * n + col];
}
C[row * n + col] = value;
}
}
int main() {
int size = N * N * sizeof(float);
float *h_A, *h_B, *h_C; // 主机内存
float *d_A, *d_B, *d_C; // 设备内存
// 分配主机内存
h_A = (float*)malloc(size);
h_B = (float*)malloc(size);
h_C = (float*)malloc(size);
// 初始化矩阵A和B
for (int i = 0; i < N * N; i++) {
h_A[i] = 1.0f; // 简单起见，将A和B初始化为全1矩阵
h_B[i] = 1.0f;
}
// 分配设备内存
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);
// 将主机内存拷贝到设备内存
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 定义线程块和网格维度
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid((N + threadsPerBlock.x - 1) / threadsPerBlock.x,
(N + threadsPerBlock.y - 1) / threadsPerBlock.y);
// 调用CUDA内核
matrixMulKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// 将结果从设备内存拷贝到主机内存
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 验证结果
bool success = true;
for (int i = 0; i < N * N; i++) {
if (h_C[i] != N) {
success = false;
break;
}
}
if (success) {
printf("矩阵相乘成功!\n");
} else {
printf("矩阵相乘失败!\n");
}
// 释放内存
free(h_A);
free(h_B);
free(h_C);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
return 0;
}
5. 注入共享库
在运行时通过 LD_PRELOAD 机制将共享库注入到目标应用程序中。
LD_PRELOAD=nvbit-release/tools/mem_trace.so ./matrixmul