我如何以编程方式检查在一个隔间中有哪些目标可用?



假设我有一个库文件,或者为了方便起见,我把一个库文件加载到内存中(这样我就有了数据的void*)。

使用模块的CUDA驱动程序API,我可以尝试将数据加载到当前上下文中的模块中;如果编译后的代码不可用于相关目标(并且没有可以被jit代替的PTX),则此操作将失败。但是-我实际上想做的是检查哪些目标有模块数据(或模块文件)中的代码?

非编程,我知道我可以调用:

cuobjdump my.fatbin

并获得其中内容的列表。但是我想在我的应用程序的代码中做到这一点。

这可能吗?

您可以在程序中调用cuobjdump并解析其输出。

#include <cstdlib>
#include <string>
__global__
void kernel(){
}
int main(int argc, char** argv){
std::string command{};
command += "cuobjdump ";
command += argv[0];
command += " > out.txt";
int sysret = system(command.c_str());
kernel<<<1,1>>>();
cudaDeviceSynchronize();
return sysret;
}

您可以使用ELF解析器来完成此操作。

看起来cubin文件实际上是稍微不标准的ELF文件。具体来说,它们有一个.nv_fatbinELF节,其中包含针对不同目标编译的代码区域;

如果您使用了一个ELF库,并让它接受一些无效的/不同的魔术数字/版本号,它可能会以一种您可以轻松从中提取元数据的方式解析库文件,包括每个区域的目标体系结构。

请参阅如何使用ELF库生成(但请注意,这是带有许多幻数的Python代码,有些难以理解)。

现在只剩下"简单的编码问题"了,调整ELF解析器来收集和公开您需要的内容。剩下的问题是,是否所有相关信息都在ELF元数据中,或者是否还需要对内容进行进一步解析。

最新更新