 // Replace each pixel with a smoothed value.
 // This one uses shared memory, with 2d thread/block indexing:
 __global__ void avg_filter_shared(uchar *old_image_G, uchar *new_image_G) {
-  extern __shared__ uchar old_image_shared[];
+  extern __shared__ uchar old_image_shared[];  // extern means we define the size in kernel launch
 
   // blocks are 32x32 to process a 30x30 region, so we adjust size and offset to
   // handle the overlap:

  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:
