Kernel Configuration

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

to avoid using little executing functions like

void Launchkernel(int gs, int bs, params...)

{

kernelcall<<<gs,bs>>>(params...);

}

I tried to use cudaConfigureCall, which doesn’t return an error. When i start my kernel like

cudaConfigureCall(gs,bs,sm,stream);

kernelcall(params...);

the cudaGetLastError function returns cudaInvalidDeviceConfig …

this error code actually should have been returned by cudaConfigureCall. Also, all parameters are ok.

Is it actually possible to use cudaConfigureCall like this, or do i need to use multiple cudaSetupArgument + cudaLaunch ?

2

Yes.

I’m surprised the compiler doesn’t catch this. Then again it is like using a device (function) pointer on the host, which isn’t caught either.

3

It’s not only that the compiler doesn’t catch it - there isn’t actually an exception thrown. Though i tried to do it the way you proposed - I get struck on “cudaInvalidDeviceFunction”, probably some sort of c++ mangling, but i really don’t figure out how to solve it.

Way of least resistance - executing shell function.

4

Oh - I didn’t mean “catch” in the C++ sense. I just wanted to say that the compiler could print an error message.

Probably. You did notice that cudaLaunch takes a char* with the function name, and not a function pointer?

??? You aren’t saying here that, in order to avoid a C++ function call, you are starting a shell instead, are you?

5

shell - not the thing you may consider as a shell…bad term. just an additional, superfluous function in the .cu file that does the <<<gs,bs>>> stuff, taking everything as a parameter. Currently i tried to avoid packing anything else than the kernels into the .cu file - looked some kind of clean.
char * - yes, i know… cudaLaunch(“kernelcall”);

6

Yes, I put most of my CUDA stuff into wrapper functions as well, as some of it’s C++ magic doesn’t go well together with other libraries.

The simplest way to figure out the mangled function names is to compile with [font=“Courier New”]-Xptxas=-v[/font], which will print them together with the memory use statistics for each kernel.

7

I don’t know this for sure, but in case it’s possible that the name extension (the character gobbledegoob named mangling) changes after some updates of nvcc, i stay away from hardcoded cudaLaunch(const char*) calls, since i don’t want to be responsible for the colleagues that are eating the carpet after spilling their coffee because of a sudden rage.

8

OK. Another option would be to turn off name mangling altogether by declaring the link style as [font=“Courier New”]extern “C”[/font].