   grids.z = nOutputs;
 
   size_t SHM_SIZE = getSuitableShmConfig(grids, threads, imgOut.m, imgOut.n);
+  // printf("grids = (%lu, %lu, %lu), threads = (%lu, %lu, %lu) \n", grids.x, grids.y, grids.z, threads.x, threads.y, threads.z);
 
   for (size_t i=0; i<nInputs; ++i)
     for (size_t b=0; b<batch_size; ++b) {

    return;

  float sum = 0; 
  for (int i = 0; i < kW; ++i)
    for(int j = 0; j < kH; ++j)
      sum += K[ i * kH + j ] * D[ (tx + i) * HEIGHT_STEP + (ty + j) ]; 

  output[ x * fH + y ] = sum;
}


  error.resize(imgIn.area() * getNumInputMaps() + 1, delta.getCols());

  for (size_t i=0; i<delta.getCols(); ++i)
    upsample_kernel<<<grids, threads>>>(
	error.getData() + i * error.getRows(),
	delta.getData() + i * delta.getRows(),
	imgOut.m, imgOut.n, imgIn.m, imgIn.n);

  error *= 1 / (_scale * _scale);

  CCE(cudaDeviceSynchronize());
}

