Add doxygen (#33)

* Warnings as errors

* Doxygen

* Use conda in CI

.gitignore

@@ -1 +1,7 @@
-build/
+build/
+*.pkl
+*.json
+*.pdf
+*.csv
+*.png
+*.pyc

CMakeLists.txt

@@ -18,7 +18,7 @@ add_library(GPUTreeShap::GPUTreeShap ALIAS GPUTreeShap)
 target_sources(GPUTreeShap INTERFACE ${GPUTreeShap_SOURCE_DIR}/GPUTreeShap/gpu_treeshap.h)
 target_include_directories(GPUTreeShap INTERFACE ${GPUTreeShap_SOURCE_DIR})
 
-set(COMPILE_OPTIONS --expt-extended-lambda -lineinfo)
+set(COMPILE_OPTIONS --expt-extended-lambda -lineinfo --Werror all-warnings)
 set(GCC_COMPILE_OPTIONS -Xcompiler -Werror,-Wall,-Wextra)
 if(BUILD_GTEST)
   set(gtest_force_shared_crt ON CACHE BOOL "Always use msvcrt.dll")
@@ -60,5 +60,13 @@ if(BUILD_BENCHMARKS)
   add_executable(BenchmarkGPUTreeShap benchmark/benchmark_gpu_treeshap.cu)
   target_link_libraries(BenchmarkGPUTreeShap PRIVATE GPUTreeShap)
   target_compile_options(BenchmarkGPUTreeShap PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: ${COMPILE_OPTIONS}>)
+  if(NOT MSVC)
+    target_compile_options(BenchmarkGPUTreeShap PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:${GCC_COMPILE_OPTIONS}>)
+  endif()
   target_link_libraries(BenchmarkGPUTreeShap  PRIVATE benchmark::benchmark)
 endif()
+
+include(cmake/doxygen.cmake)
+add_doxygen_target(IN_DOXYFILE Doxyfile.in
+  OUT_DOXYFILE ${CMAKE_CURRENT_BINARY_DIR}/Doxyfile
+  CWD ${CMAKE_CURRENT_BINARY_DIR})

GPUTreeShap/gpu_treeshap.h

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -161,6 +161,7 @@ __host__ __device__ inline size_t IndexPhiInteractions(size_t row_idx,
   return matrix_offset + i * (num_columns + 1) + j;
 }
 
+#ifndef DOXYGEN_SHOULD_SKIP_THIS
 namespace detail {
 
 // Shorthand for creating a device vector with an appropriate allocator type
@@ -1206,18 +1207,29 @@ void ComputeBias(const PathVectorT& device_paths, DoubleVectorT* bias) {
 
 };  // namespace detail
 
+#endif /* DOXYGEN_SHOULD_SKIP_THIS */
+
+/** \defgroup GPUTreeShap
+ *  @{
+ */
+
+
 /*!
  * Compute feature contributions on the GPU given a set of unique paths through
  * a tree ensemble and a dataset. Uses device memory proportional to the tree
  * ensemble size.
  *
  * \exception std::invalid_argument Thrown when an invalid argument error
- * condition occurs. \tparam  PathIteratorT     Thrust type iterator, may be
+ * condition occurs. 
+ * \tparam  PathIteratorT     Thrust type iterator, may be
  * thrust::device_ptr for device memory, or stl iterator/raw pointer for host
- * memory. \tparam  PhiIteratorT      Thrust type iterator, may be
+ * memory. 
+ * \tparam  PhiIteratorT      Thrust type iterator, may be
  * thrust::device_ptr for device memory, or stl iterator/raw pointer for host
- * memory. Value type must be floating point. \tparam  DatasetT User-specified
- * dataset container. \tparam  DeviceAllocatorT  Optional thrust style
+ * memory. Value type must be floating point. 
+ * \tparam  DatasetT User-specified
+ * dataset container. 
+ * \tparam  DeviceAllocatorT  Optional thrust style
  * allocator.
  *
  * \param X           Thin wrapper over a dataset allocated in device memory. X
@@ -1230,10 +1242,14 @@ void ComputeBias(const PathVectorT& device_paths, DoubleVectorT* bias) {
  * root with feature_idx = -1 and zero_fraction = 1.0. The ordering of path
  * elements inside a unique path does not matter - the result will be the same.
  * Paths may contain duplicate features. See the PathElement class for more
- * information. \param end         Path end iterator. \param num_groups  Number
+ * information. 
+ * \param end         Path end iterator. 
+ * \param num_groups  Number
  * of output groups. In multiclass classification the algorithm outputs feature
- * contributions per output class. \param phis_begin  Begin iterator for output
- * phis. \param phis_end    End iterator for output phis.
+ * contributions per output class. 
+ * \param phis_begin  Begin iterator for output
+ * phis. 
+ * \param phis_end    End iterator for output phis.
  */
 template <typename DeviceAllocatorT = thrust::device_allocator<int>,
           typename DatasetT, typename PathIteratorT, typename PhiIteratorT>
@@ -1526,4 +1542,7 @@ void GPUTreeShapInterventional(DatasetT X, DatasetT R, PathIteratorT begin,
                                     temp_phi.data().get());
   thrust::copy(temp_phi.begin(), temp_phi.end(), phis_begin);
 }
+
+/** @}*/
+
 }  // namespace gpu_treeshap

