   // 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();


  CUBE_ASYNC_END(ENTIRE_PIPELINE);

}

  //allocated exactly as many thread blocks as are needed
  dim3 dimGrid(((Nblock/2+1)*(Nblock/2))/2, NFREQUENCY);

  // Create events used to record the completion of a given kernel
  cudaEvent_t copyCompletion;
  cudaEventCreate(&copyCompletion);
  checkCudaError();

  CUBE_ASYNC_START(ENTIRE_PIPELINE);

  // 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
  checkCudaError();

  CUBE_ASYNC_START(PIPELINE_LOOP);

