CUDA Overlap Data is not working

Question

Using steams to overlap data transfer with kernel execution is not working in my system.

Hello I want to use Overlapping computation and data transfers in CUDA ,but I can't. NVIDIA help document say Overlapping computation and data transfers is possible if you use streams. but my system has not being working Please help me.

My system is below

OS : Window 7 64bit
CUDA : ver 5.0.7
Develp kit : Visual studion 2008
GPU : GTX 680

I get a profile View is Like this enter image description here

I am not getting overlapping, the code is below:

    -new pinned memory 
        cudaHostAlloc((void **)&apBuffer, sizeof(BYTE)*lBufferSize,cudaHostAllocDefault);
    -call function

   //Input Data
    for(int i=0;i<m_n3DChannelCnt*m_nBucket;++i)
    {
        cudaErrorChk_Return(cudaMemcpyAsync(d_ppbImg[i],ppbImg[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyHostToDevice,m_pStream[i/m_nBucket]));
    }
   //Call Function
    for(int i=0;i<m_n3DChannelCnt ;++i)
    {KernelGetVis8uObjPhsPhs<<<nBlockCnt,nThreadCnt,0,m_pStream[i]>>>(d_ppbVis[i],d_ppbAvg[i],d_ppfPhs[i],d_ppfObj[i],d_ppbAmp[i]
                                            ,nTotalSize,d_ppstRefData[i],d_ppbImg[i*m_nBucket],d_ppbImg[i*m_nBucket+1],d_ppbImg[i*m_nBucket+2],d_ppbImg[i*m_nBucket+3]
                                            ,fSclFloatVis2ByteVis);

    }
   //OutputData
    for(int i=0;i<m_n3DChannelCnt;++i)
    {
        if(ppbVis && ppbVis[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbVis[i],d_ppbVis[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
        if(ppbAvg && ppbAvg[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbAvg[i],d_ppbAvg[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
        if(ppfPhs && ppfPhs[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppfPhs[i],d_ppfPhs[i],sizeof(float)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
        if(ppfObj && ppfObj[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppfObj[i],d_ppfObj[i],sizeof(float)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
        if(ppbAmp && ppbAmp[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbAmp[i],d_ppbAmp[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));

    }

Please let me know about why the profiler doesn't show overlapping of kernel execution and data transfer.

Please pay __much__ more attention to the formatting and language you use in Stack Overflow questions. As you posted it, it was a nearly unintelligible mess. — talonmies, Mar 31 '13 at 08:28

score 1 · Answer 1 · answered Mar 31 '13 at 13:17

You need to invoke cudaMemcpyAsync() and kernel launches in the right order. Before compute capability 3.5 there was only a single queue for invoking device-side operations, and they do not get reordered. Combine the "Call Function" and "OutputData" phases to something like

//Call Function and OutputData
for(int i=0;i<m_n3DChannelCnt ;++i)
{KernelGetVis8uObjPhsPhs<<<nBlockCnt,nThreadCnt,0,m_pStream[i]>>>(d_ppbVis[i],d_ppbAvg[i],d_ppfPhs[i],d_ppfObj[i],d_ppbAmp[i]
                                        ,nTotalSize,d_ppstRefData[i],d_ppbImg[i*m_nBucket],d_ppbImg[i*m_nBucket+1],d_ppbImg[i*m_nBucket+2],d_ppbImg[i*m_nBucket+3]
                                        ,fSclFloatVis2ByteVis);

    if(ppbVis && ppbVis[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbVis[i],d_ppbVis[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
    if(ppbAvg && ppbAvg[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbAvg[i],d_ppbAvg[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
    if(ppfPhs && ppfPhs[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppfPhs[i],d_ppfPhs[i],sizeof(float)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
    if(ppfObj && ppfObj[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppfObj[i],d_ppfObj[i],sizeof(float)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));
    if(ppbAmp && ppbAmp[i]) cudaErrorChk_Return(cudaMemcpyAsync(ppbAmp[i],d_ppbAmp[i],sizeof(BYTE)*m_nImgWidth*m_nImgHeight,cudaMemcpyDeviceToHost,m_pStream[i]));

}

You will only be able to overlap kernel launches with the first or last of the memory copies though, as you have five cudaMemcpyAsync() calls within the stream which again don't get reordered. Allocate all five arrays contiguously in memory so that you can transfer them with a single cudaMemcpyAsync().

Overall I notice however that the data transfers takes much longer than the kernels run, so overlapping compute and copy will provide only a minor speedup in your case.

score 0 · Answer 2 · answered May 28 '13 at 02:16

You might want to check if your code works as expected (i.e., with overlapping) in LINUX. I've just incurred the same problem and found that the WINDOWS might have some problems (either in NVIDIA's driver or Windows itself), which interfere the overlapping in CUDA streaming.

You can try and check if the "simpleStreams" example in SDK works with overlapping in your machine. For my case, the "simpleStream" running on Windows does not have overlap at all, but it runs perfectly in Linux. To be specific, I am using CUDA 5.0 + VS2010, on a Fermi GTX570.

score -2 · Answer 3 · edited May 23 '17 at 10:24

TL;DR: The issue is caused by the WDDM TDR delay option in Nsight Monitor! When set to false, the issue appears. Instead, if you set the TDR delay value to a very high number, and the "enabled" option to true, the issue goes away.

Read below for other (older) steps followed until i came to the solution above, and some other possible causes.

I just recently were able to mostly solve this problem! It is specific to windows and aero i think. Please try these steps and post your results to help others! I have tried it on GTX 650 and GT 640.

Before you do anything, consider using both onboard gpu(as display) and the discrete gpu (for computations), because there are verified issues with the nvidia driver for windows! When you use onboard gpu, said drivers don't get fully loaded, so many bugs are evaded. Also, system responsiveness is maintained while working!

Make sure your concurrency problem is not related to other issues like old drivers (incl. bios version), wrong code, incapable device, etc.
Go to computer>properties
Select advanced system settings on the left side
Go to the Advanced tab
On Performance click settings
In the Visual Effects tab, select the "adjust for best performance" bullet.

This will disable aero and almost all visual effects. If this configuration works, you can try enabling one-by-one the boxes for visual effects until you find the precise one that causes problems!

Alternatively, you can:

Right click on desktop, select personalize
Select a theme from basic themes, that doesn't have aero.

This will also work as the above, but with more visual options enabled. For my two devices, this setting also works, so i kept it.

Please, when you try these solutions, come back here and post your findings!

For me, it solved the problem for most cases (a tiled dgemm i have made),but NOTE THAT i still can't run "simpleStreams" properly and achieve concurrency...

UPDATE: The problem is fully solved with a new windows installation!! The previous steps improved the behavior for some cases, but a fresh install solved all the problems!

I will try to find a less radical way of solving this problem, maybe restoring just the registry will be enough.

CUDA Overlap Data is not working

3 Answers3