     }
     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 ///////////////////

    for (int i = 0; i < nStreams; i++) {
      checkCudaErrors( cudaStreamCreate(&stream[i]) );
    }

    // do {copy to device, launch kernel, copy to host} for each stream:
    for (int i = 0; i < nStreams; i++) {

  }

  if (mode==0) { ///////////////// Single-threaded CPU version /////////////////
    clock_t tic, toc;
    tic = clock();
    for (int row=0; row<image_height; row++) {
      for (int col=0; col<image_width; col++) {
	  new_image[row*image_width + col] = average_3x3_CPU(image.data, col, row);
      }
    }    
    toc = clock();
    printf("CPU time (single thread): %f ms\n", (double)(toc - tic)*1000 / CLOCKS_PER_SEC);
  }

  ///////////////////////////////// GPU code: //////////////////////////////////

  int devID;  // which GPU to use

  // for timing:
  cudaEvent_t start, lap1, lap2, stop;
  float todevms, kernelms, tohostms, total;

  if (mode > 1) {
    // create timers:
    cudaEventCreate(&start);
    cudaEventCreate(&lap1);  // after memcpy to GPU
    cudaEventCreate(&lap2);  // after kernel completes
    cudaEventCreate(&stop);  // after memcpy to CPU

    // Choose the fastest GPU (optional):
    devID = gpuGetMaxGflopsDeviceId();
    checkCudaErrors( cudaSetDevice(devID) );
