Is there some kind of incompatibility with Boost::thread() and Nvidia CUDA?

Question

I'm developing a generic streaming CUDA kernel execution Framework that allows parallel data copy & execution on the GPU.

Currently I'm calling the cuda kernels within a C++ static function wrapper, so I can call the kernels from a .cpp file (not .cu), like this:

//kernels.cu:

//kernel definition
__global__ void kernelCall_kernel(  dataRow* in,  dataRow* out,  void* additionalData){
    //Do something
};

//kernel handler, so I can compile this .cu and link it with the main project and call it within a .cpp file
extern "C" void kernelCall( dataRow* in,  dataRow* out,  void* additionalData){ 
    int blocksize = 256;  
    dim3 dimBlock(blocksize);
    dim3 dimGrid(ceil(tableSize/(float)blocksize)); 
    kernelCall_kernel<<<dimGrid,dimBlock>>>(in, out, additionalData);   

}

If I call the handler as a normal function, the data printed is right.

//streamProcessing.cpp
//allocations and definitions of data omitted

//copy data to GPU
cudaMemcpy(data_d,data_h,tableSize,cudaMemcpyHostToDevice);
//call:
kernelCall(data_d,result_d,null);
//copy data back
cudaMemcpy(result_h,result_d,resultSize,cudaMemcpyDeviceToHost);
//show result:
printTable(result_h,resultSize);// this just iterate and shows the data

But to allow parallel copy and execution of data on the GPU I need to create a thread, so when I call it making a new boost::thread:

//allocations, definitions of data,copy data to GPU omitted
//call:
boost::thread* kernelThreadOwner = new boost::thread(kernelCall, data_d,result_d,null); 
kernelThreadOwner->join();
//Copy data back and print ommited

I just get garbage when printing the result on the end.

Currently I'm just using one thread, for testing purpose, so there should be no much difference in calling it directly or creating a thread. I have no clue why calling the function directly gives the right result, and when creating a thread not. Is this a problem with CUDA & boost? Am I missing something? Thank you in advise.

Answer 1

The problem is that (pre CUDA 4.0) CUDA contexts are tied to the thread in which they were created. When you are using two threads, you have two contexts. The context that the main thread is allocating and reading from, and the context that the thread which runs the kernel inside are not the same. Memory allocations are not portable between contexts. They are effectively separate memory spaces inside the same GPU.

If you want to use threads in this way, you either need to refactor things so that one thread only "talks" to the GPU, and communicates with the parent via CPU memory, or use the CUDA context migration API, which allows a context to be moved from one thread to another (via cuCtxPushCurrent and cuCtxPopCurrent). Be aware that context migration isn't free, and there is latency involved, so if you plan to migrating contexts around frequently, you might find it more efficient to change to a different design which preserves context-thread affinity.