     }
 
     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;
+}
+
 
 
 //---------------------------------------------------------------------

    if (d_keys_in) CubDebugExit(g_allocator.DeviceFree(d_keys_in));
}


/* *
 * Test DeviceSelect on iterator type
 * /
template <
    Backend         BACKEND,
    typename        Key,
    typename        Value>
void TestIterator(
    int             num_items,
    char*           key_type_string,
    char*           value_type_string,
    Int2Type<false> is_number)
{}


/ **
 * Test different gen modes
 * /
template <
    Backend         BACKEND,
    typename        Key,
    typename        Value>
void Test(
    int             num_items,
    char*           type_string)
{
    for (int max_segment = 1; max_segment < CUB_MIN(num_items, (unsigned short) -1); max_segment *= 11)
    {
        TestPointer<BACKEND, Key, Value>(num_items, 0, max_segment, key_type_string, value_type_string);
        TestPointer<BACKEND, Key, Value>(num_items, 2, max_segment, key_type_string, value_type_string);
        TestPointer<BACKEND, Key, Value>(num_items, 7, max_segment, key_type_string, value_type_string);
    }

    TestIterator<BACKEND, Key, Value>(num_items, type_string, Int2Type<Traits<T>::CATEGORY != NOT_A_NUMBER>());

}


/ **
 * Test different dispatch
 * /
template <
    typename        Key,
    typename        Value>

    int             max_segment,
    ReductionOp     reduction_op,
    char*           key_type_string,
    char*           value_type_string,
    Int2Type<true>  is_primitive)
{
    // Allocate host arrays
    Key* h_keys_in        = new Key[num_items];
    Key* h_keys_reference = new Key[num_items];

    Value one_val;
    InitValue(SEQ_INC, one_val, 1);
    ConstantInputIterator<Value, int> h_values_in(one_val);
    Value* h_values_reference = new Value[num_items];

    // Initialize problem and solution
