   // do my part to populate the shared memory:
   if ((px_y < 0) || (px_y >= image_height_G)) { return; }
   if ((px_x < 0) || (px_x >= image_width_G )) { return; }
-  old_image_shared[ threadIdx.y*blockDim.x + threadIdx.x ] = px_y*image_width_G+px_x;
+  old_image_shared[ threadIdx.y*blockDim.x + threadIdx.x ] = old_image_G[ px_y*image_width_G+px_x ];
   __syncthreads();  // make sure all the data is here before we continue
 
   // kill the threads that aren't in the inner 30x30 block:

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

  if ((mode == 2) || (mode == 3)) { ////////// Non-streaming examples //////////

    // Copy input image to GPU:
    checkCudaErrors(  cudaMemcpy( old_image_G,

      avg_filter_shared<<<grid_dim,block_dim,32*32*sizeof(uchar)>>>
	(old_image_G, new_image_G);
    }
    getLastCudaError("Kernel execution failed (avg_filter).");

    // Copy result back from GPU:
