 
   ComplexInput **array_d = internal->array_d;
   cudaStream_t *streams = internal->streams;
+  cudaEvent_t *copyCompletion = internal->copyCompletion;
+  cudaEvent_t *kernelCompletion = internal->kernelCompletion;
   cudaChannelFormatDesc channelDesc = internal->channelDesc;
 
   // set pointers to the real and imaginary components of the device matrix

    cudaStreamWaitEvent(streams[1], copyCompletion[(p+1)%2], 0); // only start the kernel once the h2d transfer is complete
    CUBE_ASYNC_KERNEL_CALL(shared2x2float2, dimGrid, dimBlock, 0, streams[1], 
			   matrix_real_d, matrix_imag_d, NSTATION, writeMatrix);
    cudaEventRecord(kernelCompletion[(p+1)%2], streams[1]); // record the completion of the h2d transfer
    checkCudaError();

    // Download next chunk of input data

  cudaStreamWaitEvent(streams[1], copyCompletion[(PIPE_LENGTH+1)%2], 0);
  CUBE_ASYNC_KERNEL_CALL(shared2x2float2, dimGrid, dimBlock, 0, streams[1], matrix_real_d, matrix_imag_d,
			 NSTATION, writeMatrix);

  if(doDump) {
    checkCudaError();
    //copy the data back, employing a similar strategy as above
    CUBE_COPY_CALL(context->matrix_h + context->output_offset, internal->matrix_d, compiletime_info.matLength*sizeof(Complex), cudaMemcpyDeviceToHost);
    checkCudaError();
  } else {
    // record the completion of the kernel for next call
    cudaEventRecord(kernelCompletion[(PIPE_LENGTH+1)%2], streams[1]);
    checkCudaError();
  }

  CUBE_ASYNC_END(ENTIRE_PIPELINE);

  return XGPU_OK;
}
