       int offset = i * image_size*sizeof(uchar)/nStreams;
 
       checkCudaErrors(  cudaMemcpyAsync( &old_image_G[offset],
-					 &image.data[offset],
+					 &old_image[offset],
 					 image_size*sizeof(uchar)/nStreams,
 					 cudaMemcpyHostToDevice,
 					 stream[i] )  );

  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) {


// These are really only global to match/simplify the pseudocode in the slides.

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
