 
 // These are really only global to match/simplify the pseudocode in the slides.
 
-uchar *old_image;  // original image, CPU-side.  we don't actually use this.
+uchar *old_image;  // original image, CPU-side.  only used for pinned input.
 uchar *new_image;  // processed image, CPU-side
 int image_height;  // number of rows in image
 int image_width;   // number of columns in image

  else if (mode == 5) {
    // need to use pinned memory for streaming
    checkCudaErrors( cudaMallocHost((void**)&new_image, image_height*image_width*sizeof(uchar)) );
  }

  if (new_image == NULL) {

      int offset = i * image_size*sizeof(uchar)/nStreams;

      checkCudaErrors(  cudaMemcpyAsync( &old_image_G[offset],
					 &old_image[offset],
					 image_size*sizeof(uchar)/nStreams,
					 cudaMemcpyHostToDevice,
					 stream[i] )  );
