Indirect calls to statically defined functions from JIT compiled code

Hi,

Does CUDA support indirectly calling statically defined functions in my project, from dynamically generated (and JIT compiled) PTX code?

Here’s an example of what I’m trying to do:

---- foo.cu ----

device foo(float *a)
{

}

int main()
{
void *fooAddr;
cudaGetSymbolFromAddres(&fooAddr, foo);

char *ptx = generatePTX(fooAddr, ...);
CUmodule M;
// Compile the ptx.
cuModuleLoadDataEx(M, ptx, ...);
// Get a pointer to the compiled kernel.
CUFunction F;
cuModuleGetFunction(&F, M, “kernel”);
// Launch the kernel.
cuLaunchKernel(F, ...);

return 0;

}

—/ foo.cu /—

And the generated PTX (generatePTX() above), looks something like this:

---- PTX ----


mov.f64 %reg0, %arr
mov.f64 %reg1, %fooAddr /* value obtained from cudaGetSymbolAddress above. */
Fproto: .callprototype _ (.param .u64);
call %reg1, (%reg0), Fproto

/* indirect call as described in http://docs.nvidia.com/cuda/parallel-thread-execution/#control-flow-instructions-call */

—/ PTX /—