 }; \
 template <> \
 struct name<int> { \
-  const static int identity = ident_i; \
+  __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) { \

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_);

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