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

  // Final kernel calculation
  cudaBindTexture2D(0, tex2dfloat2, array_compute, channelDesc, NFREQUENCY*Nstation*NPOL, NTIME_PIPE, 
		    NFREQUENCY*Nstation*NPOL*sizeof(ComplexInput));
  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]));
  checkCudaError();

  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