benchmark/benchmark_gpu_treeshap.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,11 +17,12 @@
 #include <benchmark/benchmark.h>
 #include "../tests/test_utils.h"
 
-using namespace gpu_treeshap;  // NOLINT
+
+using namespace gpu_treeshap; // NOLINT
 
 class Fixture : public benchmark::Fixture {
  public:
-  void SetUp(const ::benchmark::State& state) override {
+  void SetUp(const ::benchmark::State &state) override {
     num_groups = 5;
     num_rows = state.range(0);
     num_features = state.range(1);
@@ -36,10 +37,19 @@ class Fixture : public benchmark::Fixture {
     phis.reset(new thrust::device_vector<float>(
         X.NumRows() * (X.NumCols() + 1) * num_groups));
   }
-  void TearDown(const ::benchmark::State& state) {
+  void TearDown([[maybe_unused]] const ::benchmark::State &state) {
     phis.reset();
     test_data.reset();
   }
+  // to keep compiler happy
+  void SetUp(::benchmark::State &st) override {
+    SetUp(const_cast<const ::benchmark::State &>(st));
+  }
+
+  // to keep compiler happy
+  void TearDown(::benchmark::State &st) override {
+    TearDown(const_cast<const ::benchmark::State &>(st));
+  }
   std::vector<PathElement<XgboostSplitCondition>> model;
   std::unique_ptr<TestDataset> test_data;
   DenseDatasetWrapper X;
@@ -49,7 +59,7 @@ class Fixture : public benchmark::Fixture {
   size_t num_features;
 };
 
-BENCHMARK_DEFINE_F(Fixture, GPUTreeShap)(benchmark::State& st) { // NOLINT
+BENCHMARK_DEFINE_F(Fixture, GPUTreeShap)(benchmark::State &st) { // NOLINT
   for (auto _ : st) {
     GPUTreeShap(X, model.begin(), model.end(), num_groups, phis->begin(),
                 phis->end());
@@ -62,7 +72,7 @@ BENCHMARK_REGISTER_F(Fixture, GPUTreeShap)
     ->Args({100000, 500, 20, 10000});
 
 BENCHMARK_DEFINE_F(Fixture, GPUTreeShapInterventional)
-(benchmark::State& st) {  // NOLINT
+(benchmark::State &st) { // NOLINT
   TestDataset R_test_data(1000, num_features, 1429);
   DenseDatasetWrapper R = R_test_data.GetDeviceWrapper();
   for (auto _ : st) {
@@ -75,7 +85,8 @@ BENCHMARK_REGISTER_F(Fixture, GPUTreeShapInterventional)
     ->Args({1000, 10, 6, 1000})
     ->Args({10000, 50, 10, 1000});
 
-BENCHMARK_DEFINE_F(Fixture, GPUTreeShapInteractions)(benchmark::State& st) {// NOLINT
+BENCHMARK_DEFINE_F(Fixture, GPUTreeShapInteractions)
+(benchmark::State &st) { // NOLINT
   phis.reset(new thrust::device_vector<float>(X.NumRows() * (X.NumCols() + 1) *
                                               (X.NumCols() + 1) * num_groups));
   for (auto _ : st) {
@@ -91,7 +102,7 @@ BENCHMARK_REGISTER_F(Fixture, GPUTreeShapInteractions)
     ->Args({1000, 250, 20, 10000});
 
 BENCHMARK_DEFINE_F(Fixture, GPUTreeShapTaylorInteractions)
-(benchmark::State& st) {// NOLINT
+(benchmark::State &st) { // NOLINT
   phis.reset(new thrust::device_vector<float>(X.NumRows() * (X.NumCols() + 1) *
                                               (X.NumCols() + 1) * num_groups));
   for (auto _ : st) {
@@ -110,13 +121,13 @@ std::vector<int> GenerateCounts(size_t n, size_t max_depth) {
   std::mt19937 gen(95);
   std::uniform_int_distribution<int> distrib(0, max_depth - 1);
   std::vector<int> out(n);
-  for (auto& x : out) {
+  for (auto &x : out) {
     x = distrib(gen);
   }
   return out;
 }
 
-static void BFDBinPacking(benchmark::State& state) {// NOLINT
+static void BFDBinPacking(benchmark::State &state) { // NOLINT
   size_t n = state.range(0);
   size_t max_depth = state.range(1);
   thrust::device_vector<int> counts = GenerateCounts(n, max_depth);

build.sh

@@ -0,0 +1,6 @@
+#!/bin/bash
+mkdir -p build
+cd build
+cmake .. -DCMAKE_BUILD_TYPE=Release -DBUILD_GTEST=ON -DBUILD_EXAMPLES=ON -DBUILD_BENCHMARKS=ON
+make -j4
+cmake --build . --target docs_cuml

ci/gpu/build.sh

@@ -1,5 +1,5 @@
 #!/bin/bash
-# Copyright (c) 2020, NVIDIA CORPORATION.
+# Copyright (c) 2022, NVIDIA CORPORATION.
 #########################################
 # GPUTreeShap GPU build and test script for CI #
 #########################################
@@ -25,23 +25,6 @@ cd ~
 # SETUP - Check environment
 ################################################################################
 
-gpuci_logger "Install cmake"
-mkdir cmake
-cd cmake
-wget https://github.com/Kitware/CMake/releases/download/v3.18.2/cmake-3.18.2-Linux-x86_64.sh
-sh cmake-3.18.2-Linux-x86_64.sh --skip-license
-export PATH=$PATH:$PWD/bin
-cd ..
-
-gpuci_logger "Install gtest"
-wget https://github.com/google/googletest/archive/release-1.10.0.zip
-unzip release-1.10.0.zip
-mv googletest-release-1.10.0 gtest && cd gtest
-cmake . && make
-cp -r googletest/include include
-export GTEST_ROOT=$PWD
-cd ..
-
 gpuci_logger "Check environment"
 env
 
@@ -52,16 +35,16 @@ nvidia-smi
 $CC --version
 $CXX --version
 
+gpuci_logger "Activate conda env"
+. /opt/conda/etc/profile.d/conda.sh
+conda activate rapids
 
 ################################################################################
 # BUILD - Build tests
 ################################################################################
 
 gpuci_logger "Build C++ targets"
-mkdir $WORKSPACE/build
-cd $WORKSPACE/build
-cmake .. -DBUILD_GTEST=ON -DBUILD_EXAMPLES=ON -DBUILD_BENCHMARKS=ON
-make -j
+./build.sh
 
 ################################################################################
 # TEST - Run GoogleTest

cmake/doxygen.cmake

@@ -0,0 +1,33 @@
+# Copyright (c) 2019-2021, NVIDIA CORPORATION.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+#
+
+find_package(Doxygen 1.8.12)
+
+function(add_doxygen_target)
+  if(Doxygen_FOUND)
+    set(options "")
+    set(oneValueArgs IN_DOXYFILE OUT_DOXYFILE CWD)
+    set(multiValueArgs "")
+    cmake_parse_arguments(dox "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
+    configure_file(${dox_IN_DOXYFILE} ${dox_OUT_DOXYFILE} @ONLY)
+    add_custom_target(docs_cuml
+      ${DOXYGEN_EXECUTABLE} ${dox_OUT_DOXYFILE}
+      WORKING_DIRECTORY ${dox_CWD}
+      VERBATIM
+      COMMENT "Generate doxygen docs")
+  else()
+    message("add_doxygen_target: doxygen exe not found")
+  endif()
+endfunction(add_doxygen_target)

tests/test_utils.h

@@ -32,7 +32,6 @@ class DenseDatasetWrapper {
   DenseDatasetWrapper(const float* data, int num_rows, int num_cols)
       : data(data), num_rows(num_rows), num_cols(num_cols) {}
   __device__ float GetElement(size_t row_idx, size_t col_idx) const {
-    assert(col_idx >= 0);
     return data[row_idx * num_cols + col_idx];
   }
   __host__ __device__ size_t NumRows() const { return num_rows; }
@@ -90,12 +89,13 @@ void GenerateModel(std::vector<PathElement<SplitConditionT>>* model, int group,
         upper_bound = 1.0f - bound_dis(*gen);
       }
       // Don't make the zero fraction too small
+      SplitConditionT split(lower_bound, upper_bound, bern_dis(*gen));
       std::uniform_real_distribution<float> zero_fraction_dis(0.05, 1.0);
       model->emplace_back(PathElement<SplitConditionT>{
           base_path_idx + i,
           feature_dis(*gen),
           group,
-          {lower_bound, upper_bound, bern_dis(*gen)},
+           split,
           zero_fraction_dis(*gen),
           v});
     }