     return x + 1;
 }
 
-#define _REDUCE(offset)                               \
-    if (tid < offset) {                               \
-        impl.Reduce(work[(tid + offset)], work[tid]); \
-    }
-
 template <
         typename In,
         typename Out,

        }
    }
}

#undef _REDUCE

}  // namespace reduce_detail


        }

        static_assert(kMaxReductionBlockSize == 512, "");
        if (out_block_size <= 256) {
            work[tid] = accum;
            __syncthreads();
            // NOTE: Compiler optimizes to unroll this loop
            for (int stride = 256; stride > 0; stride >>= 1) {
                if (out_block_size <= stride) {
                    if (tid < stride) {
                        impl.Reduce(work[tid + stride], work[tid]);
                    }
                    __syncthreads();
                }
            }
            accum = work[tid];
