Can a ptx file contain multiple kernels ?
yes, it can
It would be helpfull if the CUDA Driver API had a function/API to return how many kernels a module contains.
It would also be helpfull if it would automatically extract the kernel entry names… and kernel handles.
Especially kernel entry names, since they seemed mangled.
IT would be even more helpfull if the entry names weren’t mangled at all External Image.
Unfortunately, I think we’re stuck with this in the general case as long as C++ function overloading is allowed for global functions. Some kind of mangling is required to capture the type information.
You can prevent mangling if you don’t need overloading by wrapping your global function like this:
extern "C" {
__global__ void mykernel(int *stuff) {
}
// etc etc
}
This is what PyCUDA does to aid in fetching kernels from compiled modules.
Wow that’s pretty cool ! And very handy ! Because the entry names seemed to change after some code changes (?), so this will prevent that which is real nice.
I do wonder though:
Will this still allow c++ classes to be used inside the global ?
I don’t know what extern “C” normally does… it almost seems as if it’s saying: “only c allowed” “no C++” ! External Image =D (But that’s probably just me ! External Image External Image
I have seen this before sometimes… I think it’s actually for exporting dll function names as well, kinda makes sense ! External Image =D
If those dll names would change all the time that would be a real problem for dynamic loading ! External Image =D
extern “C” is only used to indicate to the compiler that you would like C-style linkage, i.e. the symbol name is the function name. You can still use C++ classes, but specifically function overloading won’t work since multiple functions with the same name and different parameters will get the same symbol name in the linker.
Just to add on a bit, these are entirely doable with an ELF library. Just reading through the section header table as well as some entries in .shstrtab and you’ll find all the names (prefixed with .text.). It would be nice, however, to have these functions straight from the driver API.