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


  // create the streams
  streams = (cudaStream_t*) malloc(2*sizeof(cudaStream_t));
  for(int i=0; i<2; i++)
    cudaStreamCreate(&(streams[i]));

  channelDesc = cudaCreateChannelDesc<COMPLEX_INPUT>();

  unsigned char timeIndex[PIPE_LENGTH*NFREQUENCY];
  for (int tf=0; tf<PIPE_LENGTH*NFREQUENCY; tf++) timeIndex[tf] = 0;
  cudaMemcpyToSymbol(tIndex, timeIndex, PIPE_LENGTH*NFREQUENCY*sizeof(unsigned char), cudaMemcpyHostToDevice);  

  //checkCudaError();

#if 0
  // check symbols are copied over

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

