         identity,
         p_prefix);
 
+    // Test reference aggregate is returned in all threads
     for (int i = 0; i < BLOCK_THREADS; ++i)
     {
         h_aggregate[i] = aggregate;

    T       *d_out = NULL;
    T       *d_aggregate = NULL;
    clock_t *d_elapsed = NULL;
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(T) * TILE_SIZE));
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_out, sizeof(T) * (TILE_SIZE + 1)));
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_aggregate, sizeof(T) * BLOCK_THREADS));
    CubDebugExit(g_allocator.DeviceAllocate((void**)&d_elapsed, sizeof(clock_t)));
    CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(T) * TILE_SIZE, cudaMemcpyHostToDevice));
    CubDebugExit(cudaMemset(d_out, 0, sizeof(T) * (TILE_SIZE + 1)));
    CubDebugExit(cudaMemset(d_aggregate, 0, sizeof(T) * BLOCK_THREADS));

    if (threadIdx.x == 0)
    {
        d_out[TILE_SIZE] = prefix_op.prefix;
    }
    else
    {
        // In a separate conditional as a workaround for SM10 Win32 bug where we write the wrong prefix
        *d_elapsed = (start > stop) ? start - stop : stop - start;
    }
}
