     return sum;
 }
 
+/* 64 < size */
+template<class InputIterator, class OffsetIterator>
+__device__ __forceinline__ static
+typename std::iterator_traits<InputIterator>::value_type
+segmentedSum_128x8(InputIterator in,
+                   OffsetIterator segOffset, sq::SizeType segLen,
+                   sq::SizeType nSegments) {
+    enum { BLOCK_DIM = 128 };
+    typedef typename std::iterator_traits<InputIterator>::value_type V;
+    typedef typename std::iterator_traits<OffsetIterator>::value_type OffsetT;
+    typedef cub::BlockReduce<V, BLOCK_DIM> BlockReduce;
+
+    int iSegment = blockIdx.x;
+    __shared__ typename BlockReduce::TempStorage temp_storage;
+
+    V sum = V();
+    if (iSegment < nSegments) {
+        OffsetT segBegin = segOffset[iSegment] + threadIdx.x;
+        V v = V();
+#pragma unroll
+        for (int idx = 0; idx < 8; ++idx) {
+            if (threadIdx.x + BLOCK_DIM * idx < segLen)
+                v += in[segBegin + BLOCK_DIM * idx];
+        }
+        sum = BlockReduce(temp_storage).Sum(v);
+    }
+    return sum;
+}
+
+template<class InputIterator, class OffsetIterator>
+__device__ __forceinline__ static
+typename std::iterator_traits<InputIterator>::value_type
+segmentedSum_64x16(InputIterator in,
+                   OffsetIterator segOffset, sq::SizeType segLen,
+                   sq::SizeType nSegments) {
+    enum { BLOCK_DIM = 64 };
+    typedef typename std::iterator_traits<InputIterator>::value_type V;
+    typedef typename std::iterator_traits<OffsetIterator>::value_type OffsetT;
+    typedef cub::BlockReduce<V, BLOCK_DIM> BlockReduce;
+
+    int iSegment = blockIdx.x;
+    __shared__ typename BlockReduce::TempStorage temp_storage;
+
+    V sum = V();
+    if (iSegment < nSegments) {
+        OffsetT segBegin = segOffset[iSegment] + threadIdx.x;
+        V v = V();
+#pragma unroll
+        for (int idx = 0; idx < 15; ++idx) {
+            v += in[segBegin + BLOCK_DIM * idx];
+        }
+        if (15 * BLOCK_DIM + threadIdx.x < segLen)
+            v += in[segBegin + 15 * BLOCK_DIM];
+        sum = BlockReduce(temp_storage).Sum(v);
+    }
+    return sum;
+}
+
+
 template<class InputIterator, class OffsetIterator>
 __device__ __forceinline__ static
 typename std::iterator_traits<InputIterator>::value_type

    dim3 blockDim(128);
    dim3 gridDim(nSegments);
    segmentedSumKernel_128<<<gridDim, blockDim, 0, stream>>>(in, out, segOffset, segLen, nSegments);
    DEBUG_SYNC;
}


    enum { BLOCK_DIM = 128 };
    typedef typename std::iterator_traits<OutputIterator>::value_type V;
    V sum = segmentedSum_128Striped(in, offsetBegin, offsetEnd, nSegments, nBlocksPerSeg);
    if (threadIdx.x == 0)
        out[blockIdx.x] = sum;
}


template<class InputIterator, class OutputIterator,
         class OffsetBeginIterator, class OffsetEndIterator>
__global__ static void
segmentedSumKernel_64x16(InputIterator in, OutputIterator out,
                         OffsetBeginIterator offsetBegin, OffsetEndIterator offsetEnd,
                         sq::SizeType nSegments) {
    enum { BLOCK_DIM = 64 };
    typedef typename std::iterator_traits<OutputIterator>::value_type V;
    V sum = segmentedSum_64x16(in, offsetBegin, offsetEnd, nSegments);
    if (threadIdx.x == 0)
        out[blockIdx.x] = sum;
}

template<class InputIterator, class OutputIterator,
         class OffsetBeginIterator, class OffsetEndIterator>
__global__ static void
segmentedSumKernel_128x8(InputIterator in, OutputIterator out,
                         OffsetBeginIterator offsetBegin, OffsetEndIterator offsetEnd,
                         sq::SizeType nSegments) {
    enum { BLOCK_DIM = 128 };
    typedef typename std::iterator_traits<OutputIterator>::value_type V;
    V sum = segmentedSum_128x8(in, offsetBegin, offsetEnd, nSegments);
    if (threadIdx.x == 0)
        out[blockIdx.x] = sum;
}
