title: Launching asynchronous memory copy opeerations on multiple-GPUs
body: <p>I want to divide &amp; copy an array of data on host to device memories of multiple gpus. Also, I want to do all these copy operations concurrently.</p>
  
  <p>For this I am using cudaMemcpyAsync which I launch in private stream of each GPU.</p>
  
  <p><strong>Here is what I am doing  (Doubts in code are marked with comments starting with ?? )</strong></p>
  
  <pre><code>#define SIZE 1000
  #define GPUCOUNT 2
  
  int* hostData = nullptr;
  int *devData[GPUCOUNT];
  cudaStream_t stream[GPUCOUNT];
  
  // Create one stream per GPU
  for ( int i=0; i != GPUCOUNT ; ++i )
  {    
      // DO I need to call cudaSetDevice before creating stream for each GPU ??
      cudaStreamCreate(&amp;stream[i]));
  }
  
  // Allocate pinned data on host
  cudaMallocHost (&amp;hostData, SIZE );
  
  // Allocate data on each device and copy part of host data to it
  for( int i=0; i != GPUCOUNT ; ++i )
  {
     cudaSetDevice(i);
     cudaMalloc( (void**) &amp;devData[i], sizeof(int) * SIZE/GPUCOUNT  ); // ?? Does blocking behavior of cudamalloc prevents asynch memcpy invoked in stream of other GPUs from running concurrently 
     cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, stream[i] );
  }
  
  // Some CPU code while copy is happening
  // ....
  
  // Wait for copy on all streams to finish
  cudaDeviceSynchronize();
  
  // Do something else
  </code></pre>
  
  <p>As I read the C Programming guide, I see that the above mem copy operations would not happen asynchronously because in between two consecutive async memory copy launches I am invoking a host operation which allocates device memory (blocking call).</p>
  
  <blockquote>
    <p><strong>3.2.5.5.4. Implicit Synchronization</strong></p>
    
    <p>Two commands from different streams cannot run concurrently if any one
    of the following operations is issued in-between them by the host
    thread:</p>
    
    <p>‣ a page-locked host memory allocation,</p>
    
    <p><strong>‣ a device memory allocation,</strong></p>
    
    <p>‣ a device memory set,</p>
    
    <p>‣ a memory copy between two addresses to the same device memory,</p>
    
    <p>‣ any CUDA command to the default stream,</p>
  </blockquote>
  
  <p>If the above reason seems to be true, then I need to split my memory allocation and copy operation</p>
  
  <pre><code>// Allocate data on each device 
  for( int i=0; i != GPUCOUNT ; ++i )
  {
     cudaSetDevice(i);
     cudaMalloc( (void**) &amp;devData[i], sizeof(int) * SIZE/GPUCOUNT  );
  }
  
  // Copy part of host data to each device
  for( int i=0; i != GPUCOUNT ; ++i )
  {
     // ?? DO I need to call cudaSetDevice before memory copy ??
     // CUDA guide says:"A memory copy will succeed even if it is issued to a stream that is not associated to the current device."
  
     cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, stream[i] ); 
  }
  </code></pre>
  
  <p><strong>Is my above analysis valid ?</strong></p>
  
  <p><strong>Also, Is it not possible to do this without creating an explicit per gpu stream by launching cudaMemcpyAsync operation in the default stream(stream id 0) of each GPU ?</strong>. I am basing this on following claims made in CUDA C programming guide:</p>
  
  <blockquote>
    <p>Each device has its own default stream (see Default Stream), so
    commands issued to the default stream of a device may execute out of
    order or concurrently with respect to commands issued to the default
    stream of any other device.</p>
  </blockquote>
  
  <p>The code would then look like this</p>
  
  <pre><code>#define SIZE 1000
  #define GPUCOUNT 2
  
  int* hostData = nullptr;
  int *devData[GPUCOUNT];
  
  // Allocate pinned data on host
  cudaMallocHost (&amp;hostData, SIZE );
  
  // Allocate data on each device
  for( int i=0; i != GPUCOUNT ; ++i )
  {
     cudaSetDevice(i);
     cudaMalloc( (void**) &amp;devData[i], sizeof(int) * SIZE/GPUCOUNT  );
  }
  
  // Copy part of host data to each device
  for( int i=0; i != GPUCOUNT ; ++i )
  {
     // ?? DO I need to call cudaSetDevice before memory copy ??
     // CUDA guide says:"A memory copy will succeed even if it is issued to a stream that is not associated to the current device."
  
     cudaMemcpyAsync( (void*) devData[i], hostData + i*SIZE/GPUCOUNT, SIZE/GPUCOUNT, cudaMemcpyHostToDevice, 0 ); 
  }
  
  // Some CPU code while copy is happening
  // ....
  
  // Wait for copy on all streams to finish
  cudaDeviceSynchronize();
  
  // Do something else
  </code></pre>