 	}
 };
 
-////////////////////////////////////////////////////////////////////////////////
-// CTASegReduce
-// Core segmented reduction code. Supports fast-path and slow-path for intra-CTA
-// segmented reduction. Stores partials to global memory.
-// Callers feed CTASegReduce::ReduceToGlobal values in thread order.
-template<int NT, int VT, bool HalfCapacity, typename T, typename Op>
-struct CTASegReduce {
-	typedef CTASegScan<NT, Op> SegScan;
-
-	enum {
-		NV = NT * VT,
-		Capacity = HalfCapacity ? (NV / 2) : NV
-	};
-
-	union Storage {
-		typename SegScan::Storage segScanStorage;
-		T values[Capacity];
-	};
-
-	template<typename DestIt>
-	SGPU_DEVICE static void ReduceToGlobal(const int rows[VT + 1], int total,
-		int tidDelta, int startRow, int block, int tid, T data[VT],
-		DestIt dest_global, T* carryOut_global, T identity, Op op,
-		Storage& storage) {
-
-		// Run a segmented scan within the thread.
-		T x, localScan[VT];
-		#pragma unroll
-		for(int i = 0; i < VT; ++i) {
-			x = i ? op(x, data[i]) : data[i];
-			localScan[i] = x;
-			if(rows[i] != rows[i + 1]) x = identity;
-		}
-
-		// Run a parallel segmented scan over the carry-out values to compute
-		// carry-in.
-		T carryOut;
-		T carryIn = SegScan::SegScanDelta(tid, tidDelta, x,
-			storage.segScanStorage, &carryOut, identity, op);
-
-		// Store the carry-out for the entire CTA to global memory.
-		if(!tid) carryOut_global[block] = carryOut;
-
-		dest_global += startRow;
-		if(HalfCapacity && total > Capacity) {
-			// Add carry-in to each thread-local scan value. Store directly
-			// to global.
-			#pragma unroll
-			for(int i = 0; i < VT; ++i) {
-				// Add the carry-in to the local scan.
-				T x2 = op(carryIn, localScan[i]);
-
-				// Store on the end flag and clear the carry-in.
-				if(rows[i] != rows[i + 1]) {
-					carryIn = identity;
-					dest_global[rows[i]] = x2;
-				}
-			}
-		} else {
-			// All partials fit in shared memory. Add carry-in to each thread-
-			// local scan value.
-			#pragma unroll
-			for(int i = 0; i < VT; ++i) {
-				// Add the carry-in to the local scan.
-				T x2 = op(carryIn, localScan[i]);
-
-				// Store reduction when the segment changes and clear the
-				// carry-in.
-				if(rows[i] != rows[i + 1]) {
-					storage.values[rows[i]] = x2;
-					carryIn = identity;
-				}
-			}
-			__syncthreads();
-
-			// Cooperatively store reductions to global memory.
-			for(int index = tid; index < total; index += NT)
-				dest_global[index] = storage.values[index];
-			__syncthreads();
-		}
-	}
-};
-
 } // namespace sgpu
 

		}
	}

	// Load elements from a single segment and store in thread order.
	template<typename CsrIt, typename InputIt, typename SourcesIt>
	SGPU_DEVICE static void LoadIndirectFast(int tid, int gid,
		int startSeg, CsrIt csr_global, InputIt data_global,

					identity;
			}
			if(HalfCapacity)
				HalfSmemRegToThread<NT, VT>(stridedData, tid, storage.data, data);
			else {
				DeviceRegToShared<NT, VT>(stridedData, tid, storage.data);
				DeviceSharedToThread<VT>(storage.data, tid, data);
