/******************************************************************************
   * Copyright (c) 2011, Duane Merrill.  All rights reserved.
   * Copyright (c) 2011-2018, NVIDIA CORPORATION.  All rights reserved.
   *
   * Redistribution and use in source and binary forms, with or without
   * modification, are permitted provided that the following conditions are met:
   *     * Redistributions of source code must retain the above copyright
   *       notice, this list of conditions and the following disclaimer.
   *     * Redistributions in binary form must reproduce the above copyright
   *       notice, this list of conditions and the following disclaimer in the
   *       documentation and/or other materials provided with the distribution.
   *     * Neither the name of the NVIDIA CORPORATION nor the
   *       names of its contributors may be used to endorse or promote products
   *       derived from this software without specific prior written permission.
   *
   * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
   * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
   * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
   * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
   * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
   * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
   * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
   * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
   * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
   * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
   *
   ******************************************************************************/
  
  /******************************************************************************
   * Test of BlockScan utilities
   ******************************************************************************/
  
  // Ensure printing of CUDA runtime errors to console
  #define CUB_STDERR
  
  #include <stdio.h>
  #include <iostream>
  #include <limits>
  #include <typeinfo>
  
  #include <cub/block/block_scan.cuh>
  #include <cub/block/block_load.cuh>
  #include <cub/block/block_store.cuh>
  #include <cub/util_ptx.cuh>
  #include <cub/util_allocator.cuh>
  
  #include "test_util.h"
  
  
  using namespace cub;
  
  //---------------------------------------------------------------------
  // Globals, constants and typedefs
  //---------------------------------------------------------------------
  
  bool                    g_verbose       = false;
  int                     g_repeat        = 0;
  CachingDeviceAllocator  g_allocator(true);
  
  
  /**
   * Primitive variant to test
   */
  enum TestMode
  {
      BASIC,
      AGGREGATE,
      PREFIX,
  };
  
  
  /**
   * Scan mode to test
   */
  enum ScanMode
  {
      EXCLUSIVE,
      INCLUSIVE
  };
  
  
  /**
   * \brief WrapperFunctor (for precluding test-specialized dispatch to *Sum variants)
   */
  template<typename OpT>
  struct WrapperFunctor
  {
      OpT op;
  
      WrapperFunctor(OpT op) : op(op) {}
  
      template <typename T>
      __host__ __device__ __forceinline__ T operator()(const T &a, const T &b) const
      {
          return op(a, b);
      }
  };
  
  
  /**
   * Stateful prefix functor
   */
  template <
      typename T,
      typename ScanOpT>
  struct BlockPrefixCallbackOp
  {
      int     linear_tid;
      T       prefix;
      ScanOpT  scan_op;
  
      __device__ __forceinline__
      BlockPrefixCallbackOp(int linear_tid, T prefix, ScanOpT scan_op) :
          linear_tid(linear_tid),
          prefix(prefix),
          scan_op(scan_op)
      {}
  
      __device__ __forceinline__
      T operator()(T block_aggregate)
      {
          // For testing purposes
          T retval = (linear_tid == 0) ? prefix  : T();
          prefix = scan_op(prefix, block_aggregate);
          return retval;
      }
  };
  
  
  //---------------------------------------------------------------------
  // Exclusive scan
  //---------------------------------------------------------------------
  
  /// Exclusive scan (BASIC, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data[0], data[0], initial_value, scan_op);
  }
  
  /// Exclusive scan (BASIC, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data, data, initial_value, scan_op);
  }
  
  /// Exclusive scan (AGGREGATE, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data[0], data[0], initial_value, scan_op, block_aggregate);
  }
  
  /// Exclusive scan (AGGREGATE, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data, data, initial_value, scan_op, block_aggregate);
  }
  
  /// Exclusive scan (PREFIX, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data[0], data[0], scan_op, prefix_op);
  }
  
  /// Exclusive scan (PREFIX, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.ExclusiveScan(data, data, scan_op, prefix_op);
  }
  
  
  //---------------------------------------------------------------------
  // Exclusive sum
  //---------------------------------------------------------------------
  
  /// Exclusive sum (BASIC, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data[0], data[0]);
  }
  
  /// Exclusive sum (BASIC, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data, data);
  }
  
  /// Exclusive sum (AGGREGATE, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data[0], data[0], block_aggregate);
  }
  
  /// Exclusive sum (AGGREGATE, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data, data, block_aggregate);
  }
  
  /// Exclusive sum (PREFIX, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data[0], data[0], prefix_op);
  }
  
  /// Exclusive sum (PREFIX, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<EXCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.ExclusiveSum(data, data, prefix_op);
  }
  
  
  //---------------------------------------------------------------------
  // Inclusive scan
  //---------------------------------------------------------------------
  
  /// Inclusive scan (BASIC, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data[0], data[0], scan_op);
  }
  
  /// Inclusive scan (BASIC, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data, data, scan_op);
  }
  
  /// Inclusive scan (AGGREGATE, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data[0], data[0], scan_op, block_aggregate);
  }
  
  /// Inclusive scan (AGGREGATE, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data, data, scan_op, block_aggregate);
  }
  
  /// Inclusive scan (PREFIX, 1)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data[0], data[0], scan_op, prefix_op);
  }
  
  /// Inclusive scan (PREFIX, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename ScanOpT, typename PrefixCallbackOp, int ITEMS_PER_THREAD, typename IsPrimitiveT>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, ScanOpT &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, IsPrimitiveT is_primitive)
  {
      block_scan.InclusiveScan(data, data, scan_op, prefix_op);
  }
  
  
  //---------------------------------------------------------------------
  // Inclusive sum
  //---------------------------------------------------------------------
  
  /// Inclusive sum (BASIC, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data[0], data[0]);
  }
  
  /// Inclusive sum (BASIC, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<BASIC> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data, data);
  }
  
  /// Inclusive sum (AGGREGATE, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data[0], data[0], block_aggregate);
  }
  
  /// Inclusive sum (AGGREGATE, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<AGGREGATE> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data, data, block_aggregate);
  }
  
  /// Inclusive sum (PREFIX, 1)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[1], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data[0], data[0], prefix_op);
  }
  
  /// Inclusive sum (PREFIX, ITEMS_PER_THREAD)
  template <typename BlockScanT, typename T, typename PrefixCallbackOp, int ITEMS_PER_THREAD>
  __device__ __forceinline__ void DeviceTest(
      BlockScanT &block_scan, T (&data)[ITEMS_PER_THREAD], T &initial_value, Sum &scan_op, T &block_aggregate, PrefixCallbackOp &prefix_op,
      Int2Type<INCLUSIVE> scan_mode, Int2Type<PREFIX> test_mode, Int2Type<true> is_primitive)
  {
      block_scan.InclusiveSum(data, data, prefix_op);
  }
  
  
  
  //---------------------------------------------------------------------
  // Test kernels
  //---------------------------------------------------------------------
  
  /**
   * BlockScan test kernel.
   */
  template <
      int                 BLOCK_DIM_X,
      int                 BLOCK_DIM_Y,
      int                 BLOCK_DIM_Z,
      int                 ITEMS_PER_THREAD,
      ScanMode            SCAN_MODE,
      TestMode            TEST_MODE,
      BlockScanAlgorithm  ALGORITHM,
      typename            T,
      typename            ScanOpT>
  __launch_bounds__ (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)
  __global__ void BlockScanKernel(
      T                   *d_in,
      T                   *d_out,
      T                   *d_aggregate,
      ScanOpT              scan_op,
      T                   initial_value,
      clock_t             *d_elapsed)
  {
      const int BLOCK_THREADS     = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z;
      const int TILE_SIZE         = BLOCK_THREADS * ITEMS_PER_THREAD;
  
      // Parameterize BlockScan type for our thread block
      typedef BlockScan<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockScanT;
  
      // Allocate temp storage in shared memory
      __shared__ typename BlockScanT::TempStorage temp_storage;
  
      int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);
  
      // Per-thread tile data
      T data[ITEMS_PER_THREAD];
      LoadDirectBlocked(linear_tid, d_in, data);
  
      __threadfence_block();      // workaround to prevent clock hoisting
      clock_t start = clock();
      __threadfence_block();      // workaround to prevent clock hoisting
  
      // Test scan
      T                                   block_aggregate;
      BlockScanT                          block_scan(temp_storage);
      BlockPrefixCallbackOp<T, ScanOpT>   prefix_op(linear_tid, initial_value, scan_op);
  
      DeviceTest(block_scan, data, initial_value, scan_op, block_aggregate, prefix_op,
          Int2Type<SCAN_MODE>(), Int2Type<TEST_MODE>(), Int2Type<Traits<T>::PRIMITIVE>());
  
      // Stop cycle timer
      __threadfence_block();      // workaround to prevent clock hoisting
      clock_t stop = clock();
      __threadfence_block();      // workaround to prevent clock hoisting
  
      // Store output
      StoreDirectBlocked(linear_tid, d_out, data);
  
      // Store block_aggregate
      if (TEST_MODE != BASIC)
          d_aggregate[linear_tid] = block_aggregate;
  
      // Store prefix
      if (TEST_MODE == PREFIX)
      {
          if (linear_tid == 0)
              d_out[TILE_SIZE] = prefix_op.prefix;
      }
  
      // Store time
      if (linear_tid == 0)
          *d_elapsed = (start > stop) ? start - stop : stop - start;
  }
  
  
  
  //---------------------------------------------------------------------
  // Host utility subroutines
  //---------------------------------------------------------------------
  
  /**
   * Initialize exclusive-scan problem (and solution)
   */
  template <typename T, typename ScanOpT>
  T Initialize(
      GenMode     gen_mode,
      T           *h_in,
      T           *h_reference,
      int         num_items,
      ScanOpT     scan_op,
      T           initial_value,
      Int2Type<EXCLUSIVE>)
  {
      InitValue(gen_mode, h_in[0], 0);
  
      T block_aggregate   = h_in[0];
      h_reference[0]      = initial_value;
      T inclusive         = scan_op(initial_value, h_in[0]);
  
      for (int i = 1; i < num_items; ++i)
      {
          InitValue(gen_mode, h_in[i], i);
          h_reference[i] = inclusive;
          inclusive = scan_op(inclusive, h_in[i]);
          block_aggregate = scan_op(block_aggregate, h_in[i]);
      }
  
      return block_aggregate;
  }
  
  
  /**
   * Initialize inclusive-scan problem (and solution)
   */
  template <typename T, typename ScanOpT>
  T Initialize(
      GenMode     gen_mode,
      T           *h_in,
      T           *h_reference,
      int         num_items,
      ScanOpT      scan_op,
      T           initial_value,
      Int2Type<INCLUSIVE>)
  {
      InitValue(gen_mode, h_in[0], 0);
  
      T block_aggregate   = h_in[0];
      T inclusive         = scan_op(initial_value, h_in[0]);
      h_reference[0]      = inclusive;
  
      for (int i = 1; i < num_items; ++i)
      {
          InitValue(gen_mode, h_in[i], i);
          inclusive = scan_op(inclusive, h_in[i]);
          block_aggregate = scan_op(block_aggregate, h_in[i]);
          h_reference[i] = inclusive;
      }
  
      return block_aggregate;
  }
  
  
  /**
   * Test thread block scan.  (Specialized for sufficient resources)
   */
  template <
      int                 BLOCK_DIM_X,
      int                 BLOCK_DIM_Y,
      int                 BLOCK_DIM_Z,
      int                 ITEMS_PER_THREAD,
      ScanMode            SCAN_MODE,
      TestMode            TEST_MODE,
      BlockScanAlgorithm  ALGORITHM,
      typename            ScanOpT,
      typename            T>
  void Test(
      GenMode             gen_mode,
      ScanOpT             scan_op,
      T                   initial_value,
      Int2Type<true>      sufficient_resources)
  {
      const int BLOCK_THREADS     = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z;
      const int TILE_SIZE         = BLOCK_THREADS * ITEMS_PER_THREAD;
  
      // Allocate host arrays
      T *h_in = new T[TILE_SIZE];
      T *h_reference = new T[TILE_SIZE];
      T *h_aggregate = new T[BLOCK_THREADS];
  
      // Initialize problem
      T block_aggregate = Initialize(
          gen_mode,
          h_in,
          h_reference,
          TILE_SIZE,
          scan_op,
          initial_value,
          Int2Type<SCAN_MODE>());
  
      // Test reference block_aggregate is returned in all threads
      for (int i = 0; i < BLOCK_THREADS; ++i)
      {
          h_aggregate[i] = block_aggregate;
      }
  
      // Run kernel
      printf("Test-mode %d, gen-mode %d, policy %d, %s %s BlockScan, %d (%d,%d,%d) thread block threads, %d items per thread, %d tile size, %s (%d bytes) elements:
  ",
          TEST_MODE, gen_mode, ALGORITHM,
          (SCAN_MODE == INCLUSIVE) ? "Inclusive" : "Exclusive", typeid(ScanOpT).name(),
          BLOCK_THREADS, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z,
          ITEMS_PER_THREAD,  TILE_SIZE,
          typeid(T).name(), (int) sizeof(T));
      fflush(stdout);
  
      // Initialize/clear device arrays
      T       *d_in = NULL;
      T       *d_out = NULL;
      T       *d_aggregate = NULL;
      clock_t *d_elapsed = NULL;
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_elapsed, sizeof(unsigned long long)));
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(T) * TILE_SIZE));
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_out, sizeof(T) * (TILE_SIZE + 2)));
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_aggregate, sizeof(T) * BLOCK_THREADS));
      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));
  
      // Display input problem data
      if (g_verbose)
      {
          printf("Input data: ");
          for (int i = 0; i < TILE_SIZE; i++)
          {
              std::cout << CoutCast(h_in[i]) << ", ";
          }
          printf("
  
  ");
      }
  
      // Run block_aggregate/prefix kernel
      dim3 block_dims(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);
      BlockScanKernel<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, ALGORITHM><<<1, block_dims>>>(
          d_in,
          d_out,
          d_aggregate,
          scan_op,
          initial_value,
          d_elapsed);
  
      CubDebugExit(cudaPeekAtLastError());
      CubDebugExit(cudaDeviceSynchronize());
  
      // Copy out and display results
      printf("\tScan results: ");
      int compare = CompareDeviceResults(h_reference, d_out, TILE_SIZE, g_verbose, g_verbose);
      printf("%s
  ", compare ? "FAIL" : "PASS");
      AssertEquals(0, compare);
  
      if (TEST_MODE == AGGREGATE)
      {
          // Copy out and display block_aggregate
          printf("\tScan block aggregate: ");
          compare = CompareDeviceResults(h_aggregate, d_aggregate, BLOCK_THREADS, g_verbose, g_verbose);
          printf("%s
  ", compare ? "FAIL" : "PASS");
          AssertEquals(0, compare);
      }
  
      if (TEST_MODE == PREFIX)
      {
          // Copy out and display updated prefix
          printf("\tScan running total: ");
          T running_total = scan_op(initial_value, block_aggregate);
          compare = CompareDeviceResults(&running_total, d_out + TILE_SIZE, 1, g_verbose, g_verbose);
          printf("%s
  ", compare ? "FAIL" : "PASS");
          AssertEquals(0, compare);
      }
  
      printf("\tElapsed clocks: ");
      DisplayDeviceResults(d_elapsed, 1);
  
      // Cleanup
      if (h_in) delete[] h_in;
      if (h_reference) delete[] h_reference;
      if (h_aggregate) delete[] h_aggregate;
      if (d_in) CubDebugExit(g_allocator.DeviceFree(d_in));
      if (d_out) CubDebugExit(g_allocator.DeviceFree(d_out));
      if (d_aggregate) CubDebugExit(g_allocator.DeviceFree(d_aggregate));
      if (d_elapsed) CubDebugExit(g_allocator.DeviceFree(d_elapsed));
  }
  
  
  /**
   * Test thread block scan.  (Specialized for insufficient resources)
   */
  template <
      int                 BLOCK_DIM_X,
      int                 BLOCK_DIM_Y,
      int                 BLOCK_DIM_Z,
      int                 ITEMS_PER_THREAD,
      ScanMode            SCAN_MODE,
      TestMode            TEST_MODE,
      BlockScanAlgorithm  ALGORITHM,
      typename            ScanOpT,
      typename            T>
  void Test(
      GenMode             gen_mode,
      ScanOpT             scan_op,
      T                   initial_value,
      Int2Type<false>     sufficient_resources)
  {}
  
  
  /**
   * Test thread block scan.
   */
  template <
      int                 BLOCK_DIM_X,
      int                 BLOCK_DIM_Y,
      int                 BLOCK_DIM_Z,
      int                 ITEMS_PER_THREAD,
      ScanMode            SCAN_MODE,
      TestMode            TEST_MODE,
      BlockScanAlgorithm  ALGORITHM,
      typename            ScanOpT,
      typename            T>
  void Test(
      GenMode             gen_mode,
      ScanOpT             scan_op,
      T                   initial_value)
  {
      // Check size of smem storage for the target arch to make sure it will fit
      typedef BlockScan<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockScanT;
  
      enum
      {
  #if defined(SM100) || defined(SM110) || defined(SM130)
          sufficient_smem         = (sizeof(typename BlockScanT::TempStorage)     <= 16 * 1024),
          sufficient_threads      = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)    <= 512),
  #else
          sufficient_smem         = (sizeof(typename BlockScanT::TempStorage)     <= 16 * 1024),
          sufficient_threads      = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)    <= 1024),
  #endif
  
  #if defined(_WIN32) || defined(_WIN64)
          // Accommodate ptxas crash bug (access violation) on Windows
          special_skip            = ((TEST_ARCH <= 130) && (Equals<T, TestBar>::VALUE) && (BLOCK_DIM_Z > 1)),
  #else
          special_skip            = false,
  #endif
          sufficient_resources    = (sufficient_smem && sufficient_threads && !special_skip),
      };
  
      Test<BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, ALGORITHM>(
          gen_mode, scan_op, initial_value, Int2Type<sufficient_resources>());
  }
  
  
  
  /**
   * Run test for different thread block dimensions
   */
  template <
      int                 BLOCK_THREADS,
      int                 ITEMS_PER_THREAD,
      ScanMode            SCAN_MODE,
      TestMode            TEST_MODE,
      BlockScanAlgorithm  ALGORITHM,
      typename            ScanOpT,
      typename            T>
  void Test(
      GenMode     gen_mode,
      ScanOpT     scan_op,
      T           initial_value)
  {
      Test<BLOCK_THREADS, 1, 1, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, ALGORITHM>(gen_mode, scan_op, initial_value);
      Test<BLOCK_THREADS, 2, 2, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, ALGORITHM>(gen_mode, scan_op, initial_value);
  }
  
  
  /**
   * Run test for different policy types
   */
  template <
      int         BLOCK_THREADS,
      int         ITEMS_PER_THREAD,
      ScanMode    SCAN_MODE,
      TestMode    TEST_MODE,
      typename    ScanOpT,
      typename    T>
  void Test(
      GenMode     gen_mode,
      ScanOpT     scan_op,
      T           initial_value)
  {
  #ifdef TEST_RAKING
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, BLOCK_SCAN_RAKING>(gen_mode, scan_op, initial_value);
  #endif
  #ifdef TEST_RAKING_MEMOIZE
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, BLOCK_SCAN_RAKING_MEMOIZE>(gen_mode, scan_op, initial_value);
  #endif
  #ifdef TEST_WARP_SCANS
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, SCAN_MODE, TEST_MODE, BLOCK_SCAN_WARP_SCANS>(gen_mode, scan_op, initial_value);
  #endif
  }
  
  
  /**
   * Run tests for different primitive variants
   */
  template <
      int         BLOCK_THREADS,
      int         ITEMS_PER_THREAD,
      typename    ScanOpT,
      typename    T>
  void Test(
      GenMode     gen_mode,
      ScanOpT     scan_op,
      T           identity,
      T           initial_value)
  {
      // Exclusive (use identity as initial value because it will dispatch to *Sum variants that don't take initial values)
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, BASIC>(gen_mode, scan_op, identity);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, AGGREGATE>(gen_mode, scan_op, identity);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, PREFIX>(gen_mode, scan_op, identity);
  
      // Exclusive (non-specialized, so we can use initial-value)
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, BASIC>(gen_mode, WrapperFunctor<ScanOpT>(scan_op), initial_value);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, AGGREGATE>(gen_mode, WrapperFunctor<ScanOpT>(scan_op), initial_value);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, EXCLUSIVE, PREFIX>(gen_mode, WrapperFunctor<ScanOpT>(scan_op), initial_value);
  
      // Inclusive
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, INCLUSIVE, BASIC>(gen_mode, scan_op, identity);      // This scan doesn't take an initial value
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, INCLUSIVE, AGGREGATE>(gen_mode, scan_op, identity);  // This scan doesn't take an initial value
      Test<BLOCK_THREADS, ITEMS_PER_THREAD, INCLUSIVE, PREFIX>(gen_mode, scan_op, initial_value);
  }
  
  
  /**
   * Run tests for different problem-generation options
   */
  template <
      int         BLOCK_THREADS,
      int         ITEMS_PER_THREAD,
      typename    ScanOpT,
      typename    T>
  void Test(
      ScanOpT     scan_op,
      T           identity,
      T           initial_value)
  {
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(UNIFORM, scan_op, identity, initial_value);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(INTEGER_SEED, scan_op, identity, initial_value);
  
      // Don't test randomly-generated floats b/c of stability
      if (Traits<T>::CATEGORY != FLOATING_POINT)
          Test<BLOCK_THREADS, ITEMS_PER_THREAD>(RANDOM, scan_op, identity, initial_value);
  }
  
  
  /**
   * Run tests for different data types and scan ops
   */
  template <
      int BLOCK_THREADS,
      int ITEMS_PER_THREAD>
  void Test()
  {
      // Get ptx version
      int ptx_version;
      CubDebugExit(PtxVersion(ptx_version));
  
      // primitive
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), (unsigned char) 0, (unsigned char) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), (unsigned short) 0, (unsigned short) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), (unsigned int) 0, (unsigned int) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), (unsigned long long) 0, (unsigned long long) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), (float) 0, (float) 99);
  
      // primitive (alternative scan op)
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Max(), std::numeric_limits<char>::min(), (char) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Max(), std::numeric_limits<short>::min(), (short) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Max(), std::numeric_limits<int>::min(), (int) 99);
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Max(), std::numeric_limits<long long>::min(), (long long) 99);
  
      if (ptx_version > 120)                          // Don't check doubles on PTX120 or below because they're down-converted
          Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Max(), std::numeric_limits<double>::max() * -1, (double) 99);
  
      // vec-1
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_uchar1(0), make_uchar1(17));
  
      // vec-2
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_uchar2(0, 0), make_uchar2(17, 21));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_ushort2(0, 0), make_ushort2(17, 21));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_uint2(0, 0), make_uint2(17, 21));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_ulonglong2(0, 0), make_ulonglong2(17, 21));
  
      // vec-4
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_char4(0, 0, 0, 0), make_char4(17, 21, 32, 85));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_short4(0, 0, 0, 0), make_short4(17, 21, 32, 85));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_int4(0, 0, 0, 0), make_int4(17, 21, 32, 85));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), make_longlong4(0, 0, 0, 0), make_longlong4(17, 21, 32, 85));
  
      // complex
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), TestFoo::MakeTestFoo(0, 0, 0, 0), TestFoo::MakeTestFoo(17, 21, 32, 85));
      Test<BLOCK_THREADS, ITEMS_PER_THREAD>(Sum(), TestBar(0, 0), TestBar(17, 21));
  
  }
  
  
  /**
   * Run tests for different items per thread
   */
  template <int BLOCK_THREADS>
  void Test()
  {
      Test<BLOCK_THREADS, 1>();
      Test<BLOCK_THREADS, 2>();
      Test<BLOCK_THREADS, 9>();
  }
  
  
  
  /**
   * Main
   */
  int main(int argc, char** argv)
  {
      // Initialize command line
      CommandLineArgs args(argc, argv);
      g_verbose = args.CheckCmdLineFlag("v");
      args.GetCmdLineArgument("repeat", g_repeat);
  
      // Print usage
      if (args.CheckCmdLineFlag("help"))
      {
          printf("%s "
              "[--device=<device-id>] "
              "[--repeat=<repetitions of entire test suite>]"
              "[--v] "
              "
  ", argv[0]);
          exit(0);
      }
  
      // Initialize device
      CubDebugExit(args.DeviceInit());
  
  #ifdef QUICK_TEST
  
      Test<128, 1, 1, 1, EXCLUSIVE, AGGREGATE, BLOCK_SCAN_WARP_SCANS>(UNIFORM, Sum(), int(0));
  
      // Compile/run quick tests
      Test<128, 1, 1, 4, EXCLUSIVE, AGGREGATE, BLOCK_SCAN_WARP_SCANS>(UNIFORM, Sum(), int(0));
      Test<128, 1, 1, 4, EXCLUSIVE, AGGREGATE, BLOCK_SCAN_RAKING>(UNIFORM, Sum(), int(0));
      Test<128, 1, 1, 4, EXCLUSIVE, AGGREGATE, BLOCK_SCAN_RAKING_MEMOIZE>(UNIFORM, Sum(), int(0));
  
      Test<128, 1, 1, 2, INCLUSIVE, PREFIX, BLOCK_SCAN_RAKING>(INTEGER_SEED, Sum(), TestFoo::MakeTestFoo(17, 21, 32, 85));
      Test<128, 1, 1, 1, EXCLUSIVE, AGGREGATE, BLOCK_SCAN_WARP_SCANS>(UNIFORM, Sum(), make_longlong4(17, 21, 32, 85));
  
  
  #else
  
      // Compile/run thorough tests
      for (int i = 0; i <= g_repeat; ++i)
      {
          // Run tests for different thread block sizes
          Test<17>();
          Test<32>();
          Test<62>();
          Test<65>();
  //            Test<96>();             // TODO: file bug for UNREACHABLE error for Test<96, 9, BASIC, BLOCK_SCAN_RAKING>(UNIFORM, Sum(), NullType(), make_ulonglong2(17, 21));
          Test<128>();
      }
  
  #endif
  
      return 0;
  }