     printf("Usage: %s <input image> <output image> <mode>\n", argv[0]);
     printf("       mode 0 = CPU, single thread\n");
     printf("       mode 1 = CPU, 4 threads\n");
-    printf("       mode 2 = GPU, global memory\n");
-    printf("       mode 3 = CPU, global memory, 2d grid\n");
+    printf("       mode 2 = GPU, global memory, 1d grid\n");
+    printf("       mode 3 = GPU, global memory, 2d grid\n");
+    printf("       mode 4 = GPU, global memory, 1d grid, 4 streams\n");
     exit(EXIT_FAILURE);
   }
 

  checkCudaErrors( cudaMalloc((void**)&old_image_G,image_size*sizeof(uchar)) );
  checkCudaErrors( cudaMalloc((void**)&new_image_G,image_size*sizeof(uchar)) );

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

  checkCudaErrors( cudaMemcpyToSymbol(image_height_G, &image_height, sizeof(int)) );
  checkCudaErrors( cudaMemcpyToSymbol(image_width_G,  &image_width,  sizeof(int)) );

  // Set GPU output to zero (optional):
  //checkCudaErrors( cudaMemset(new_image_G, 0, image_size) );

  // Run the kernel.  ceil() so we don't miss the end(s) of the image:
  if (mode == 2) {
    avg_filter<<< ceil((float)image_size/256), 256 >>>(old_image_G, new_image_G);
  }
  if (mode == 3) {
    dim3 block_dim = dim3(32, 32);
    dim3 grid_dim = dim3( ceil((float)image_width / block_dim.x),
    			  ceil((float)image_height / block_dim.y) );
    avg_filter_2D<<<grid_dim,block_dim>>>(old_image_G, new_image_G);
  }
  getLastCudaError("Kernel execution failed (avg_filter).");
  // checkCudaErrors( cudaDeviceSynchronize() );  // not needed

  // Copy result back:
  checkCudaErrors(  cudaMemcpy( new_image,
				new_image_G,
				image_size*sizeof(uchar),
				cudaMemcpyDeviceToHost )  );

  ////////////////////////// Save the output to disk: //////////////////////////

  Mat result = Mat(image_height, image_width, CV_8UC1, new_image);

  //int chunk_size = image_size / N;

  // Prepare CPU output array:
  if (mode != 4) {
    new_image = (uchar*)malloc( image_height*image_width*sizeof(uchar) );
  }
  else if (mode == 4) {
    // need to use pinned memory for streaming
    checkCudaErrors( cudaMallocHost((void**)&new_image, image_height*image_width*sizeof(uchar)) );
  }

  if (new_image == NULL) {
    fprintf(stderr, "Can't malloc space for new_image.  Exiting abruptly.\n");
    exit(EXIT_FAILURE);
