   //allocated exactly as many thread blocks as are needed
   dim3 dimGrid(((Nblock/2+1)*(Nblock/2))/2, compiletime_info.nfrequency);
 
-  // Create events used to record the completion of the device-host transfer and kernels
-  // TODO Store in XGPUInternalContext and move to xInit()?
-  cudaEvent_t copyCompletion[2], kernelCompletion[2];
-  for (int i=0; i<2; i++) {
-    cudaEventCreate(&kernelCompletion[i]);
-    cudaEventCreate(&copyCompletion[i]);
-  }
-  checkCudaError();
-
   CUBE_ASYNC_START(ENTIRE_PIPELINE);
 
   // Need to fill pipeline before loop
   long long unsigned int vecLengthPipe = compiletime_info.vecLengthPipe;
   ComplexInput *array_hp = context->array_h + context->input_offset;
+  // Only start the transfer once the kernel has completed
+  // (no-op unless xgpuCudaXengine() is called back-to-back with doDump=0)
+  cudaStreamWaitEvent(streams[0], kernelCompletion[0], 0);
   CUBE_ASYNC_COPY_CALL(array_d[0], array_hp, vecLengthPipe*sizeof(ComplexInput), cudaMemcpyHostToDevice, streams[0]);
   cudaEventRecord(copyCompletion[0], streams[0]); // record the completion of the h2d transfer
   checkCudaError();


  ComplexInput **array_d = internal->array_d;
  cudaStream_t *streams = internal->streams;
  cudaChannelFormatDesc channelDesc = internal->channelDesc;

  // set pointers to the real and imaginary components of the device matrix

    //assign the device
    cudaSetDevice(internal->device);

    for(int i=0; i<2; i++) {
      cudaStreamDestroy(internal->streams[i]);
      cudaEventDestroy(internal->copyCompletion[i]);
      cudaEventDestroy(internal->kernelCompletion[i]);
    }

    if(internal->free_array_h) {
      cudaFreeHost(internal->free_array_h);
