   cudaUnbindTexture(&textureReadDepth);
   cudaGraphicsUnmapResources(1, &inputImage);
 
-  callStages(inout, ybar, vhat, ysum, compute_width, compute_height,
+  callStages(inout, ybar, vhat, ysum, computeWidth, computeHeight,
              compute_m_size, compute_n_size);
 }
 


thrust::host_vector<float> algSAT(float *h_inout, int w, int h)
{
  int compute_width = w;
  int compute_height = w;
  if (compute_width % 32 > 0)
    compute_width += (32 - (compute_width % 32));
  if (compute_height % 32 > 0)
    compute_height += (32 - (compute_height % 32));
  int compute_m_size = (compute_width + WS - 1) / WS;
  int compute_n_size = (compute_height + WS - 1) / WS;

  // alg_setup algs;
  thrust::device_vector<float> inout(h_inout, h_inout + w * h);


void SummedAreaTable::runKernel()
{
  int computeWidth = inputImage->getWidth();
  int computeHeight = inputImage->getHeight();
  if (computeWidth % 32 > 0)
    computeWidth += (32 - (computeWidth % 32));
  if (computeHeight % 32 > 0)
    computeHeight += (32 - (computeHeight % 32));
  int compute_m_size = (computeWidth + WS - 1) / WS;
  int compute_n_size = (computeHeight + WS - 1) / WS;

  std::cout << computeWidth << "/" << computeHeight << std::endl;

  // set data structure sizes
  resizeIfNecessary(inout, computeWidth * computeHeight);
  resizeIfNecessary(ybar, compute_n_size * computeWidth);
  resizeIfNecessary(vhat, compute_m_size * computeHeight);
  resizeIfNecessary(ysum, compute_m_size * compute_n_size);

  // initialize occupancy function from inputImage
