Draft for catch2 testing framework usage

[gevtushenko]

Cub uses testing facilities that are far from perfect. This PR contains Catch2 integration and a few convenience wrappers. Main advantages of Catch2:

• Readable way of specifying cartesian products of types/compile-time parameters. Instead of having

  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>();
  }
  
  template <
      typename                    SampleT,
      int                         BINS,
      int                         BLOCK_THREADS>
  void Test()
  {
      Test<SampleT, BINS, BLOCK_THREADS, 1>();
      Test<SampleT, BINS, BLOCK_THREADS, 5>();
  }
  
  
  template <
      typename                    SampleT,
      int                         BINS>
  void Test()
  {
      Test<SampleT, BINS, 32>();
      Test<SampleT, BINS, 96>();
      Test<SampleT, BINS, 128>();
  }
  

  Catch2 allows writing:

  using types = catch2_helper::type_list<unsigned char, unsigned short>;
  using threads_in_block = catch2_helper::enum_type_list<int, 32, 256, 1024>;
  using items_per_thread = catch2_helper::enum_type_list<int, 1, 5>;
  using bins = catch2_helper::enum_type_list<int, 32, 96, 128>;
  using algorithms = catch2_helper::enum_type_list<BlockHistogramAlgorithm,
                                                   BLOCK_HISTO_SORT,
                                                   BLOCK_HISTO_ATOMIC>;
  
  using cases = catch2_helper::cartesian_product<
    catch2_helper::type_list<
      types,
      items_per_thread,
      threads_in_block,
      bins,
      algorithms>>;
  
  TEMPLATE_LIST_TEST_CASE("Block histogram can be computed with uniform input",
                          "[histogram][block]",
                          cases)
  

• Readable generators. Instead of traversing parameter space in multiple nested loops:

  void Test(thrust::default_random_engine &rng)
  {
    for (unsigned int num_items = ItemsPerThread * ThreadsInBlock;
         num_items > 1;
         num_items /= 2)
    {
  

  Catch2 provides readable generator expressions:

    const int num_keys = GENERATE_COPY(take(10, random(0, ItemsPerThread * ThreadsInBlock)));
  

• Display of variables' content in case of failure. Currently, only the place of failure is displayed. Catch2 displays the content of compared variables:

  thrust::host_vector<KeyT> h_reference = /* ... */;
  REQUIRE( h_reference == d_keys );
  

  ctest_block_merge_sort.cu(180): FAILED:
  REQUIRE( h_reference == d_keys ) with expansion:
  { 120, -2147483648, -2147483648, -2147483648 }
  ==
  { -2147483648, -2147483648, -2147483648, -2147483648 }
  

• Breaking into the debugger. Catch2 automatically creates a breakpoint in case of failure. It's not the case with the current approach.
• Ability to execute only part of unit-tests in a file. It's also possible to perform tests by wildcard or tags. I aggregate all files that start with ctest_ into a single test (per configuration). So it'll be possible to perform all the histogram tests (across scopes) or all of the block-scope tests depending on the command line parameters.

It's not the full list of features. I'd like to discuss this draft in case there are any corner cases I'm missing. If there's a better unit-testing framework I'd also like to discuss it here.

[jrhemstad]

Huge +1 from me. I experimented with using Catch2 in cuCollections and I have loved it.

I know CUB doesn't currently use GTest, but many of my favorite features are in comparison to what you can do with GTest:

• Compared to GTest, it cuts out a lot of the boilerplate of writing test fixture clases. The independent, depth-first execution of nested of SECTIONs makes it natural and intuitive to reuse common code but maintain independent execution paths.
• This may seem silly, but I also really love that you can include plain text description of what the intent of a particular test/section is, e.g.:

TEST_CASE( "vectors can be sized and resized", "[vector]" ) {
    std::vector<int> v( 5 );

    REQUIRE( v.size() == 5 );
    REQUIRE( v.capacity() >= 5 );

    SECTION( "resizing bigger changes size and capacity" ) {
        v.resize( 10 );

        REQUIRE( v.size() == 10 );
        REQUIRE( v.capacity() >= 10 );
    }
}

• Another big advantage compared to Gtest is you can use __device__ lambdas in the body of a catch2 test. With GTest where a test section is expanded into a private member function of a class. nvcc forbids a __device__ lambda in such a function.

[gevtushenko]

Huge +1 from me. I experimented with using Catch2 in cuCollections and I have loved it.

I know CUB doesn't currently use GTest, but many of my favorite features are in comparison to what you can do with GTest:

• Compared to GTest, it cuts out a lot of the boilerplate of writing test fixture clases. The independent, depth-first execution of nested of SECTIONs makes it natural and intuitive to reuse common code but maintain independent execution paths.
• This may seem silly, but I also really love that you can include plain text description of what the intent of a particular test/section is, e.g.:

TEST_CASE( "vectors can be sized and resized", "[vector]" ) {
    std::vector<int> v( 5 );

    REQUIRE( v.size() == 5 );
    REQUIRE( v.capacity() >= 5 );

    SECTION( "resizing bigger changes size and capacity" ) {
        v.resize( 10 );

        REQUIRE( v.size() == 10 );
        REQUIRE( v.capacity() >= 10 );
    }
}

• Another big advantage compared to Gtest is you can use __device__ lambdas in the body of a catch2 test. With GTest where a test section is expanded into a private member function of a class. nvcc forbids a __device__ lambda in such a function.

Initially, I've integrated GTest. Both generators and sections convinced me that it's better to use Catch2 😄 Have you encountered an issue related to Catch2 usage in .cu files? Currently I have to suppress the following warning to get it working:

#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC
#pragma diag_suppress 177
#endif

[jrhemstad]

Have you encountered an issue related to Catch2 usage in .cu files?

All the test files in cuCollections are .cu files: https://github.com/NVIDIA/cuCollections/tree/dev/tests

The is one warning that's generated when using a __device__ lambda in a test due to the fact that a Catch2 test expands out to a function in an anonymous namespace and the capture generated by nvcc for __device__ lambda takes the address of the enclosing function and uses it as a template parameter. So it complains about using the address of an anonymous function as part of the type signature.

What warning are you seeing?

[gevtushenko]

Have you encountered an issue related to Catch2 usage in .cu files?

All the test files in cuCollections are .cu files: https://github.com/NVIDIA/cuCollections/tree/dev/tests

The is one warning that's generated when using a __device__ lambda in a test due to the fact that a Catch2 test expands out to a function in an anonymous namespace and the capture generated by nvcc for __device__ lambda takes the address of the enclosing function and uses it as a template parameter. So it complains about using the address of an anonymous function as part of the type signature.

What warning are you seeing?

I get a bunch of unused variables warnings:

catch.hpp(15860): error #177-D: variable "colour" was declared but never referenced

I get it with MSVC 16.11.1 + C++17. This issue appears only in .cu files.

[gevtushenko]

gpuCI: NVIDIA/thrust#1559