 
 
 /// Generic reduction (full, 1)
-template <typename BlockReduce, typename T, typename ReductionOp>
+template <typename BlockReduceT, typename T, typename ReductionOp>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T (&data)[1], ReductionOp &reduction_op)
-{
-    return BlockReduce(temp_storage).Reduce(data[0], reduction_op);
+    BlockReduceT &block_reduce, T (&data)[1], ReductionOp &reduction_op)
+{
+    return block_reduce.Reduce(data[0], reduction_op);
 }
 
 /// Generic reduction (full, ITEMS_PER_THREAD)
-template <typename BlockReduce, typename T, int ITEMS_PER_THREAD, typename ReductionOp>
+template <typename BlockReduceT, typename T, int ITEMS_PER_THREAD, typename ReductionOp>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T (&data)[ITEMS_PER_THREAD], ReductionOp &reduction_op)
-{
-    return BlockReduce(temp_storage).Reduce(data, reduction_op);
+    BlockReduceT &block_reduce, T (&data)[ITEMS_PER_THREAD], ReductionOp &reduction_op)
+{
+    return block_reduce.Reduce(data, reduction_op);
 }
 
 /// Generic reduction (partial, 1)
-template <typename BlockReduce, typename T, typename ReductionOp>
+template <typename BlockReduceT, typename T, typename ReductionOp>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T &data, ReductionOp &reduction_op, int valid_threads)
-{
-    return BlockReduce(temp_storage).Reduce(data, reduction_op, valid_threads);
+    BlockReduceT &block_reduce, T &data, ReductionOp &reduction_op, int valid_threads)
+{
+    return block_reduce.Reduce(data, reduction_op, valid_threads);
 }
 
 /// Sum reduction (full, 1)
-template <typename BlockReduce, typename T>
+template <typename BlockReduceT, typename T>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T (&data)[1], Sum &reduction_op)
-{
-    return BlockReduce(temp_storage).Sum(data[0]);
+    BlockReduceT &block_reduce, T (&data)[1], Sum &reduction_op)
+{
+    return block_reduce.Sum(data[0]);
 }
 
 /// Sum reduction (full, ITEMS_PER_THREAD)
-template <typename BlockReduce, typename T, int ITEMS_PER_THREAD>
+template <typename BlockReduceT, typename T, int ITEMS_PER_THREAD>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T (&data)[ITEMS_PER_THREAD], Sum &reduction_op)
-{
-    return BlockReduce(temp_storage).Sum(data);
+    BlockReduceT &block_reduce, T (&data)[ITEMS_PER_THREAD], Sum &reduction_op)
+{
+    return block_reduce.Sum(data);
 }
 
 /// Sum reduction (partial, 1)
-template <typename BlockReduce, typename T>
+template <typename BlockReduceT, typename T>
 __device__ __forceinline__ T DeviceTest(
-    typename BlockReduce::TempStorage &temp_storage, T &data, Sum &reduction_op, int valid_threads)
-{
-    return BlockReduce(temp_storage).Sum(data, valid_threads);
+    BlockReduceT &block_reduce, T &data, Sum &reduction_op, int valid_threads)
+{
+    return block_reduce.Sum(data, valid_threads);
 }
 
 

    clock_t start = clock();

    // Cooperatively reduce the tile's aggregate
    T tile_aggregate = DeviceTest<BlockReduce>(temp_storage, partial, reduction_op, num_items);

    // Stop cycle timer
#if CUB_PTX_ARCH == 100

    clock_t                 *d_elapsed)
{
    // Cooperative threadblock reduction utility type (returns aggregate only in thread-0)
    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduceT;

    // Allocate temp storage in shared memory
    __shared__ typename BlockReduceT::TempStorage temp_storage;

    int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);

