     }
 
     return cudaSuccess;
-
-}
+}
+
+
+/**
+ * Dispatch to run-length encode entrypoint
+ */
+template <
+    typename                    KeyInputIterator,
+    typename                    KeyOutputIterator,
+    typename                    ValueOutputIterator,
+    typename                    NumSegmentsIterator,
+    typename                    Offset>
+cudaError_t Dispatch(
+    Int2Type<THRUST>            dispatch_to,
+    int                         timing_timing_iterations,
+    size_t                      *d_temp_storage_bytes,
+    cudaError_t                 *d_cdp_error,
+
+    void                        *d_temp_storage,
+    size_t                      &temp_storage_bytes,
+    KeyInputIterator            d_keys_in,
+    KeyOutputIterator           d_keys_out,
+    ConstantInputIterator<typename std::iterator_traits<ValueOutputIterator>::value_type, Offset> d_values_in,
+    ValueOutputIterator         d_values_out,
+    NumSegmentsIterator         d_num_segments,
+    cub::Equality               equality_op,
+    cub::Sum                    reduction_op,
+    Offset                      num_items,
+    cudaStream_t                stream,
+    bool                        debug_synchronous)
+{
+    typedef typename std::iterator_traits<KeyInputIterator>::value_type Key;
+    typedef typename std::iterator_traits<ValueOutputIterator>::value_type Value;
+
+    if (d_temp_storage == 0)
+    {
+        temp_storage_bytes = 1;
+    }
+    else
+    {
+        thrust::device_ptr<Key>     d_keys_in_wrapper(d_keys_in);
+        thrust::device_ptr<Key>     d_keys_out_wrapper(d_keys_out);
+        thrust::device_ptr<Value>   d_values_out_wrapper(d_values_out);
+
+        thrust::pair<thrust::device_ptr<Key>, thrust::device_ptr<Value> > d_out_ends;
+
+        Value one_val;
+        InitValue(SEQ_INC, one_val, 1);
+        thrust::constant_iterator<Value> constant_one(one_val);
+
+        for (int i = 0; i < timing_timing_iterations; ++i)
+        {
+            d_out_ends = thrust::reduce_by_key(
+                d_keys_in_wrapper,
+                d_keys_in_wrapper + num_items,
+//                d_values_in,
+                constant_one,
+                d_keys_out_wrapper,
+                d_values_out_wrapper);
+        }
+
+        Offset num_segments = d_out_ends.first - d_keys_out_wrapper;
+        CubDebugExit(cudaMemcpy(d_num_segments, &num_segments, sizeof(Offset), cudaMemcpyHostToDevice));
+    }
+
+    return cudaSuccess;
+}
+
 
 
 //---------------------------------------------------------------------



/**
 * Dispatch to unique entrypoint
 */
template <
    typename                    KeyInputIterator,

    int sm_version;
    CubDebugExit(SmVersion(sm_version, device_ordinal));

    TestPointer<CUB, int, int>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(int));
//    TestOp<int, int>(num_items, CUB_TYPE_STRING(int), CUB_TYPE_STRING(int));
//    TestOp<int, uint4>(num_items, CUB_TYPE_STRING(int), CUB_TYPE_STRING(uint4));

/*
#ifdef QUICK_TEST

    // Compile/run quick tests
    if (num_items < 0) num_items = 32000000;

    printf("int\n");
    TestPointer<THRUST, int, int>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(int));
    TestPointer<CUB, int, int>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(int));

    printf("RLE");
    TestIterator<THRUST, int, int>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(int), Int2Type<Traits<int>::PRIMITIVE>());
    TestIterator<CUB, int, int>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(int), Int2Type<Traits<int>::PRIMITIVE>());

    printf("float\n");
    TestPointer<THRUST, int, float>(num_items, entropy_reduction, maxseg, cub::Sum(), CUB_TYPE_STRING(int), CUB_TYPE_STRING(float));
