     }
     getLastCudaError("Kernel execution failed (avg_filter).");
 
+    cudaEventRecord(lap2);
+
     // Copy result back from GPU:
     checkCudaErrors(  cudaMemcpy( new_image,
 				  new_image_G,
 				  image_size*sizeof(uchar),
 				  cudaMemcpyDeviceToHost )  );
+    cudaEventRecord(stop);
+    cudaEventSynchronize(stop);
+
+    cudaEventElapsedTime(&todevms,  start, lap1);
+    cudaEventElapsedTime(&kernelms, lap1,  lap2);
+    cudaEventElapsedTime(&tohostms, lap2,  stop);
+    cudaEventElapsedTime(&total,    start, stop);
+    // printf("GPU time (memcpy, kernel, memcpy, total): %f, %f, %f, %f ms\n",
+    // 	   todevms, kernelms, tohostms, total);
+    printf("GPU time (incl. main memcpy's): %f ms\n", total);
   }
 
   else if (mode == 5) { ////////////////// Streaming example ///////////////////

      (mode == 3) || 
      (mode == 4)) { ///////////////// Non-streaming examples //////////////////

    // Copy input image to GPU:
    checkCudaErrors(  cudaMemcpy( old_image_G,
				  image.data,
				  image_size*sizeof(uchar),
				  cudaMemcpyHostToDevice )  );

    // Run the kernel.  ceil() so we don't miss the end(s) of the image:
    if (mode == 2) {

    checkCudaErrors( cudaFree(new_image_G) );
    checkCudaErrors( cudaFree(old_image_G) );
    checkCudaErrors( cudaDeviceReset() );

    cudaEventDestroy(start);
    cudaEventDestroy(lap1);
    cudaEventDestroy(lap2);
    cudaEventDestroy(stop);
  }
  exit(EXIT_SUCCESS);
}
