     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> BlockReduce;
+    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduceT;
 
     // Allocate temp storage in shared memory
-    __shared__ typename BlockReduce::TempStorage temp_storage;
+    __shared__ typename BlockReduceT::TempStorage temp_storage;
 
     int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);
 



/// Generic reduction (full, 1)
template <typename BlockReduce, 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);
}

/// Generic reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduce, 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);
}

/// Generic reduction (partial, 1)
template <typename BlockReduce, 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);
}

/// Sum reduction (full, 1)
template <typename BlockReduce, typename T>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T (&data)[1], Sum &reduction_op)
{
    return BlockReduce(temp_storage).Sum(data[0]);
}

/// Sum reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduce, 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);
}

/// Sum reduction (partial, 1)
template <typename BlockReduce, 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);
}



    const int TILE_SIZE         = BLOCK_THREADS * ITEMS_PER_THREAD;

    // Cooperative threadblock reduction utility type (returns aggregate 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);

