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

				stridedData, identity);

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

}

////////////////////////////////////////////////////////////////////////////////
// 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();
		}
	}
};

////////////////////////////////////////////////////////////////////////////////
// CTASegReduceLoad
// Load logic for segmented reduce and interval reduce. Supports direct and
// indirect loads, and has fast- and slow-path indirect loads.
