-1

My work uses 4 streams and I wish they can be run concurrently. The code is something like this:

for (int i=0; i<N; i++)   //batch numbers
{
    for (int j=0; j<4; j++)
       myCudaCode(stream[j]);    // working codes using the specified stream
}

However, from nvvp profiler I see the streams are actually not concurrently running because the CPU are fully occupied by the kernel launches. I did not use any cudaDeviceSynchronize. You may see the figure from the following link.

enter image description here

I understand that all my kernels on GPU are pretty small, comparable to the kernel launch times on CPU. But so far we do not intend to change them. From the above figure, I see most kernel launches on CPU take around 5~10 us, which is considered normal. The whole processing time for one batch is around 0.4 ms (as shown in the gray)

An intuitive thinking to optimize the code is to use multi-threading to parallelize the CUDA kernel launches on CPU. Here is what I did by use of openMP:

for (int i=0; i<N; i++)
{
    #pragma omp parallel num_threads(4)
       myCudaCode(stream[omp_get_thread_num()]);
}

Now the nvvp profiler shows like this:

enter image description here

The four streams are seemingly running concurrently. However, for each CPU thread, the CUDA kernel launches are not no longer as compact as before, and also are significantly stretched (typically 20~30 us). The resulting times required for one batch processing (shown in gray) are now around 0.5ms, even longer than the single thread case.

I also tried pthread method. It shows the similar problem.

So I'd like to ask for an effective way to parallize the kernel launches on CPU. Ideally, the times are expected be reduced by one fourth.

I'm pretty sure every kernel is small enough, far away from the full GPU computing resources. And I'm using Linux, i7 8 cores CPU and GTX 1070 GPU.

Update: based on my experiments, it seems the use of multi CPU threads do not reduce the total kernel launch times at all. Suppose the single thread code for processing N streams requires time T, then using openMP and the N-thread code for processing N streams (one-stream for one-thread) will also approximately require time T. As the figure I posted, even though the streams are now seemingly concurrent, but every kernel launch latency becomes also significant. Interestingly, the total time (for one batch or N streams) therefore remains approximately unchanged.

2
  • 1
    It sounds like you have an application which is poorly suited to CUDA. When kernel launch latency is your bottleneck, you are doing something wrong
    – talonmies
    Jan 19, 2018 at 6:16
  • Can't that GPU do dynamic parallelism from multiple GPU threads concurrently? Is device side launch overhead greater than host side sometimes? developer.download.nvidia.com/assets/cuda/docs/… Jan 19, 2018 at 11:29

1 Answer 1

-2

I think you need to create multiple streams. Check per-thread default stream if you are using cuda 7.0 + .

https://devblogs.nvidia.com/gpu-pro-tip-cuda-7-streams-simplify-concurrency/

Check the above link for a detailed example.

4
  • Actually I've already been using streams but the problem lies in the kernel launch. Hope you can read my post again to see what is my question. Thanks
    – wisdompeak
    Feb 8, 2018 at 9:51
  • Did you try per-thread default stream ? Compile using this : nvcc –default-stream per-thread ./stream_test.cu -o stream_per-thread Feb 8, 2018 at 11:10
  • Sorry. That does not work. Also, your suggestion is totally a different situation from my case. In my code I did not use the default stream that could potentially block other streams. All the streams in my code are truly supposed to be parallel.
    – wisdompeak
    Feb 9, 2018 at 21:39
  • Oh my bad, i thought you were using default stream Feb 11, 2018 at 6:30

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.