                       temp.values.begin(),
                       thrust::multiplies<ValueType>());
     thrust::transform(temp.values.begin(), temp.values.begin() + S.num_entries,
-                      thrust::constant_iterator<ValueType>(lambda),
+                      thrust::constant_iterator<ValueType>(-lambda),
                       temp.values.begin(),
                       thrust::multiplies<ValueType>());
     // temp <- D^-1

                     cusp::array1d<ValueType,cusp::device_memory>& x) const
    {
        // TODO check sizes
        cusp::coo_matrix<IndexType,ValueType,MemorySpace> & A = levels[0].A;

        cusp::array1d<ValueType,MemorySpace> residual(A.num_rows);  // TODO eliminate temporaries
            

            cusp::blas::axpby(b, residual, residual, 1.0f, -1.0f);
            ValueType norm = cusp::blas::nrm2(residual);

            //printf("%10.8f  %6.4f\n", norm, norm/last_norm);

            last_norm = norm;
        }
