                            inPtr, d_sum, size, stream_, CUB_DEBUG);
 }
 
-
-namespace {
-/* Functors for offsets */
-
-struct Linear {
-    Linear(IdxType _a, IdxType _b) : a(_a), b(_b) { }
-    __device__
-    IdxType operator[](IdxType idx) const { return a * idx + b; }
-    IdxType a, b;
-};
-
-}
 
 template<class real> void DeviceMathKernelsType<real>::
 sumBatched(real *d_sum, real alpha, const real *d_A, SizeType size, SizeType nBatch) {

    DEBUG_SYNC;
}


namespace {

template<class real>
struct InDotPtr {
    typedef InDotPtr<real> SelfType;
    
    __host__ __device__
    InDotPtr(const real *_d_x, const real *_d_y) : d_x(_d_x), d_y(_d_y) { }
    __device__
    real operator[](IdxType idx) const {
        return d_x[idx] * d_y[idx];
    }
    __device__
    SelfType operator+(IdxType idx) const {
        return SelfType(&d_x[idx], &d_y[idx]);
    }
    
    const real *d_x, *d_y;
};

}

namespace std {

template<class real>
struct iterator_traits<InDotPtr<real>> : sqaod_cuda::base_iterator_traits<real> { };

}


template<class real> void DeviceMathKernelsType<real>::

#include "cudafuncs.h"
#include "DeviceKernels.h"
#include "cub_iterator.cuh"
#include <cub/cub.cuh>

using sqaod::SizeType;
using sqaod::IdxType;
using sqaod::PackedBits;
using namespace sqaod_cuda;


template<class OutType, class real>  static __global__
