   // Need to fill pipeline before loop
   ComplexInput *array_hp = &array_h[0*NTIME_PIPE * NFREQUENCY * Nstation * NPOL];
   CUBE_ASYNC_COPY_CALL(array_d[0], array_hp, vecLengthPipe*sizeof(ComplexInput), cudaMemcpyHostToDevice, streams[0]);
-  cudaEventRecord(copyCompletion, streams[0]); // record the completion of the h2d transfer
+  cudaEventRecord(copyCompletion[0], streams[0]); // record the completion of the h2d transfer
   checkCudaError();
 
   CUBE_ASYNC_START(PIPELINE_LOOP);

    // Kernel Calculation
    cudaBindTexture2D(0, tex2dfloat2, array_compute, channelDesc, NFREQUENCY*Nstation*NPOL, NTIME_PIPE, 
		      NFREQUENCY*Nstation*NPOL*sizeof(ComplexInput));
    cudaStreamWaitEvent(streams[1], copyCompletion, 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);
    checkCudaError();

    // Download input data
    ComplexInput *array_hp = &array_h[p*NTIME_PIPE * NFREQUENCY * Nstation * NPOL];
    CUBE_ASYNC_COPY_CALL(array_load, array_hp, vecLengthPipe*sizeof(ComplexInput), cudaMemcpyHostToDevice, streams[0]);
    cudaEventRecord(copyCompletion, streams[0]); // record the completion of the h2d transfer
    checkCudaError();
  }



  CUBE_ASYNC_END(ENTIRE_PIPELINE);

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