赞
踩
NVBit (NVIDIA Binary Instrumentation Tool) 是一个专门用于 CUDA 编程环境的工具,用于在程序运行时动态插入和修改 CUDA 二进制代码(SASS 机器代码)。这对性能分析、错误检测和调试非常有用。
编写 CUDA 程序:
.cu
结尾。simple_add.cu
。前端编译器 (NVCC):
后端编译器:
直接与 CUDA 驱动程序交互:
LD_PRELOAD
机制在运行时注入库,使其能够在加载任何其他库之前加载指定的共享库。应用程序二进制接口 (ABI):
CUDA API 回调:
开发 .cu 文件:
.cu
文件中定义一个设备函数 incr_counter
,用于在每次指令执行时计数。编译 .cu 文件:
.cu
文件编译成目标文件。链接生成共享库:
libnvbit.a
链接,生成一个共享库(通常是 .so
文件)。libmy_nvbit_tool.so
。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 }
CUDA 驱动程序调用回调:
nvbit_at_cuda_driver_call
,每次 CUDA 驱动程序调用时触发。is_exit
标识是否在退出时触发。cuLaunchKernel_params
类型。nvbit_get_instrs
获取指令列表,并在每条指令之前插入 incr_counter
调用。应用程序终止回调:
counter
变量的值。NVBit 框架提供了五类主要的用户级 API:回调(Callback)、检查(Inspection)、插桩(Instrumentation)、控制(Control)和设备(Device)。
回调 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);
这些 API 允许用户在特定事件发生时插入自定义代码,例如在内核启动时,nvbit_at_cuda_driver_call
回调会触发,并提供 CUfunction(即内核)作为参数。NVBit 的回调接口使用与 CUPTI 相同的事件枚举,使得 NVBit 易于使用。
检查 API 允许用户检索和检查组成 CUfunction 的指令。提供了两种视图方式:
以下是检查 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);
此外,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); };
插桩 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);
控制 API 允许用户在应用程序运行时控制插桩,例如动态选择执行插桩或非插桩版本的函数。用户可以随时重置应用的插桩,以便应用新的插桩选择。
以下是控制 API 的主要函数:
/* 基于标志值运行插桩或原始代码 */
void nvbit_enable_instrumented(CUcontext ctx, CUfunction func, bool flag);
/* 重置函数的插桩,允许重新应用插桩 */
void nvbit_reset_instrumentation(CUcontext ctx, CUfunction func);
设备 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);
图 3 显示了 NVBit 核心的高层次组件,包括驱动拦截器、工具函数加载器、硬件抽象层、指令提升器和代码生成器。下面我们详细描述这些组件的功能和作用。
驱动拦截器位于 NVBit 核心层的底部,使用 LD_PRELOAD
提供的函数重载机制拦截 CUDA 驱动 API。当 CUDA 驱动加载应用程序函数(CUfunction)时,驱动拦截器记录其属性,包括:
这些属性供 NVBit 核心库的其他组件使用。例如,计算跳转到插桩函数之前需要保存的寄存器数量时会用到最大寄存器消耗。驱动拦截器还负责将 CUDA 驱动回调 API 传播到 NVBit 用户级回调 API。
工具函数加载器负责加载 NVBit 工具动态库中的所有设备函数。这个过程不会在应用程序启动时自动发生,因为 CUDA 驱动不了解 NVBit 工具库中包含的设备和全局函数。
一些加载的设备函数(使用宏 NVBIT_EXPORT_DEV_FUNCTION
导出的)记录在一个映射中,函数名称与包含函数属性的结构体相关联,例如寄存器使用数量、请求的堆栈大小和代码在 GPU 内存中的位置。代码生成器在创建跳转到插桩函数所需的代码时会使用这些信息。
工具函数加载器还负责加载其他预构建的设备函数(嵌入在 libnvbit.a
中),例如在跳转到用户注入的函数之前用于保存和恢复寄存器的函数。NVBit 实现了一组固定的保存和恢复函数,每个函数针对特定数量的一般用途寄存器。
硬件抽象层在 CUcontext 在特定设备上启动时初始化。在 HAL 初始化期间,记录设备特定的信息,例如:
在一个 GPU 系列中,指令大小是唯一且固定的。Kepler、Maxwell 和 Pascal 具有 64 位宽的编码,而 Volta 具有 128 位宽的编码。ABI 版本指定在进入和退出插桩函数之前必须保存和恢复的寄存器和特殊寄存器(例如在 Volta 中保存收敛障碍状态的寄存器)。HAL 还初始化设备特定的汇编/反汇编函数。这些函数用于在代码生成器中汇编代码或在指令提升器中反汇编代码。使用 HAL 提高了 NVBit 在不同 GPU 代际间的可移植性,因为 SASS ISA 不是固定不变的。
指令提升器负责检索每个应用程序级 CUfunction 的“原始” SASS 指令缓冲区。当用户请求检查 CUfunction 的指令时(使用 nvbit_get_instrs
或 nvbit_get_basic_blocks
),指令提升器将每条指令转换为 Instr
类的对象。Instr 类是机器独立的,表示单个 SASS 指令。反汇编的指令可以排列成一个向量或细分成向量(表示基本块),具体取决于用户的 API 使用情况。
在 CUDA 驱动回调退出时,如果应用了插桩,代码生成器开始工作。图 4 展示了 NVBit 插桩代码生成的过程。
生成的 trampoline 通常包含以下指令:
每条插桩指令都有一个 trampoline,但出于效率考虑,这些 trampoline 的空间分配由自定义内存分配器批量处理。trampoline 的内容可能会有所不同,具体取决于在相同 GPU 指令之前或之后插入了多少注入函数,以及注入发生在之前、之后或两者之间。如果使用 nvbit_remove_orig
(见前面的插桩 API),“重新定位”的原始指令也必须转换为 NOP。
在运行时,用户可以决定是否为特定 CUfunction 启用或禁用插桩。代码加载器/卸载器根据传递给控制 API nvbit_enable_instrumented
的值按需交换原始代码和插桩代码。这个操作的成本与从主机到设备的 cudaMemcpy 操作相同,字节数等于原始代码的大小。为了允许交换,原始代码和插桩代码必须具有相同的字节数,并占用 GPU 内存中的相同位置。只有这样,NVBit 才能保证针对 CUfunction 的绝对跳转在无论运行哪种版本(插桩或非插桩)时继续工作。由于 trampoline 仅在设备内存中创建,因此除非使用控制 API nvbit_reset_instrumented
或卸载特定 CUfunction 的 CUmodule,否则不需要移除它们。代码加载器/卸载器还根据将要执行的代码版本计算内核启动的堆栈和寄存器需求。
当然可以。为了更好地理解 NVBit 的工作原理和各个组件的功能,我们来看一些具体的例子。
假设我们有一个简单的 CUDA 内核 simple_add
,它只是将两个数组的元素逐一相加。我们想要使用 NVBit 工具来计算每个线程执行的指令数量。
__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];
}
}
我们将编写一个 NVBit 工具来插桩这个内核,以便计算每个线程执行的指令数量。
编写插桩函数:
我们需要一个设备函数来计数。每次执行指令时,这个函数都会增加一个计数器。
__managed__ long counter = 0;
extern "C" __device__ __noinline__ void incr_counter() {
atomicAdd(&counter, 1);
}
NVBIT_EXPORT_DEV_FUNC(incr_counter);
编写回调函数:
我们需要在 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);
}
}
编写初始化和终止回调:
我们需要在应用程序启动和终止时初始化和打印计数器的值。
void nvbit_at_init() {
counter = 0;
}
void nvbit_at_term() {
printf("Total thread instructions: %ld\n", counter);
}
编译和运行:
将上述代码编译成 NVBit 工具库,并在运行时注入到 CUDA 程序中。
nvcc -o simple_add simple_add.cu
LD_PRELOAD=./libnvbit_tool.so ./simple_add
当程序运行时,simple_add
内核将被插桩,计数器将记录每个线程执行的指令数量,并在程序结束时打印出来。
图 4 显示了 NVBit 插桩代码生成的过程。我们通过一个具体的例子来解释这个过程。
假设我们有以下原始代码片段:
...
SHL R8, R0, 0x1
STS [R15], R8
LDG [R15 + 0x8], R12
...
我们希望在 STS [R15], R8
指令之前插入一个计数函数 foo
。下面是具体的插桩过程:
复制原始代码:
将原始代码复制到系统内存中,作为插桩代码。
生成 trampoline:
在 GPU 内存中分配一个新的代码区域,命名为 trampoline。trampoline 是用于保存和恢复状态,并跳转到插桩函数的代码段。
修改插桩代码:
将 STS [R15], R8
指令替换为跳转到 trampoline 的指令,例如 JMP L1
。
trampoline 的内容:
生成的 trampoline 包含以下指令:
foo
。foo
的程序计数器。STS [R15], R8
指令。如果这是相对控制流指令,需要调整偏移量以考虑新位置和原始目标位置。具体的代码结构如下:
L1:
JCAL save_thread_state // 保存线程状态
MOV32I R4, arg // 传递参数
JCAL "foo" // 跳转到插桩函数
JCAL restore_thread_state // 恢复线程状态
STS [R15], R8 // 执行原始指令
JMP NPC // 跳转回插桩代码
理解内存访问模式对于优化应用程序或设计内存子系统非常重要。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);
这个设备函数 ifunc
用于计数每个 warp 级全局内存指令请求的唯一缓存行数量。函数接受四个参数:一个断言值、两个寄存器值和一个立即数。
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]); } } }
这个回调函数在 CUDA 内核启动时插入 ifunc
函数:
nvbit_get_instrs
获取 CUfunction 的指令。ifunc
函数:在每个全局内存操作指令之前插入 ifunc
函数,并传递四个参数(断言值和三个操作数)。void nvbit_at_term() {
printf("Average cache lines requests per memory instruction: %f\n", uniq_lines / mem_instrs);
}
这个回调函数在程序终止时触发,打印每个内存指令请求的平均缓存行数量。
图 6 显示了对各种机器学习工作负载(如 AlexNet、ENet、GoogLeNet、ResNet 和 VGG)进行内存访问地址分歧分析的结果。这些工作负载使用了 NVIDIA 开发的预编译库(如 cuBLAS 和 cuDNN)。
结果表明,未插桩预编译库会导致内存分歧分析的不准确,并显著高估应用程序的内存分歧。因为这些预编译库包含大量不同的内核,而编译器方法无法捕获这些库内的内存引用,导致分析不完整。
通过 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);
这个设备函数 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());
}
}
这个回调函数在 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()); } }
这个回调函数在每次 CUDA 内核启动时检查网格维度:
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);
}
这个回调函数在程序终止时触发,打印前五大执行指令的统计信息。
我们在一系列 OpenACC SpeccAccel 基准测试上运行这个工具,以分析所有执行的指令并构建直方图。
图 8 显示了全插桩方法和采样方法相对于原生执行的减速:
尽管采样方法可以显著减少性能开销,但可能会导致准确性的下降。图 9 显示了采样方法的误差,每个基准测试的误差以单个数字报告,平均误差小于 0.6%。
这种采样技术的误差取决于内核执行的控制流特性。如果内核的控制流仅是网格维度的函数而不依赖于计算值,那么采样误差为 0%。
在进行体系结构探索和预硅编译器测试时,指令模拟是一种常见的技术。NVBit 提供了修改可见状态的设备 API,使我们可以模拟不存在的指令。本例中,我们演示如何使用 NVBit 模拟一个假想的 warp 级(32 点)FFT 指令 WFFT32。
首先,我们定义一个设备函数 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); } } }
设备函数 wfft32_emu
:
nvbit_read_reg64
读取输入寄存器的值。shuffle_fft_warp
实现 warp 级 FFT 功能。nvbit_write_reg64
将结果写入目标寄存器。CUDA 驱动程序回调函数:
nvbit_get_instrs
获取 CUfunction 的指令。LOP32I.OR
并且立即数操作数为 0xfefefefe
,则认为是代理指令。wfft32_emu
函数,并传递源和目标寄存器编号作为参数。在 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]));
}
获取线程标识符:
int tid = blockIdx.x * blockDim.x + threadIdx.x;
计算当前线程在网格中的全局索引。插入代理指令:
or.b64
指令,该指令用作代理指令,表示假想的 WFFT32。0xfefefefe
是用于区分的魔数。当使用 NVBit 工具对上述内核进行插桩时,内联汇编 PTX 指令将被替换为 wfft32_emu
函数。这使得我们可以结合指令模拟和指令跟踪来跟踪不存在的指令集,从而启用基于跟踪的 GPU 模拟器。
fft32_kernel
内核,该内核使用 WFFT32 计算每个 warp 的 32 点 FFT。尽管 NVBit 设计上具有广泛的适应性,允许任意注入任何 CUDA 设备函数,但它也有一些限制。以下是对这些限制的详细解释:
Copyright © 2003-2013 www.wpsshop.cn 版权所有,并保留所有权利。