 
   CUBE_ASYNC_END(ENTIRE_PIPELINE);
 
-  cudaEventDestroy(copyCompletion);
-}
+  for (int i=0; i<2; i++) {
+    cudaEventDestroy(copyCompletion[i]);
+    cudaEventDestroy(kernelCompletion[i]);
+  }
+  checkCudaError();
+}

  // Final kernel calculation
  cudaBindTexture2D(0, tex2dfloat2, array_compute, channelDesc, NFREQUENCY*Nstation*NPOL, NTIME_PIPE, 
		    NFREQUENCY*Nstation*NPOL*sizeof(ComplexInput));
  cudaStreamWaitEvent(streams[1], copyCompletion, 0);
  CUBE_ASYNC_KERNEL_CALL(shared2x2float2, dimGrid, dimBlock, 0, streams[1], (float4*)matrix_real_d, (float4*)matrix_imag_d,
			 Nstation, writeMatrix);
  checkCudaError();

    // Kernel Calculation
    cudaBindTexture2D(0, tex2dfloat2, array_compute, channelDesc, NFREQUENCY*Nstation*NPOL, NTIME_PIPE, 
		      NFREQUENCY*Nstation*NPOL*sizeof(ComplexInput));
    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], 
			   (float4*)matrix_real_d, (float4*)matrix_imag_d, Nstation, writeMatrix);
    cudaEventRecord(kernelCompletion[(p+1)%2], streams[1]); // record the completion of the h2d transfer
    checkCudaError();

    // Download input data
    ComplexInput *array_hp = &array_h[p*NTIME_PIPE * NFREQUENCY * Nstation * NPOL];
    cudaStreamWaitEvent(streams[0], kernelCompletion[p%2], 0); // only start the transfer once the kernel has completed
    CUBE_ASYNC_COPY_CALL(array_load, array_hp, vecLengthPipe*sizeof(ComplexInput), cudaMemcpyHostToDevice, streams[0]);
    cudaEventRecord(copyCompletion[p%2], streams[0]); // record the completion of the h2d transfer
    checkCudaError();
  }

