   return mask;
 }
 
-__device__ __forceinline__ int inclusive_segscan_block64(
+__device__ __forceinline__ int inclusive_segscan_warp(
     int *shmem, const int packed_value, int &dist_block, int &nseg)
 {
   const int  flag = packed_value < 0;

__device__ __forceinline__ int inclusive_segscan_array(int *shmem_in, const int N)
{
  int dist, nseg = 0;
  int y  = inclusive_segscan_block64(shmem_in, shmem_in[laneId], dist, nseg);
  if (N <= WARP_SIZE) return nseg;

  for (int p = WARP_SIZE; p < N; p += WARP_SIZE)
  {
    int *shmem = shmem_in + p;
    int y1  = inclusive_segscan_block64(shmem, shmem[laneId], dist, nseg);
    shmem[laneId] += y & BTEST(laneId < dist);
    y = y1;
  }

  return mysum;
}

/* inclusive prefix sum for a warp */
__device__ __forceinline__ int inclusive_scan_warp(int* prefix, int value) 
{
  prefix[laneId] = inclusive_scan_warp<WARP_SIZE2>(value);
  return prefix[WARP_SIZE-1];
}


/* inclusive prefix sum for an array */
__device__ int inclusive_scan_array(int N, int* prefix_in) 
{

  int y = inclusive_scan_warp(prefix_in, prefix_in[laneId]);
  if (N <= WARP_SIZE) return y;

  for (int p = WARP_SIZE; p < N; p += WARP_SIZE) 
  {
    int *prefix = &prefix_in[p];
    const int y1 = inclusive_scan_warp(prefix, prefix[laneId]);
    prefix[laneId] += y;
    y += y1;
  }

  return y;
} 

/**** binary scans ****/

__device__ __forceinline__ int lanemask_lt()
{
  int mask;
  asm("mov.u32 %0, %lanemask_lt;" : "=r" (mask));
  return mask;
}


__device__ int warp_exclusive_scan(const bool p, int &psum)
{
  const unsigned int b = __ballot(p);
  psum = __popc(b & lanemask_lt());
  return __popc(b);
}

__device__ int warp_exclusive_scan(const bool p)
{
  const int b = __ballot(p);
  return __popc(b & lanemask_lt());
}

/************************************/
/********* SEGMENTED SCAN ***********/
