Cannot find the instrumentation function

[francis0407]

I'm writing a simple program to test the functionality of NVBit.

Here is a simplified version of my code :

extern "C" __device__ __noinline__ void my_function() {
    // do something here
}

int main() {
    // load a cubin file which contains the function that need to be instrumented.
    // ...
    CUfunction myFunc;
    cuModuleGetFunction(&myFunc, myModule, myFuncName);
    auto instrs = nvbit_get_instrs(cuContext, func);
    nvbit_insert_call(instrs[18], "my_function", IPOINT_BEFORE);
    nvbit_enable_instrumented(cuContext, func, true);
    // ...
}

The error is :

ASSERT FAIL: function.cpp:764:void Function::gen_new_code(std::unordered_map<std::__cxx11::basic_string<char>, Function*>&): FAIL !(instr_func_map.find(c.instr_func_name) != instr_func_map.end()) MSG: instrumentation function merged_kernel0 not found in binary!

I don't understand how nvbit_insert_call looks for the instrumentation function.

[x-y-z]

Have you checked the examples like instr_count in the tools folder?

Basically, you need to write your tool, compile it as a dynamic library .so file, and run LD_PRELOAD=<tool>.so <your program> to perform the binary instrumentation.

[francis0407]

@x-y-z Thanks for your reply.
I understand the example instr_count, but I want to use NVBit for other purposes, such as modify the SASS of a kernel.

I don't want to use it as a shared library. So my question is what kind of device function can be found by nvbit_insert_call ? Must they be loaded as a .so file?

[x-y-z]

To modify the SASS of a kernel, you can refer to mov_replace as an example.

The current implementation of NVBit only allows using it as a shared library. Device functions used in your binary could be found by NVBit.

From the error above, your instrumentation function merged_kernel0 was not found. You should put that function in inject_functs.cu and compile the file as part of NVBit tool.