       b += idx;
     }
 
-    Ta a_ = Op::identity;
+    Ta a_ = Op::identity();
     int idx_ = 0;
     if (reduce_leading) {
       for (unsigned int i = 0; i < m; ++i) {

template <typename Ta, typename Tb, typename Op>
__global__ void kernel_reduce(const Ta *a, unsigned int n, Tb *b) {
  CUDA_GRID_STRIDE_LOOP(idx, 1) {
    Ta a_ = Op::identity;
    int idx_ = 0;
    for (unsigned int i = 0; i < n; ++i) {
      Op::reduce(*a, i, a_, idx_);

}; \
template <> \
struct name<int> { \
  __device__ inline static int identity() { \
    return ident_i; \
  } \
  template <typename Ta, typename Tb> \
  __device__ inline static void reduce(volatile Ta a, volatile int idx, \
                                       volatile Tb &b, volatile int &b_idx) { \
