title: CuPY: Not seeing kernel concurrency
body: <p>I'm currently working on parallelizing huge matrix computation using CuPY's RawKernels using async streams.</p>
  <p>It seems like each RawKernel call is waiting for prevous kernel to finish eventhough I specify that stream is non-blocking.</p>
  <p><img src="https://i.stack.imgur.com/8GhTD.png" alt="Visual profiler" /></p>
  <p>Does anyone have an idea on what I'm doing wrong?</p>
  <p>Here's a simple example that creates 32 streams. Each stream should copy single slice of 3D input array to 3D output array.</p>
  <pre><code>import cupy
  
  kernel = cupy.RawKernel(
      '''
      extern &quot;C&quot;
      __global__ void simple_copy(float* iArr, float* oArr, int rows, int cols, int slice){
          unsigned int col = blockDim.x*blockIdx.x + threadIdx.x;
          unsigned int row = blockDim.y*blockIdx.y + threadIdx.y;
      
          if(row &lt; rows &amp;&amp; col &lt; cols){
  //this for loop is just additional work to see kernel launches in visual profiler more easily
              for(int i=0; i&lt;1000; i++){
                  oArr[rows*cols*slice + row*cols + col] = iArr[rows*cols*slice + row*cols + col];
              }
          }        
      } 
      '''
      , 'simple_copy')
  
  
  device = cupy.cuda.Device()
  # [x, y, z]
  iArr1 = cupy.ones((32*32, 32*32, 32), dtype=cupy.float32)
  oArr1 = cupy.zeros((32*32, 32*32, 32), dtype=cupy.float32)
  
  n = 32
  map_streams = []
  for i in range(n):
      map_streams.append(cupy.cuda.stream.Stream(non_blocking=True))
  
  # I want to run kernel on individual z-axis slice asynchronous
  for i, stream in enumerate(map_streams):
      with stream:
          kernel((32, 32), (32, 32), (iArr1, oArr1, 32*32, 32*32, i))
  device.synchronize()
  </code></pre>