 // The 2D version seems to be about the same speed as the 1D version.
 // Apparently cache isn't helping us.
 
+// Replace each pixel with a smoothed value.  Start at offset into image, not 0.
+// This one uses global memory, and 1d thread/block indexing:
+__global__ void avg_filter_part(uchar *old_image_G, uchar *new_image_G, int offset) {
+  int idx = blockIdx.x * blockDim.x + threadIdx.x + offset;  // 1d index into image
+  if (idx >= image_height_G*image_width_G) { return; }  // off image
+  int px_x = idx % image_width_G;  // column of image
+  int px_y = idx / image_width_G;  // row of image
+  new_image_G[idx] = average_3x3(old_image_G, px_x, px_y);
+}
+
 /*
 // Replace each pixel with a smoothed value.  (shared memory version,
 // each thread probably has many pixels to process)

    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");
    exit(EXIT_FAILURE);
  }



  ///////////////////////////// Clean up and exit: /////////////////////////////

  if (mode != 4) {
    free(new_image);
  }
  else if (mode == 4) {
    checkCudaErrors( cudaFreeHost(new_image) );
  }
  
  checkCudaErrors( cudaFree(new_image_G) );
  checkCudaErrors( cudaFree(old_image_G) );
  checkCudaErrors( cudaDeviceReset() );

  exit(EXIT_SUCCESS);
