How can I associate my NVRTC program source with a file?

Question

I'm using NVRTC to compile a kernel. The relevant API call is:

nvrtcResult nvrtcCreateProgram ( 
    nvrtcProgram* prog,
    const char* src,
    const char* name, 
    int  numHeaders, 
    const char** headers,
    const char** includeNames )

As you can see, the source is a raw string, and not associated with a file. That means that when you --generate-line-info, you get line numbers, but no related filename. And that means that if you then use, say, NSight Compute - you won't be able to see your kernel source code.

Obviously, neither NSight Compute itself, nor NVRTC itself, can figure out that the raw source is mirrored in some file. But there has to be some way to get around this:

Perhaps I'm missing something in the NVRTC API which can make the source <-> file association?
Perhaps we can manipulate the resulting compiled program (reasonably, not manually, or write-my-own-new-API) to make the association?
Perhaps we can shove the source code into the compiled program somehow?

einpoklum · Answer 1 · 2021-03-10T11:09:11.513

Here's my initial workaround:

Place your source in a file, say my_kernel.cuh.
Create the string:
```
#include "my_kernel.cuh"
```
Compile just this string using NVRTC

Now, NVRTC is able to associate included files' sources with the files, so it's only a stub that will be missing in terms of source<->file association.

Caveat: You will need to be careful about paths - NVRTC's include paths, the working directory from which you invoke your program vs the directory of the source file etc.

einpoklum · Accepted Answer · 2022-02-12T10:14:51.010

It seems NVRTC does provides a default filename, such that if you place your source in the file with that name - NSight Compute may be able to find it.

The name is the one you passed to nvrtcCreateProgram() as the name argument.

So, if your kernel function (i.e. your __global__ function) is in my_kernel.cuh, and you place this file in the working directory of the profile program (which you tell NSight Compute about), or in one of the include directories you built your program with, you'll be able to read your source. If the original file's own directory is also one of the include directories, then you're in luck and you don't even have to make a copy.

How can I associate my NVRTC program source with a file?

2 Answers2