         clock_t start = clock();
 
         // Cooperatively reduce the tile's aggregate
-        T tile_aggregate = DeviceTest<BlockReduce>(temp_storage, data, reduction_op);
+        T tile_aggregate = DeviceTest(block_reduce, data, reduction_op);
 
         // Stop cycle timer
     #if CUB_PTX_ARCH == 100

    clock_t start = clock();

    // Cooperative reduce first tile
    T block_aggregate = DeviceTest<BlockReduce>(temp_storage, data, reduction_op);

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

