/******************************************************************************
   * 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 BlockHistogram utilities
   ******************************************************************************/
  
  // Ensure printing of CUDA runtime errors to console
  #define CUB_STDERR
  
  #include <stdio.h>
  #include <limits>
  #include <string>
  #include <typeinfo>
  
  #include <cub/block/block_histogram.cuh>
  #include <cub/block/block_load.cuh>
  #include <cub/block/block_store.cuh>
  #include <cub/util_allocator.cuh>
  
  #include "test_util.h"
  
  using namespace cub;
  
  
  //---------------------------------------------------------------------
  // Globals, constants and typedefs
  //---------------------------------------------------------------------
  
  bool                    g_verbose           = false;
  int                     g_timing_iterations = 0;
  int                     g_repeat            = 0;
  CachingDeviceAllocator  g_allocator(true);
  
  
  //---------------------------------------------------------------------
  // Test kernels
  //---------------------------------------------------------------------
  
  /**
   * BlockHistogram test kernel.
   */
  template <
      int                     BINS,
      int                     BLOCK_THREADS,
      int                     ITEMS_PER_THREAD,
      BlockHistogramAlgorithm ALGORITHM,
      typename                T,
      typename                HistoCounter>
  __global__ void BlockHistogramKernel(
      T                       *d_samples,
      HistoCounter            *d_histogram)
  {
      // Parameterize BlockHistogram type for our thread block
      typedef BlockHistogram<T, BLOCK_THREADS, ITEMS_PER_THREAD, BINS, ALGORITHM> BlockHistogram;
  
      // Allocate temp storage in shared memory
      __shared__ typename BlockHistogram::TempStorage temp_storage;
  
      // Per-thread tile data
      T data[ITEMS_PER_THREAD];
      LoadDirectStriped<BLOCK_THREADS>(threadIdx.x, d_samples, data);
  
      // Test histo (writing directly to histogram buffer in global)
      BlockHistogram(temp_storage).Histogram(data, d_histogram);
  }
  
  
  /**
   * Initialize problem (and solution)
   */
  template <
      int             BINS,
      typename        SampleT>
  void Initialize(
      GenMode         gen_mode,
      SampleT         *h_samples,
      int             *h_histograms_linear,
      int             num_samples)
  {
      // Init bins
      for (int bin = 0; bin < BINS; ++bin)
      {
          h_histograms_linear[bin] = 0;
      }
  
      if (g_verbose) printf("Samples: 
  ");
  
      // Initialize interleaved channel samples and histogram them correspondingly
      for (int i = 0; i < num_samples; ++i)
      {
          InitValue(gen_mode, h_samples[i], i);
          h_samples[i] %= BINS;
  
          if (g_verbose) std::cout << CoutCast(h_samples[i]) << ", ";
  
          h_histograms_linear[h_samples[i]]++;
      }
  
      if (g_verbose) printf("
  
  ");
  }
  
  
  /**
   * Test BlockHistogram
   */
  template <
      typename                    SampleT,
      int                         BINS,
      int                         BLOCK_THREADS,
      int                         ITEMS_PER_THREAD,
      BlockHistogramAlgorithm     ALGORITHM>
  void Test(
      GenMode                     gen_mode)
  {
      int num_samples = BLOCK_THREADS * ITEMS_PER_THREAD;
  
      printf("cub::BlockHistogram %s %d %s samples (%dB), %d bins, %d threads, gen-mode %s
  ",
          (ALGORITHM == BLOCK_HISTO_SORT) ? "BLOCK_HISTO_SORT" : "BLOCK_HISTO_ATOMIC",
          num_samples,
          typeid(SampleT).name(),
          (int) sizeof(SampleT),
          BINS,
          BLOCK_THREADS,
          (gen_mode == RANDOM) ? "RANDOM" : (gen_mode == INTEGER_SEED) ? "SEQUENTIAL" : "HOMOGENOUS");
      fflush(stdout);
  
      // Allocate host arrays
      SampleT         *h_samples          = new SampleT[num_samples];
      int   *h_reference = new int[BINS];
  
      // Initialize problem
      Initialize<BINS>(gen_mode, h_samples, h_reference, num_samples);
  
      // Allocate problem device arrays
      SampleT         *d_samples = NULL;
      int             *d_histogram = NULL;
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_samples,             sizeof(SampleT) * num_samples));
      CubDebugExit(g_allocator.DeviceAllocate((void**)&d_histogram,   sizeof(int) * BINS));
  
      // Initialize/clear device arrays
      CubDebugExit(cudaMemcpy(d_samples, h_samples, sizeof(SampleT) * num_samples, cudaMemcpyHostToDevice));
      CubDebugExit(cudaMemset(d_histogram, 0, sizeof(int) * BINS));
  
      // Run kernel
      BlockHistogramKernel<BINS, BLOCK_THREADS, ITEMS_PER_THREAD, ALGORITHM><<<1, BLOCK_THREADS>>>(
          d_samples,
          d_histogram);
  
      // Check for correctness (and display results, if specified)
      int compare = CompareDeviceResults((int*) h_reference, d_histogram, BINS, g_verbose, g_verbose);
      printf("\t%s
  
  ", compare ? "FAIL" : "PASS");
  
      // Flush any stdout/stderr
      CubDebugExit(cudaPeekAtLastError());
      CubDebugExit(cudaDeviceSynchronize());
      fflush(stdout);
      fflush(stderr);
  
      // Cleanup
      if (h_samples) delete[] h_samples;
      if (h_reference) delete[] h_reference;
      if (d_samples) CubDebugExit(g_allocator.DeviceFree(d_samples));
      if (d_histogram) CubDebugExit(g_allocator.DeviceFree(d_histogram));
  
      // Correctness asserts
      AssertEquals(0, compare);
  }
  
  
  /**
   * Test different sample distributions
   */
  template <
      typename                    SampleT,
      int                         BINS,
      int                         BLOCK_THREADS,
      int                         ITEMS_PER_THREAD,
      BlockHistogramAlgorithm     ALGORITHM>
  void Test()
  {
      Test<SampleT, BINS, BLOCK_THREADS, ITEMS_PER_THREAD, ALGORITHM>(UNIFORM);
      Test<SampleT, BINS, BLOCK_THREADS, ITEMS_PER_THREAD, ALGORITHM>(INTEGER_SEED);
      Test<SampleT, BINS, BLOCK_THREADS, ITEMS_PER_THREAD, ALGORITHM>(RANDOM);
  }
  
  
  /**
   * Test different ALGORITHM
   */
  template <
      typename                    SampleT,
      int                         BINS,
      int                         BLOCK_THREADS,
      int                         ITEMS_PER_THREAD>
  void Test()
  {
      Test<SampleT, BINS, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_HISTO_SORT>();
      Test<SampleT, BINS, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_HISTO_ATOMIC>();
  }
  
  
  /**
   * Test different ITEMS_PER_THREAD
   */
  template <
      typename                    SampleT,
      int                         BINS,
      int                         BLOCK_THREADS>
  void Test()
  {
      Test<SampleT, BINS, BLOCK_THREADS, 1>();
      Test<SampleT, BINS, BLOCK_THREADS, 5>();
  }
  
  
  /**
   * Test different BLOCK_THREADS
   */
  template <
      typename                    SampleT,
      int                         BINS>
  void Test()
  {
      Test<SampleT, BINS, 32>();
      Test<SampleT, BINS, 96>();
      Test<SampleT, BINS, 128>();
  }
  
  
  
  
  
  //---------------------------------------------------------------------
  // Main
  //---------------------------------------------------------------------
  
  /**
   * 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 "
              "[--n=<total input samples across all channels> "
              "[--device=<device-id>] "
              "[--repeat=<repetitions of entire test suite>]"
              "[--v] "
              "
  ", argv[0]);
          exit(0);
      }
  
      // Initialize device
      CubDebugExit(args.DeviceInit());
  
  #ifdef QUICK_TEST
  
      // Compile/run quick tests
      Test<unsigned char, 256, 128, 4, BLOCK_HISTO_SORT>(RANDOM);
      Test<unsigned char, 256, 128, 4, BLOCK_HISTO_ATOMIC>(RANDOM);
  
  #else
  
      // Compile/run thorough tests
      for (int i = 0; i <= g_repeat; ++i)
      {
          Test<unsigned char, 32>();
          Test<unsigned char, 256>();
          Test<unsigned short, 1024>();
      }
  
  #endif
  
      return 0;
  }