     partialReduceResidual<<<rGrid, rBlock>>>(N*N, c_u, c_u2, c_res);
 
     // Finish reduce in host
-    cudaMemcpy(h_res, c_res, rGrid.x*sizeof(float), cudaMemcpyDeviceToHost);
+    cudaMemcpy(h_res, c_res, rGrid.x*sizeof(datafloat), cudaMemcpyDeviceToHost);
 
     res = 0;
     for(int i = 0; i < rGrid.x; ++i)

  s_red[threadIdx.x] = 0;

  if(id < entries){
    const float diff = u[id] - newu[id];
    s_red[threadIdx.x] = diff*diff;
  }



#define BDIM 256

#define datafloat double

/* 

   Poisson problem:  diff(u, x, 2) + diff(u, y, 2) = f
