cudaMemcpy takes 30% of my project time.

[CUDAkk]

Hi All,

In my program , after doing all the calculations , I have to return back resultant image of size 2400 * 1800 from device to host using cudaMemcpy. But it takes 21 ms , which is very expensive for my program because it takes 30% of over all execution time.

Could any one tell why it happens ?

and How to reduce cudaMemcpy execution time?

[Sarnath]

Try using pinned memory

[eyalhir74]

Hi All,

In my program , after doing all the calculations , I have to return back resultant image of size 2400 * 1800 from device to host using cudaMemcpy. But it takes 21 ms , which is very expensive for my program because it takes 30% of over all execution time.

Could any one tell why it happens ?

and How to reduce cudaMemcpy execution time?

You could also try to overlap copying of the current result and launching the next loop’s kernel iteration (well actually the other way round).

Since kernel invocations are async you could do something like this:

// Prepare output space 1.

  // Run kernel iteration 1.

  // Run kernel iteration 2 --> this is async

  // Copy kernel iteration 1 output to output space 1.

  cudaThreadSynchronize();   --> Very Important..

  // Run kernel iteration 3 --> this is async.

  // Copy kernel iteration 2 output...

  cudaThreadSynchronize();  --> Very important.

  ...

  ...

Look for overlapping in google or maybe in the SDKs/samples.

eyal

[CUDAkk]

Try using pinned memory

I am using pinned memory in my kernel and declaring like :

int *a_h, *x;

cudaDeviceProp prop;

cudaGetDeviceProperties(&prop, 0);

if (!prop.canMapHostMemory)

	exit(0);

cudaSetDeviceFlags(cudaDeviceMapHost);

cudaHostAlloc((void **)&a_h, sizeof(int) * 256, cudaHostAllocMapped);

cudaHostGetDevicePointer((void **)&x, (void *)a_h, 0);

when I wirte

cudaError error = cudaHostGetDevicePointer((void **)&x, (void *)a_h, 0);

I got the error cudaErrorPriorLaunchFailure .

What is the cause of that error? please help.

[CUDAkk]

You could also try to overlap copying of the current result and launching the next loop’s kernel iteration (well actually the other way round).

Since kernel invocations are async you could do something like this:

// Prepare output space 1.

  // Run kernel iteration 1.

  // Run kernel iteration 2 --> this is async

  // Copy kernel iteration 1 output to output space 1.

  cudaThreadSynchronize();   --> Very Important..

  // Run kernel iteration 3 --> this is async.

  // Copy kernel iteration 2 output...

  cudaThreadSynchronize();  --> Very important.

  ...

  ...

Look for overlapping in google or maybe in the SDKs/samples.

eyal

Actually my cudaMemcpy is the last call of my program i.e. copy the final output image from device to host. So,

How can use overlap copy?

[Keither]

Hi All,

In my program , after doing all the calculations , I have to return back resultant image of size 2400 * 1800 from device to host using cudaMemcpy. But it takes 21 ms , which is very expensive for my program because it takes 30% of over all execution time.

Could any one tell why it happens ?

and How to reduce cudaMemcpy execution time?

try using:

cudaMallocHost instead of cudaMalloc

and

cudaFreeHost instead of cudaFree

declarations are exactly the same you just add “Host” to call diffrent function

it really speeded up my cudaMemcpy operations by a ?_? x100+ times (3ms → ~0.018ms this is just rough figures). it should do wonders for you as your operation takes 21 ms.

cudaMallocHost()

cudaError_t cudaMallocHost(void** hostPtr, size_t size);

allocates size bytes of host memory that is page-locked and accessible to the device. The driver tracks the virtual memory ranges allocated with this function and automatically accelerates calls to functions such as cudaMemcpy*(). Since the memory can be accessed directly by the device, it can be read or written with much higher bandwidth than pageable memory obtained with functions such as malloc(). Allocating excessive amounts of memory with cudaMallocHost() may degrade system performance, since it reduces the amount of memory available to the system for paging. As a result, this function is best used sparingly to allocate staging areas for data exchange between host and device.

cudaFreeHost()

cudaError_t cudaFreeHost(void* hostPtr);

frees the memory space pointed to by hostPtr, which must have been returned by a previous call to cudaMallocHost().