 	dim3 blockSize(1024);
 	dim3 gridSize((numCells * numLeftFields * numRightFields / 1024) + 1);
 
-	cudaDocontractFieldFieldScalar_kernel<<<gridSize, blockSize>>>(d_left,
+  timespec tic;
+  timespec toc;
+  if (time != NULL)
+	  clock_gettime(CLOCK_MONOTONIC, &tic);
+	
+  cudaDocontractFieldFieldScalar_kernel<<<gridSize, blockSize>>>(d_left,
 			d_right, d_out, numCells, numLeftFields, numRightFields, numPoints);
+
+  cudaDeviceSynchronize();
+  if (time != NULL) {
+    clock_gettime(CLOCK_MONOTONIC, &toc);
+    *time += getElapsedTime(tic, toc);
+  }
+  
 
 	cudaMemcpy(h_out, d_out, sizeof(double) * numCells * numLeftFields * numRightFields, cudaMemcpyDeviceToHost);
 

	clock_gettime(CLOCK_MONOTONIC, &tic);

	//repeat the calculation 5 times so we can average out some randomness
	for(int i = 0; i < 5; ++i){
		contractFieldFieldScalarKokkosCuda<Kokkos::Cuda, cuda_input_view_t,
			cuda_output_view_t, cuda_input_host_t, cuda_output_host_t>
				(cuda_hostOut, cuda_hostLeft, cuda_hostRight, cuda_kokkosOut,
				 cuda_kokkosLeft, cuda_kokkosRight);
	}

	clock_gettime(CLOCK_MONOTONIC, &toc);
	const double elapsedTime_kokkos_cuda = getElapsedTime(tic, toc);

// No checking right now
  /*

	}
  */

	std::cout << "kokkos cuda speedup of " << elapsedTime_serial/elapsedTime_kokkos_cuda_nocopy << std::endl;

	Kokkos::finalize();

