Compile kernel to DLL and use it

Question

I managed to make whole thing with function pointers work and now I want to dynamically load such a kernel. My code:

cuh:

ifndef customkernel_cuh
define customkernel_cuh

extern "C" pfunctionWhere __declspec(dllexport) getHostPointer();

endif

cu:

__device__
    bool myWhere2(PapayaColumnValue *values)
{
    return ((int)values[1]) == 1 || ((int)values[1]) == 3;
}
__device__ pfunctionWhere pMyWhere2 = myWhere2;

pfunctionWhere __declspec(dllexport) getHostPointer()
{
    cudaError_t cudaStatus;
    pfunctionWhere h_pMyWhere2;
    cudaStatus = cudaMemcpyFromSymbol(&h_pMyWhere2, pMyWhere2, sizeof(pfunctionWhere));
    cudaDeviceSynchronize();
    return h_pMyWhere2;
}

main.cpp:

HINSTANCE hGetProcIDDLL = LoadLibrary("xxx.dll");
    if (hGetProcIDDLL == NULL) {
        std::cout << "could not load the dynamic library" << std::endl;
    }
    dll_func dll_getHostPointer = (dll_func)GetProcAddress(hGetProcIDDLL, "getHostPointer");
    DWORD dw = GetLastError(); 
    if (!dll_getHostPointer) {
        std::cout << "could not locate the function" << std::endl;
    }
    pfunctionWhere h_pMyWhere2 = (*dll_getHostPointer)();

And if I debug into dll cudaStatus = cudaSuccess, but pointer to function is null and it is returned from dll invocation. My question is: is it possible to write kernel functions in DLL and then get pointer to such kernels and pass it to main program? I need it to be able to change the kernel while main program is working.

Ivan Solntsev · Answer 1 · 2014-12-02T14:13:08.223

1

You could compile your kernel code to PTX and run it using CUDA driver API, see CUDA C Programming Guide / Driver Api / Module.

If you invoke nvcc with -ptx option instead of --compile, it will generate ptx files. It is not linked with your exe program, and you can change ptx files at any time.

edited Dec 02 '14 at 14:13

answered Jun 11 '13 at 15:07

Ivan Solntsev

2,081
2
31
40

I am trying to load some code that can be compiled after compiling main program and then it could be changed and reload. Would it be possible with PTX? – user2390724 Jun 12 '13 at 09:42
Yes, you compiele ptx files separately from cpu program. – Ivan Solntsev Jun 12 '13 at 09:59

score 0 · Answer 2 · answered Jun 11 '13 at 11:17

0

The whole code doesn't make sense.

First, you are not checking the cudaStatus.

Second you are copying from constant memory, but why? surely you didn't update the constant memory in your kernel. You are probably looking for cudaMemcpy not cudaMemcpyFromSymbol

Have a Google on "Pinned Memory", it might be useful in your case.

answered Jun 11 '13 at 11:17

Adam

3,872
6
36
66

this is based on first answer: http://stackoverflow.com/questions/15644261/cuda-function-pointers I am checking cudaStatus via debugger. Copying from constant memory via cudaMemcpyFromSymbol is a way for passing pointer to device function to kernel. You have it in a link above – user2390724 Jun 12 '13 at 09:39

Compile kernel to DLL and use it

2 Answers2