CUDA memcpy async is not returning immediately

Question

Following is my code which tries to achieve non-blocking cuda memory copy host to device.

for (i = 0; i < ldu; ++i)
{
     cudaMemcpyAsync(dA+i*num_row, &A+i*LDA,
         num_row*sizeof(double), cudaMemcpyHostToDevice,streams[0]) ; 

}

Average time for each such call is around 10 microseconds. I have tried blocking version which takes 30 microseconds. 10 microseconds, for a non-blocking call seems a lot. A is allocated using cudaHostalloc. I run my code on machine equipped with 1 single Tesla C2050 , and I use cuda version 5.5 to compile the code. I have read that gpu PCI-transfer latency ( a bit irrelevant to non-blocking call, but to give an idea about the order of time) is around 5us. So it return time for a non-blocking call to be 10 us is a bit on higher side. Any thing I can do to speed it up?

Couple of things that I tried were, putting an openmp pragma (which resulted in slow down), sending data using different streams ( which gave around the same average time)

Is this in windows or linux? What time do you measure if you issue just a single call rather than a sequence of calls? — Robert Crovella
It is in linux. I'll get back to you about 2nd part in a moment — arbitUser1401
for a single call i.e. if I change ldu=1, then the return time is around 3.5 us — arbitUser1401
There seems to be a number of async copies that can be queued into a single stream, above which the queueing performance will slow down quite a bit. This number seems to be in the 1000-10000 range. So I am guessing, for your test that ldu is in this range. Here is my test. Below 1000 queued copies, the average overhead is less than 3us. Above 10000 copies, the average overhead is over 20us. However I believe you can replace your copy loop with a single call to cudaMemcpy2DAsync. — Robert Crovella

Robert Crovella Robert Crovella · Accepted Answer · 2013-12-26T17:25:50

Any thing I can do to speed it up?

I believe you can replace your copy loop with:

cudaMemcpy2DAsync(dA, num_row*sizeof(double), &A, LDA*sizeof(double), num_row*sizeof(double), ldu, cudaMemcpyHostToDevice, streams[0]);

which should speed things up (at least from a call overhead standpoint) significantly.

You may have to play with your parameters a bit, as your names are somewhat confusing to me (maybe you are using column-major storage). The cudaMemcpy2DAsync function is documented here.

CUDA memcpy async is not returning immediately

2 Answers