From e977f3e4194313a857da96fdd810ee5ef1d742b3 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 10 Aug 2021 14:06:57 +1000
Subject: [PATCH] Warnings are errors (#4075)

Depends on rapidsai/raft#299
Depends on rapidsai/raft#311

Fixes #4086

This PR enables -Wall (warnings are errors) and fixes all current compiler warnings.

Compiler warnings should almost never be ignored. -Wall helps prevent hard-to-find bugs.

Tested with both Release and Debug builds since more warnings are typically reported in debug builds.

In any file I touched, I also cleaned up the includes so that they are grouped and sorted and ordered from "near" to "far" (relative to the file that is including).

Also cleaned up many instances of `size_t` --> `std::size_t`, however I will leave a global search and replace for a separate PR to make reviewing easier.

Authors:
  - Mark Harris (https://github.com/harrism)

Approvers:
  - Robert Maynard (https://github.com/robertmaynard)
  - William Hicks (https://github.com/wphicks)
  - Dante Gama Dessavre (https://github.com/dantegd)

URL: https://github.com/rapidsai/cuml/pull/4075
---
 cpp/bench/prims/gram_matrix.cu                |   2 +-
 cpp/cmake/modules/ConfigureCUDA.cmake         |   6 +-
 cpp/include/cuml/ensemble/randomforest.hpp    |   5 +
 cpp/include/cuml/random_projection/rproj_c.h  |   3 +-
 cpp/include/cuml/tree/decisiontree.hpp        |   7 +-
 cpp/src/common/allocatorAdapter.hpp           |   9 +-
 cpp/src/common/cumlHandle.hpp                 |   1 +
 cpp/src/common/cuml_api.cpp                   |  10 +-
 cpp/src/common/tensor.hpp                     |   7 +-
 cpp/src/dbscan/adjgraph/algo.cuh              |  14 +-
 cpp/src/dbscan/dbscan.cuh                     |  23 +--
 cpp/src/dbscan/mergelabels/runner.cuh         |   1 +
 cpp/src/dbscan/mergelabels/tree_reduction.cuh |   4 +
 cpp/src/dbscan/runner.cuh                     |  86 +++++-----
 cpp/src/decisiontree/decisiontree.cuh         |  29 ++--
 cpp/src/fil/fil.cu                            |  52 +++---
 cpp/src/fil/infer.cu                          |  56 ++++---
 cpp/src/glm/ols_mg.cu                         |  11 +-
 cpp/src/glm/ridge_mg.cu                       |  23 ++-
 cpp/src/hdbscan/detail/extract.cuh            |  23 ++-
 cpp/src/hdbscan/detail/reachability_faiss.cuh |  25 +--
 cpp/src/holtwinters/internal/hw_eval.cuh      |  10 +-
 cpp/src/holtwinters/internal/hw_optim.cuh     |   9 +-
 cpp/src/kmeans/common.cuh                     |  32 ++--
 cpp/src/knn/knn.cu                            |   6 +-
 cpp/src/knn/knn_opg_common.cuh                |  64 ++++----
 cpp/src/pca/pca_mg.cu                         |  22 +--
 cpp/src/pca/sign_flip_mg.cu                   |  15 +-
 cpp/src/random_projection/rproj.cuh           |  55 ++++---
 cpp/src/random_projection/rproj_utils.cuh     |   6 +-
 cpp/src/randomforest/randomforest.cu          |  22 ++-
 cpp/src/randomforest/randomforest.cuh         |  26 +--
 cpp/src/solver/cd_mg.cu                       |  33 ++--
 cpp/src/solver/lars_impl.cuh                  |   4 +-
 cpp/src/solver/shuffle.h                      |   7 +-
 cpp/src/svm/kernelcache.cuh                   |  15 +-
 cpp/src/svm/smosolver.cuh                     |  36 ++---
 cpp/src/svm/workingset.cuh                    |  48 +++---
 cpp/src/tsvd/tsvd_mg.cu                       |  17 +-
 .../simpl_set_embed/optimize_batch_kernel.cuh |  10 +-
 cpp/src_prims/cache/cache.cuh                 |  13 +-
 cpp/src_prims/cache/cache_util.cuh            |   6 +-
 cpp/src_prims/linalg/batched/matrix.cuh       | 148 +++++++++---------
 cpp/src_prims/linalg/init.h                   |   1 +
 cpp/src_prims/linalg/reduce_rows_by_key.cuh   |  15 +-
 .../metrics/batched/information_criterion.cuh |   7 +-
 cpp/src_prims/metrics/contingencyMatrix.cuh   |  28 ++--
 cpp/src_prims/selection/haversine_knn.cuh     |   4 +-
 cpp/src_prims/selection/knn.cuh               |  50 +++---
 cpp/src_prims/sparse/batched/csr.cuh          |  74 ++++-----
 cpp/src_prims/stats/minmax.cuh                |  25 ++-
 cpp/test/prims/batched/csr.cu                 |  25 +--
 .../prims/batched/information_criterion.cu    |  16 +-
 cpp/test/prims/batched/matrix.cu              |  30 ++--
 cpp/test/prims/rand_index.cu                  |  20 ++-
 cpp/test/sg/fil_test.cu                       |  23 ++-
 cpp/test/sg/multi_sum_test.cu                 |  18 ++-
 cpp/test/sg/rf_test.cu                        |  30 ++--
 cpp/test/sg/rf_treelite_test.cu               |  22 ++-
 cpp/test/sg/shap_kernel.cu                    |  11 +-
 cpp/test/sg/svc_test.cu                       |  67 ++++----
 cpp/test/sg/umap_parametrizable_test.cu       |  42 ++---
 62 files changed, 845 insertions(+), 664 deletions(-)

diff --git a/cpp/bench/prims/gram_matrix.cu b/cpp/bench/prims/gram_matrix.cu
index 5ce6a21032..c561a875c2 100644
--- a/cpp/bench/prims/gram_matrix.cu
+++ b/cpp/bench/prims/gram_matrix.cu
@@ -58,7 +58,7 @@ struct GramMatrix : public Fixture {
       std::unique_ptr<GramMatrixBase<T>>(KernelFactory<T>::create(p.kernel_params, cublas_handle));
   }
 
-  ~GramMatrix() { CUBLAS_CHECK(cublasDestroy(cublas_handle)); }
+  ~GramMatrix() { CUBLAS_CHECK_NO_THROW(cublasDestroy(cublas_handle)); }
 
  protected:
   void allocateBuffers(const ::benchmark::State& state) override
diff --git a/cpp/cmake/modules/ConfigureCUDA.cmake b/cpp/cmake/modules/ConfigureCUDA.cmake
index 0d43f702d2..8dbab5fbfc 100644
--- a/cpp/cmake/modules/ConfigureCUDA.cmake
+++ b/cpp/cmake/modules/ConfigureCUDA.cmake
@@ -25,8 +25,10 @@ endif()
 list(APPEND CUML_CUDA_FLAGS --expt-extended-lambda --expt-relaxed-constexpr)
 
 # set warnings as errors
-# list(APPEND CUML_CUDA_FLAGS -Werror=cross-execution-space-call)
-# list(APPEND CUML_CUDA_FLAGS -Xcompiler=-Wall,-Werror,-Wno-error=deprecated-declarations)
+if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 11.2.0)
+    list(APPEND CUML_CUDA_FLAGS -Werror=all-warnings)
+endif()
+list(APPEND CUML_CUDA_FLAGS -Xcompiler=-Wall,-Werror,-Wno-error=deprecated-declarations)
 
 if(DISABLE_DEPRECATION_WARNING)
     list(APPEND CUML_CXX_FLAGS -Wno-deprecated-declarations)
diff --git a/cpp/include/cuml/ensemble/randomforest.hpp b/cpp/include/cuml/ensemble/randomforest.hpp
index 20a95d1c1b..4f82396dfc 100644
--- a/cpp/include/cuml/ensemble/randomforest.hpp
+++ b/cpp/include/cuml/ensemble/randomforest.hpp
@@ -19,8 +19,13 @@
 #include <cuml/common/logger.hpp>
 #include <cuml/ensemble/treelite_defs.hpp>
 #include <cuml/tree/decisiontree.hpp>
+
 #include <map>
 
+namespace raft {
+class handle_t;  // forward decl
+}
+
 namespace ML {
 
 enum RF_type {
diff --git a/cpp/include/cuml/random_projection/rproj_c.h b/cpp/include/cuml/random_projection/rproj_c.h
index 20a8d2cc62..d4f1702b54 100644
--- a/cpp/include/cuml/random_projection/rproj_c.h
+++ b/cpp/include/cuml/random_projection/rproj_c.h
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuml/common/device_buffer.hpp>
+
 #include <raft/handle.hpp>
 #include <raft/mr/device/allocator.hpp>
 
@@ -95,4 +96,4 @@ void RPROJtransform(const raft::handle_t& handle,
 
 size_t johnson_lindenstrauss_min_dim(size_t n_samples, double eps);
 
-}  // namespace ML
\ No newline at end of file
+}  // namespace ML
diff --git a/cpp/include/cuml/tree/decisiontree.hpp b/cpp/include/cuml/tree/decisiontree.hpp
index 54020c45ec..827859f816 100644
--- a/cpp/include/cuml/tree/decisiontree.hpp
+++ b/cpp/include/cuml/tree/decisiontree.hpp
@@ -15,13 +15,12 @@
  */
 
 #pragma once
-#include <vector>
+
 #include "algo_helper.h"
 #include "flatnode.h"
 
-namespace raft {
-class handle_t;
-}
+#include <string>
+#include <vector>
 
 namespace ML {
 
diff --git a/cpp/src/common/allocatorAdapter.hpp b/cpp/src/common/allocatorAdapter.hpp
index f0f41d9d28..ca3f5c2cc5 100644
--- a/cpp/src/common/allocatorAdapter.hpp
+++ b/cpp/src/common/allocatorAdapter.hpp
@@ -16,14 +16,15 @@
 
 #pragma once
 
-#include <limits>
-
-#include <thrust/system/cuda/execution_policy.h>
-
 #include <raft/cudart_utils.h>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/mr/host/allocator.hpp>
 
+#include <thrust/system/cuda/execution_policy.h>
+
+#include <cstddef>
+#include <limits>
+
 namespace ML {
 
 template <typename T>
diff --git a/cpp/src/common/cumlHandle.hpp b/cpp/src/common/cumlHandle.hpp
index 4120df9e41..4b0a4793fc 100644
--- a/cpp/src/common/cumlHandle.hpp
+++ b/cpp/src/common/cumlHandle.hpp
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuml/cuml_api.h>
+
 #include <raft/handle.hpp>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/mr/host/allocator.hpp>
diff --git a/cpp/src/common/cuml_api.cpp b/cpp/src/common/cuml_api.cpp
index e5fe9a6646..cca2793bca 100644
--- a/cpp/src/common/cuml_api.cpp
+++ b/cpp/src/common/cuml_api.cpp
@@ -14,17 +14,17 @@
  * limitations under the License.
  */
 
+#include "cumlHandle.hpp"
+
 #include <cuml/cuml_api.h>
-#include <raft/mr/device/allocator.hpp>
-#include <raft/mr/host/allocator.hpp>
+#include <cuml/common/utils.hpp>
 
 #include <raft/cudart_utils.h>
-#include <cuml/common/utils.hpp>
-#include <functional>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/mr/host/allocator.hpp>
-#include "cumlHandle.hpp"
 
+#include <cstddef>
+#include <functional>
 namespace ML {
 namespace detail {
 
diff --git a/cpp/src/common/tensor.hpp b/cpp/src/common/tensor.hpp
index 8bb4b17221..8578556199 100644
--- a/cpp/src/common/tensor.hpp
+++ b/cpp/src/common/tensor.hpp
@@ -19,6 +19,7 @@
 #include <raft/cudart_utils.h>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/mr/host/allocator.hpp>
+
 #include <vector>
 
 namespace ML {
@@ -171,7 +172,7 @@ class Tensor {
   std::shared_ptr<raft::mr::host::allocator> _hAllocator;
 
   /// Raw pointer to where the tensor data begins
-  DataPtrT _data;
+  DataPtrT _data{};
 
   /// Array of strides (in sizeof(T) terms) per each dimension
   IndexT _stride[Dim];
@@ -179,9 +180,9 @@ class Tensor {
   /// Size per each dimension
   IndexT _size[Dim];
 
-  AllocState _state;
+  AllocState _state{};
 
-  cudaStream_t _stream;
+  cudaStream_t _stream{};
 };
 
 };  // end namespace ML
diff --git a/cpp/src/dbscan/adjgraph/algo.cuh b/cpp/src/dbscan/adjgraph/algo.cuh
index effafb6c7f..13cbf3eae6 100644
--- a/cpp/src/dbscan/adjgraph/algo.cuh
+++ b/cpp/src/dbscan/adjgraph/algo.cuh
@@ -16,15 +16,17 @@
 
 #pragma once
 
-#include <thrust/device_ptr.h>
-#include <thrust/scan.h>
-#include <common/allocatorAdapter.hpp>
-#include <raft/cuda_utils.cuh>
 #include "../common.cuh"
 #include "pack.h"
 
+#include <common/allocatorAdapter.hpp>
+
+#include <raft/cuda_utils.cuh>
 #include <raft/sparse/convert/csr.cuh>
 
+#include <thrust/device_ptr.h>
+#include <thrust/scan.h>
+
 using namespace thrust;
 
 namespace ML {
@@ -32,8 +34,6 @@ namespace Dbscan {
 namespace AdjGraph {
 namespace Algo {
 
-using namespace MLCommon;
-
 static const int TPB_X = 256;
 
 /**
@@ -61,4 +61,4 @@ void launcher(const raft::handle_t& handle,
 }  // namespace Algo
 }  // namespace AdjGraph
 }  // namespace Dbscan
-}  // namespace ML
\ No newline at end of file
+}  // namespace ML
diff --git a/cpp/src/dbscan/dbscan.cuh b/cpp/src/dbscan/dbscan.cuh
index 5250536aae..467ecb0839 100644
--- a/cpp/src/dbscan/dbscan.cuh
+++ b/cpp/src/dbscan/dbscan.cuh
@@ -16,13 +16,16 @@
 
 #pragma once
 
+#include "runner.cuh"
+
 #include <common/nvtx.hpp>
+
 #include <cuml/cluster/dbscan.hpp>
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/common/logger.hpp>
-#include "runner.cuh"
 
 #include <algorithm>
+#include <cstddef>
 
 namespace ML {
 namespace Dbscan {
@@ -65,7 +68,7 @@ size_t compute_batch_size(size_t& estimated_memory,
 
   // To avoid overflow, we need: batch_size <= MAX_LABEL / n_rows (floor div)
   Index_ MAX_LABEL = std::numeric_limits<Index_>::max();
-  if (batch_size > MAX_LABEL / n_rows) {
+  if (batch_size > static_cast<std::size_t>(MAX_LABEL / n_rows)) {
     Index_ new_batch_size = MAX_LABEL / n_rows;
     CUML_LOG_WARN(
       "Batch size limited by the chosen integer type (%d bytes). %d -> %d. "
@@ -77,7 +80,8 @@ size_t compute_batch_size(size_t& estimated_memory,
   }
 
   // Warn when a smaller index type could be used
-  if (sizeof(Index_) > sizeof(int) && batch_size < std::numeric_limits<int>::max() / n_rows) {
+  if ((sizeof(Index_) > sizeof(int)) &&
+      (batch_size < std::numeric_limits<int>::max() / static_cast<std::size_t>(n_rows))) {
     CUML_LOG_WARN(
       "You are using an index type of size (%d bytes) but a smaller index "
       "type (%d bytes) would be sufficient. Using the smaller integer type "
@@ -110,8 +114,11 @@ void dbscanFitImpl(const raft::handle_t& handle,
   int algo_adj = 1;
   int algo_ccl = 2;
 
-  int my_rank, n_rank;
-  Index_ start_row, n_owned_rows;
+  int my_rank{0};
+  int n_rank{1};
+  Index_ start_row{0};
+  Index_ n_owned_rows{n_rows};
+
   if (opg) {
     const auto& comm     = handle.get_comms();
     my_rank              = comm.get_rank();
@@ -122,10 +129,6 @@ void dbscanFitImpl(const raft::handle_t& handle,
     n_owned_rows         = max(Index_(0), end_row - start_row);
     // Note: it is possible for a node to have no work in theory. It won't
     // happen in practice (because n_rows is much greater than n_rank)
-  } else {
-    my_rank      = 0;
-    n_rank       = 1;
-    n_owned_rows = n_rows;
   }
 
   CUML_LOG_DEBUG("#%d owns %ld rows", (int)my_rank, (unsigned long)n_owned_rows);
@@ -200,4 +203,4 @@ void dbscanFitImpl(const raft::handle_t& handle,
 }
 
 }  // namespace Dbscan
-}  // namespace ML
\ No newline at end of file
+}  // namespace ML
diff --git a/cpp/src/dbscan/mergelabels/runner.cuh b/cpp/src/dbscan/mergelabels/runner.cuh
index e43ba382d1..506239e04a 100644
--- a/cpp/src/dbscan/mergelabels/runner.cuh
+++ b/cpp/src/dbscan/mergelabels/runner.cuh
@@ -18,6 +18,7 @@
 
 #include <label/merge_labels.cuh>
 
+#include <raft/handle.hpp>
 namespace ML {
 namespace Dbscan {
 namespace MergeLabels {
diff --git a/cpp/src/dbscan/mergelabels/tree_reduction.cuh b/cpp/src/dbscan/mergelabels/tree_reduction.cuh
index 5c0f12a6e9..98ec5f8cb0 100644
--- a/cpp/src/dbscan/mergelabels/tree_reduction.cuh
+++ b/cpp/src/dbscan/mergelabels/tree_reduction.cuh
@@ -18,6 +18,10 @@
 
 #include "runner.cuh"
 
+#include <common/nvtx.hpp>
+
+#include <cuml/common/logger.hpp>
+
 namespace ML {
 namespace Dbscan {
 namespace MergeLabels {
diff --git a/cpp/src/dbscan/runner.cuh b/cpp/src/dbscan/runner.cuh
index 12bde65df9..178a29887d 100644
--- a/cpp/src/dbscan/runner.cuh
+++ b/cpp/src/dbscan/runner.cuh
@@ -16,13 +16,6 @@
 
 #pragma once
 
-#include <raft/cudart_utils.h>
-#include <common/nvtx.hpp>
-#include <cuml/common/device_buffer.hpp>
-#include <label/classlabels.cuh>
-#include <raft/cuda_utils.cuh>
-#include <raft/mr/device/allocator.hpp>
-#include <raft/sparse/csr.cuh>
 #include "adjgraph/runner.cuh"
 #include "corepoints/compute.cuh"
 #include "corepoints/exchange.cuh"
@@ -30,8 +23,20 @@
 #include "mergelabels/tree_reduction.cuh"
 #include "vertexdeg/runner.cuh"
 
+#include <cuml/common/device_buffer.hpp>
 #include <cuml/common/logger.hpp>
 
+#include <common/nvtx.hpp>
+
+#include <label/classlabels.cuh>
+
+#include <raft/cudart_utils.h>
+#include <raft/cuda_utils.cuh>
+#include <raft/mr/device/allocator.hpp>
+#include <raft/sparse/csr.cuh>
+
+#include <cstddef>
+
 namespace ML {
 namespace Dbscan {
 
@@ -101,25 +106,25 @@ void final_relabel(Index_* db_cluster,
  * @return In case the workspace pointer is null, this returns the size needed.
  */
 template <typename Type_f, typename Index_ = int, bool opg = false>
-size_t run(const raft::handle_t& handle,
-           const Type_f* x,
-           Index_ N,
-           Index_ D,
-           Index_ start_row,
-           Index_ n_owned_rows,
-           Type_f eps,
-           Index_ min_pts,
-           Index_* labels,
-           Index_* core_indices,
-           int algo_vd,
-           int algo_adj,
-           int algo_ccl,
-           void* workspace,
-           size_t batch_size,
-           cudaStream_t stream)
+std::size_t run(const raft::handle_t& handle,
+                const Type_f* x,
+                Index_ N,
+                Index_ D,
+                Index_ start_row,
+                Index_ n_owned_rows,
+                Type_f eps,
+                Index_ min_pts,
+                Index_* labels,
+                Index_* core_indices,
+                int algo_vd,
+                int algo_adj,
+                int algo_ccl,
+                void* workspace,
+                std::size_t batch_size,
+                cudaStream_t stream)
 {
-  const size_t align = 256;
-  Index_ n_batches   = raft::ceildiv((size_t)n_owned_rows, batch_size);
+  const std::size_t align = 256;
+  Index_ n_batches        = raft::ceildiv((std::size_t)n_owned_rows, batch_size);
 
   int my_rank;
   if (opg) {
@@ -139,16 +144,16 @@ size_t run(const raft::handle_t& handle,
    * elements in their neighborhood, so any IdxType can be safely used, so long as N doesn't
    * overflow.
    */
-  size_t adj_size      = raft::alignTo<size_t>(sizeof(bool) * N * batch_size, align);
-  size_t core_pts_size = raft::alignTo<size_t>(sizeof(bool) * N, align);
-  size_t m_size        = raft::alignTo<size_t>(sizeof(bool), align);
-  size_t vd_size       = raft::alignTo<size_t>(sizeof(Index_) * (batch_size + 1), align);
-  size_t ex_scan_size  = raft::alignTo<size_t>(sizeof(Index_) * batch_size, align);
-  size_t labels_size   = raft::alignTo<size_t>(sizeof(Index_) * N, align);
+  std::size_t adj_size      = raft::alignTo<std::size_t>(sizeof(bool) * N * batch_size, align);
+  std::size_t core_pts_size = raft::alignTo<std::size_t>(sizeof(bool) * N, align);
+  std::size_t m_size        = raft::alignTo<std::size_t>(sizeof(bool), align);
+  std::size_t vd_size       = raft::alignTo<std::size_t>(sizeof(Index_) * (batch_size + 1), align);
+  std::size_t ex_scan_size  = raft::alignTo<std::size_t>(sizeof(Index_) * batch_size, align);
+  std::size_t labels_size   = raft::alignTo<std::size_t>(sizeof(Index_) * N, align);
 
   Index_ MAX_LABEL = std::numeric_limits<Index_>::max();
 
-  ASSERT(N * batch_size < MAX_LABEL,
+  ASSERT(N * batch_size < static_cast<std::size_t>(MAX_LABEL),
          "An overflow occurred with the current choice of precision "
          "and the number of samples. (Max allowed batch size is %ld, but was %ld). "
          "Consider using double precision for the output labels.",
@@ -278,7 +283,7 @@ size_t run(const raft::handle_t& handle,
   if (my_rank == 0) {
     ML::PUSH_RANGE("Trace::Dbscan::FinalRelabel");
     if (algo_ccl == 2) final_relabel(labels, N, stream, handle.get_device_allocator());
-    size_t nblks = raft::ceildiv<size_t>(N, TPB);
+    std::size_t nblks = raft::ceildiv<std::size_t>(N, TPB);
     relabelForSkl<Index_><<<nblks, TPB, 0, stream>>>(labels, N, MAX_LABEL);
     CUDA_CHECK(cudaPeekAtLastError());
     ML::POP_RANGE();
@@ -302,20 +307,19 @@ size_t run(const raft::handle_t& handle,
 
       // Perform stream reduction on the core points. The core_pts acts as the stencil and we use
       // thrust::counting_iterator to return the index
-      auto core_point_count =
-        thrust::copy_if(thrust_exec_policy,
-                        index_iterator,
-                        index_iterator + N,
-                        dev_core_pts,
-                        dev_core_indices,
-                        [=] __device__(const bool is_core_point) { return is_core_point; });
+      thrust::copy_if(thrust_exec_policy,
+                      index_iterator,
+                      index_iterator + N,
+                      dev_core_pts,
+                      dev_core_indices,
+                      [=] __device__(const bool is_core_point) { return is_core_point; });
 
       ML::POP_RANGE();
     }
   }
 
   CUML_LOG_DEBUG("Done.");
-  return (size_t)0;
+  return (std::size_t)0;
 }
 }  // namespace Dbscan
 }  // namespace ML
diff --git a/cpp/src/decisiontree/decisiontree.cuh b/cpp/src/decisiontree/decisiontree.cuh
index dd9849f64f..cf54531ae1 100644
--- a/cpp/src/decisiontree/decisiontree.cuh
+++ b/cpp/src/decisiontree/decisiontree.cuh
@@ -15,32 +15,37 @@
  */
 
 #pragma once
-#include <common/Timer.h>
+
+#include "batched-levelalgo/builder.cuh"
+#include "quantile/quantile.h"
+#include "treelite_util.h"
+
 #include <cuml/tree/algo_helper.h>
 #include <cuml/tree/flatnode.h>
+#include <cuml/common/logger.hpp>
+#include <cuml/tree/decisiontree.hpp>
+
+#include <common/Timer.h>
+#include <common/iota.cuh>
+#include <common/nvtx.hpp>
+
 #include <raft/cudart_utils.h>
+#include <raft/handle.hpp>
+#include <raft/mr/device/allocator.hpp>
+#include <raft/mr/host/allocator.hpp>
+
 #include <treelite/c_api.h>
 #include <treelite/tree.h>
+
 #include <algorithm>
 #include <climits>
-#include <common/iota.cuh>
-#include <cuml/common/logger.hpp>
-#include <cuml/tree/decisiontree.hpp>
 #include <iomanip>
 #include <locale>
 #include <map>
 #include <numeric>
-#include <raft/handle.hpp>
-#include <raft/mr/device/allocator.hpp>
-#include <raft/mr/host/allocator.hpp>
 #include <random>
 #include <type_traits>
 #include <vector>
-#include "batched-levelalgo/builder.cuh"
-#include "quantile/quantile.h"
-#include "treelite_util.h"
-
-#include <common/nvtx.hpp>
 
 /** check for treelite runtime API errors and assert accordingly */
 #define TREELITE_CHECK(call)                                                                     \
diff --git a/cpp/src/fil/fil.cu b/cpp/src/fil/fil.cu
index 303a771a3b..d26f6b3c16 100644
--- a/cpp/src/fil/fil.cu
+++ b/cpp/src/fil/fil.cu
@@ -16,28 +16,34 @@
 
 /** @file fil.cu implements forest inference */
 
-#include <omp.h>
+#include "common.cuh"
+
+#include <cuml/fil/fil.h>
+#include <cuml/fil/fnv_hash.h>
+#include <cuml/common/logger.hpp>
+
+#include <raft/cudart_utils.h>
+#include <raft/handle.hpp>
+#include <raft/mr/device/allocator.hpp>
+#include <raft/mr/host/allocator.hpp>
+
+#include <treelite/c_api.h>
+#include <treelite/tree.h>
+
 #include <thrust/device_ptr.h>
 #include <thrust/device_vector.h>
 #include <thrust/host_vector.h>
-#include <treelite/c_api.h>
-#include <treelite/tree.h>
+
+#include <omp.h>
+
 #include <algorithm>
 #include <cmath>
+#include <cstddef>
 #include <iomanip>
 #include <limits>
 #include <stack>
 #include <utility>
 
-#include <cuml/fil/fil.h>
-#include <cuml/fil/fnv_hash.h>
-#include <raft/cudart_utils.h>
-#include <cuml/common/logger.hpp>
-#include <raft/handle.hpp>
-#include <raft/mr/device/allocator.hpp>
-#include <raft/mr/host/allocator.hpp>
-#include "common.cuh"
-
 namespace ML {
 namespace fil {
 
@@ -266,7 +272,7 @@ struct forest {
           // for GROVE_PER_CLASS, averaging happens in infer_k
           ot                 = output_t(ot & ~output_t::AVG);
           params.num_outputs = params.num_classes;
-          do_transform = ot != output_t::RAW && ot != output_t::SOFTMAX || global_bias != 0.0f;
+          do_transform = (ot != output_t::RAW && ot != output_t::SOFTMAX) || global_bias != 0.0f;
           break;
         case leaf_algo_t::CATEGORICAL_LEAF:
           params.num_outputs = params.num_classes;
@@ -276,7 +282,7 @@ struct forest {
           // for VECTOR_LEAF, averaging happens in infer_k
           ot                 = output_t(ot & ~output_t::AVG);
           params.num_outputs = params.num_classes;
-          do_transform = ot != output_t::RAW && ot != output_t::SOFTMAX || global_bias != 0.0f;
+          do_transform = (ot != output_t::RAW && ot != output_t::SOFTMAX) || global_bias != 0.0f;
           break;
         default: ASSERT(false, "internal error: invalid leaf_algo_");
       }
@@ -633,12 +639,12 @@ void tl2fil_leaf_payload(fil_node_t* fil_node,
   auto vec = tl_tree.LeafVector(tl_node_id);
   switch (forest_params.leaf_algo) {
     case leaf_algo_t::CATEGORICAL_LEAF:
-      ASSERT(vec.size() == forest_params.num_classes,
+      ASSERT(vec.size() == static_cast<std::size_t>(forest_params.num_classes),
              "inconsistent number of classes in treelite leaves");
       fil_node->val.idx = find_class_label_from_one_hot(&vec[0], vec.size());
       break;
     case leaf_algo_t::VECTOR_LEAF: {
-      ASSERT(vec.size() == forest_params.num_classes,
+      ASSERT(vec.size() == static_cast<std::size_t>(forest_params.num_classes),
              "inconsistent number of classes in treelite leaves");
       fil_node->val.idx = *leaf_counter;
       for (int k = 0; k < forest_params.num_classes; k++) {
@@ -769,14 +775,14 @@ inline void tree_depth_hist(const tl::Tree<T, L>& tree, std::vector<level_entry>
     stack.pop();
 
     while (!tree.IsLeaf(node_id)) {
-      if (depth >= hist.size()) hist.resize(depth + 1, {0, 0});
+      if (static_cast<std::size_t>(depth) >= hist.size()) hist.resize(depth + 1, {0, 0});
       hist[depth].n_branch_nodes++;
       stack.push({tree.LeftChild(node_id), depth + 1});
       node_id = tree.RightChild(node_id);
       depth++;
     }
 
-    if (depth >= hist.size()) hist.resize(depth + 1, {0, 0});
+    if (static_cast<std::size_t>(depth) >= hist.size()) hist.resize(depth + 1, {0, 0});
     hist[depth].n_leaves++;
   }
 }
@@ -794,7 +800,7 @@ std::stringstream depth_hist_and_max(const tl::ModelImpl<T, L>& model)
   ios default_state(nullptr);
   default_state.copyfmt(forest_shape);
   forest_shape << "Depth histogram:" << endl << "depth branches leaves   nodes" << endl;
-  for (int level = 0; level < hist.size(); ++level) {
+  for (std::size_t level = 0; level < hist.size(); ++level) {
     level_entry e = hist[level];
     forest_shape << setw(5) << level << setw(9) << e.n_branch_nodes << setw(7) << e.n_leaves
                  << setw(8) << e.n_branch_nodes + e.n_leaves << endl;
@@ -928,7 +934,7 @@ void tl2fil_dense(std::vector<dense_node>* pnodes,
     vector_leaf->resize(max_leaves_per_tree * params->num_trees * params->num_classes);
   }
   pnodes->resize(num_nodes, dense_node());
-  for (int i = 0; i < model.trees.size(); ++i) {
+  for (std::size_t i = 0; i < model.trees.size(); ++i) {
     size_t leaf_counter = max_leaves_per_tree * i;
     tree2fil_dense(pnodes,
                    i * tree_num_nodes(params->depth),
@@ -976,10 +982,10 @@ struct tl2fil_sparse_check_t<sparse_node8> {
 
     // check the number of tree nodes
     const std::vector<tl::Tree<threshold_t, leaf_t>>& trees = model.trees;
-    for (int i = 0; i < trees.size(); ++i) {
+    for (std::size_t i = 0; i < trees.size(); ++i) {
       int num_nodes = trees[i].num_nodes;
       ASSERT(num_nodes <= MAX_TREE_NODES,
-             "tree %d has %d nodes, "
+             "tree %lu has %d nodes, "
              "but only %d supported for 8-byte sparse nodes",
              i,
              num_nodes,
@@ -1019,7 +1025,7 @@ void tl2fil_sparse(std::vector<int>* ptrees,
 
   // convert the nodes
 #pragma omp parallel for
-  for (int i = 0; i < num_trees; ++i) {
+  for (std::size_t i = 0; i < num_trees; ++i) {
     // Max number of leaves processed so far
     size_t leaf_counter = ((*ptrees)[i] + i) / 2;
     tree2fil_sparse(*pnodes, (*ptrees)[i], model.trees[i], *params, vector_leaf, &leaf_counter);
diff --git a/cpp/src/fil/infer.cu b/cpp/src/fil/infer.cu
index e8b097e419..1cdc9a09fe 100644
--- a/cpp/src/fil/infer.cu
+++ b/cpp/src/fil/infer.cu
@@ -14,14 +14,26 @@
  * limitations under the License.
  */
 
-#include <algorithm>
-#include <cmath>
+#include "common.cuh"
+
+#include <fil/internal.cuh>
+
+#include <cuml/fil/multi_sum.cuh>
 
 #include <raft/cudart_utils.h>
+
 #include <thrust/functional.h>
-#include <cuml/fil/multi_sum.cuh>
-#include <fil/internal.cuh>
-#include "common.cuh"
+
+#include <algorithm>
+#include <cmath>
+
+#ifndef CUDA_PRAGMA_UNROLL
+#ifdef __CUDA_ARCH__
+#define CUDA_PRAGMA_UNROLL _Pragma("unroll")
+#else
+#define CUDA_PRAGMA_UNROLL
+#endif  // __CUDA_ARCH__
+#endif  // CUDA_PRAGMA_UNROLL
 
 namespace ML {
 namespace fil {
@@ -39,7 +51,7 @@ struct Vectorized {
                                                                           vec<NITEMS, T> b) const
   {
     vec<NITEMS, T> c;
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int i = 0; i < NITEMS; i++)
       c[i] = op(a[i], b[i]);
     return c;
@@ -57,7 +69,7 @@ struct vec {
   T data[N];
   explicit __host__ __device__ vec(T t)
   {
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int i = 0; i < N; ++i)
       data[i] = t;
   }
@@ -96,7 +108,7 @@ template <int NITEMS>
 __device__ __forceinline__ vec<NITEMS, best_margin_label> to_vec(int c, vec<NITEMS, float> margin)
 {
   vec<NITEMS, best_margin_label> ret;
-#pragma unroll
+  CUDA_PRAGMA_UNROLL
   for (int i = 0; i < NITEMS; ++i)
     ret[i] = best_margin_label(c, margin[i]);
   return ret;
@@ -108,7 +120,7 @@ struct ArgMax {
     vec<NITEMS, best_margin_label> a, vec<NITEMS, best_margin_label> b) const
   {
     vec<NITEMS, best_margin_label> c;
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int i = 0; i < NITEMS; i++)
       c[i] = cub::ArgMax()(a[i], b[i]);
     return c;
@@ -125,7 +137,7 @@ __device__ __forceinline__ vec<NITEMS, output_type> tree_leaf_output(tree_type t
                                                                      int (&leaves)[NITEMS])
 {
   vec<NITEMS, output_type> out(0);
-#pragma unroll
+  CUDA_PRAGMA_UNROLL
   for (int j = 0; j < NITEMS; ++j) {
     if (FULL_NITEMS || j < n_rows) {
       /** dependent names are not considered templates by default, unless it's a
@@ -151,7 +163,7 @@ __device__ __forceinline__ vec<NITEMS, output_type> infer_one_tree(tree_type tre
   for (int j = 0; j < NITEMS; ++j)
     curr[j] = 0;
   do {
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int j = 0; j < NITEMS; ++j) {
       auto n = tree[curr[j]];
       mask &= ~(n.is_leaf() << j);
@@ -306,7 +318,7 @@ struct tree_aggregator_t {
     }
 
     if (threadIdx.x * NITEMS >= block_num_rows) return;
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int row = 0; row < NITEMS; ++row) {
       int out_preds_i = threadIdx.x * NITEMS + row;
       if (out_preds_i < block_num_rows) block_out[out_preds_i * output_stride] = acc[row];
@@ -352,7 +364,7 @@ __device__ __forceinline__ void write_best_class(
   best = block_reduce(best, vectorized(cub::ArgMax()), tmp_storage);
   // write it out to global memory
   if (threadIdx.x > 0) return;
-#pragma unroll
+  CUDA_PRAGMA_UNROLL
   for (int row = 0; row < best.NITEMS; ++row)
     if (row < num_rows) out[row] = best[row].key;
 }
@@ -393,8 +405,8 @@ __device__ __forceinline__ void normalize_softmax_and_write(Iterator begin,
       *it /= trees_per_class;
   }
   if ((transform & output_t::SOFTMAX) != 0) block_softmax(begin, end, tmp_storage);
-// write result to global memory
-#pragma unroll
+  // write result to global memory
+  CUDA_PRAGMA_UNROLL
   for (int row = 0; row < begin->NITEMS; ++row) {
     for (int c = threadIdx.x; c < end - begin; c += blockDim.x)
       if (row < num_rows) out[row * (end - begin) + c] = begin[c][row];
@@ -683,9 +695,9 @@ struct tree_aggregator_t<NITEMS, CATEGORICAL_LEAF> {
     : num_classes(params.num_classes), votes((int*)accumulate_workspace)
   {
     for (int c = threadIdx.x; c < num_classes; c += FIL_TPB * NITEMS)
-#pragma unroll
-      for (int item = 0; item < NITEMS; ++item)
-        votes[c * NITEMS + item] = 0;
+      CUDA_PRAGMA_UNROLL
+    for (int item = 0; item < NITEMS; ++item)
+      votes[c * NITEMS + item] = 0;
     // __syncthreads() is called in infer_k
   }
   __device__ __forceinline__ void accumulate(vec<NITEMS, int> single_tree_prediction,
@@ -693,7 +705,7 @@ struct tree_aggregator_t<NITEMS, CATEGORICAL_LEAF> {
                                              int thread_num_rows)
   {
     if (thread_num_rows == 0) return;
-#pragma unroll
+    CUDA_PRAGMA_UNROLL
     for (int item = 0; item < NITEMS; ++item) {
       raft::myAtomicAdd(votes + single_tree_prediction[item] * NITEMS + item, 1);
     }
@@ -704,7 +716,7 @@ struct tree_aggregator_t<NITEMS, CATEGORICAL_LEAF> {
   {
     __syncthreads();
     for (int c = threadIdx.x; c < num_classes; c += blockDim.x) {
-#pragma unroll
+      CUDA_PRAGMA_UNROLL
       for (int row = 0; row < num_rows; ++row)
         out[row * num_classes + c] = votes[c * NITEMS + row];
     }
@@ -762,7 +774,7 @@ __global__ void infer_k(storage_type forest, predict_params params)
       // cache the row for all threads to reuse
       // 2021: latest SMs still do not have >256KiB of shared memory/block required to
       // exceed the uint16_t
-#pragma unroll
+      CUDA_PRAGMA_UNROLL
       for (uint16_t input_idx = threadIdx.x; input_idx < block_num_rows * num_cols;
            input_idx += blockDim.x) {
         // for even num_cols, we need to pad sdata_stride to reduce bank conflicts
@@ -772,7 +784,7 @@ __global__ void infer_k(storage_type forest, predict_params params)
           sdata_stride == num_cols ? input_idx : input_idx + input_idx / (uint16_t)num_cols;
         sdata[sdata_idx] = block_input[input_idx];
       }
-#pragma unroll
+      CUDA_PRAGMA_UNROLL
       for (int idx = block_num_rows * sdata_stride; idx < rows_per_block * sdata_stride;
            idx += blockDim.x)
         sdata[idx] = 0.0f;
diff --git a/cpp/src/glm/ols_mg.cu b/cpp/src/glm/ols_mg.cu
index b9905dfb7f..3200eddc00 100644
--- a/cpp/src/glm/ols_mg.cu
+++ b/cpp/src/glm/ols_mg.cu
@@ -16,8 +16,10 @@
 
 #include <cuml/linear_model/ols_mg.hpp>
 #include <cuml/linear_model/preprocess_mg.hpp>
+
 #include <opg/linalg/lstsq.hpp>
 #include <opg/stats/mean.hpp>
+
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/linalg/add.cuh>
@@ -26,8 +28,11 @@
 #include <raft/matrix/matrix.cuh>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/mr/host/allocator.hpp>
+
 #include <rmm/device_uvector.hpp>
 
+#include <cstddef>
+
 using namespace MLCommon;
 
 namespace ML {
@@ -48,10 +53,6 @@ void fit_impl(raft::handle_t& handle,
               int n_streams,
               bool verbose)
 {
-  const auto& comm                   = handle.get_comms();
-  cublasHandle_t cublas_handle       = handle.get_cublas_handle();
-  cusolverDnHandle_t cusolver_handle = handle.get_cusolver_dn_handle();
-
   rmm::device_uvector<T> mu_input(0, streams[0]);
   rmm::device_uvector<T> norm2_input(0, streams[0]);
   rmm::device_uvector<T> mu_labels(0, streams[0]);
@@ -175,7 +176,7 @@ void predict_impl(raft::handle_t& handle,
   T alpha                                         = T(1);
   T beta                                          = T(0);
 
-  for (int i = 0; i < input_data.size(); i++) {
+  for (std::size_t i = 0; i < input_data.size(); i++) {
     int si = i % n_streams;
     raft::linalg::gemm(handle,
                        input_data[i]->ptr,
diff --git a/cpp/src/glm/ridge_mg.cu b/cpp/src/glm/ridge_mg.cu
index 3f90ac863e..34df73c25f 100644
--- a/cpp/src/glm/ridge_mg.cu
+++ b/cpp/src/glm/ridge_mg.cu
@@ -14,20 +14,25 @@
  * limitations under the License.
  */
 
-#include <raft/cudart_utils.h>
 #include <cuml/linear_model/preprocess_mg.hpp>
 #include <cuml/linear_model/ridge_mg.hpp>
+
 #include <opg/linalg/mv_aTb.hpp>
 #include <opg/linalg/svd.hpp>
 #include <opg/stats/mean.hpp>
+
+#include <raft/cudart_utils.h>
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/linalg/add.cuh>
 #include <raft/linalg/gemm.cuh>
 #include <raft/matrix/math.cuh>
 #include <raft/matrix/matrix.cuh>
+
 #include <rmm/device_uvector.hpp>
 
+#include <cstddef>
+
 using namespace MLCommon;
 
 namespace ML {
@@ -48,10 +53,6 @@ void ridgeSolve(const raft::handle_t& handle,
                 int n_streams,
                 bool verbose)
 {
-  auto cublasH     = handle.get_cublas_handle();
-  auto cusolverH   = handle.get_cusolver_dn_handle();
-  const auto& comm = handle.get_comms();
-
   // Implements this: w = V * inv(S^2 + λ*I) * S * U^T * b
   T* S_nnz;
   T alp   = T(1);
@@ -102,9 +103,7 @@ void ridgeEig(raft::handle_t& handle,
               int n_streams,
               bool verbose)
 {
-  const auto& comm                         = handle.get_comms();
-  const cublasHandle_t cublas_handle       = handle.get_cublas_handle();
-  const cusolverDnHandle_t cusolver_handle = handle.get_cusolver_dn_handle();
+  const auto& comm = handle.get_comms();
 
   int rank = comm.get_rank();
 
@@ -116,7 +115,7 @@ void ridgeEig(raft::handle_t& handle,
   std::vector<Matrix::RankSizePair*> partsToRanks = ADesc.blocksOwnedBy(rank);
   size_t total_size                               = 0;
 
-  for (int i = 0; i < partsToRanks.size(); i++) {
+  for (std::size_t i = 0; i < partsToRanks.size(); i++) {
     total_size += partsToRanks[i]->size;
   }
   total_size = total_size * ADesc.N;
@@ -124,7 +123,7 @@ void ridgeEig(raft::handle_t& handle,
   rmm::device_uvector<T> U_parts(total_size, streams[0]);
   T* curr_ptr = U_parts.data();
 
-  for (int i = 0; i < partsToRanks.size(); i++) {
+  for (std::size_t i = 0; i < partsToRanks.size(); i++) {
     Matrix::Data<T> d;
     d.totalSize = partsToRanks[i]->size;
     d.ptr       = curr_ptr;
@@ -132,7 +131,7 @@ void ridgeEig(raft::handle_t& handle,
     U_temp.push_back(d);
   }
 
-  for (int i = 0; i < A.size(); i++) {
+  for (std::size_t i = 0; i < A.size(); i++) {
     U.push_back(&(U_temp[i]));
   }
 
@@ -291,7 +290,7 @@ void predict_impl(raft::handle_t& handle,
   T alpha                                         = T(1);
   T beta                                          = T(0);
 
-  for (int i = 0; i < input_data.size(); i++) {
+  for (std::size_t i = 0; i < input_data.size(); i++) {
     int si = i % n_streams;
     raft::linalg::gemm(handle,
                        input_data[i]->ptr,
diff --git a/cpp/src/hdbscan/detail/extract.cuh b/cpp/src/hdbscan/detail/extract.cuh
index 13ff15c002..4514745604 100644
--- a/cpp/src/hdbscan/detail/extract.cuh
+++ b/cpp/src/hdbscan/detail/extract.cuh
@@ -16,25 +16,22 @@
 
 #pragma once
 
-#include <label/classlabels.cuh>
-
-#include <cub/cub.cuh>
-
-#include <raft/cudart_utils.h>
-
 #include "membership.cuh"
 #include "select.cuh"
 #include "stabilities.cuh"
 #include "utils.h"
 
-#include <raft/sparse/op/sort.h>
-#include <raft/sparse/convert/csr.cuh>
+#include <label/classlabels.cuh>
 
 #include <cuml/cluster/hdbscan.hpp>
 
+#include <raft/cudart_utils.h>
+#include <raft/sparse/op/sort.h>
 #include <raft/label/classlabels.cuh>
+#include <raft/sparse/convert/csr.cuh>
 
-#include <algorithm>
+#include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
@@ -42,8 +39,10 @@
 #include <thrust/sort.h>
 #include <thrust/transform.h>
 
-#include <rmm/device_uvector.hpp>
-#include <rmm/exec_policy.hpp>
+#include <cub/cub.cuh>
+
+#include <algorithm>
+#include <cstddef>
 
 namespace ML {
 namespace HDBSCAN {
@@ -234,7 +233,7 @@ value_idx extract_clusters(const raft::handle_t& handle,
   CUDA_CHECK(cudaStreamSynchronize(stream));
 
   std::set<value_idx> clusters;
-  for (int i = 0; i < is_cluster_h.size(); i++) {
+  for (std::size_t i = 0; i < is_cluster_h.size(); i++) {
     if (is_cluster_h[i] != 0) { clusters.insert(i + n_leaves); }
   }
 
diff --git a/cpp/src/hdbscan/detail/reachability_faiss.cuh b/cpp/src/hdbscan/detail/reachability_faiss.cuh
index 23eeaf8d67..a70d797b97 100644
--- a/cpp/src/hdbscan/detail/reachability_faiss.cuh
+++ b/cpp/src/hdbscan/detail/reachability_faiss.cuh
@@ -39,8 +39,13 @@
 #include <faiss/gpu/utils/Limits.cuh>
 #include <faiss/gpu/utils/MatrixMult.cuh>
 
+#include <raft/cuda_utils.cuh>
+#include <raft/handle.hpp>
+
 #include <rmm/device_uvector.hpp>
 
+#include <cstddef>
+
 namespace ML {
 namespace HDBSCAN {
 namespace detail {
@@ -306,31 +311,31 @@ void mutual_reachability_knn_l2(const raft::handle_t& handle,
   bool interrupt = false;
 
   // Tile over the input queries
-  for (int i = 0; i < m; i += tileRows) {
+  for (std::size_t i = 0; i < m; i += tileRows) {
     if (interrupt || faiss::InterruptCallback::is_interrupted()) {
       interrupt = true;
       break;
     }
 
-    int curQuerySize = std::min((size_t)tileRows, m - i);
+    int curQuerySize = std::min(static_cast<std::size_t>(tileRows), m - i);
 
     auto outDistanceView = out_dists_tensor.narrow(0, i, curQuerySize);
     auto outIndexView    = out_inds_tensor.narrow(0, i, curQuerySize);
 
-    auto queryView     = x_tensor.narrow(0, i, curQuerySize);
-    auto queryNormNiew = norms_tensor.narrow(0, i, curQuerySize);
+    auto queryView = x_tensor.narrow(0, i, curQuerySize);
+    norms_tensor.narrow(0, i, curQuerySize);
 
     auto outDistanceBufRowView = outDistanceBufs[curStream]->narrow(0, 0, curQuerySize);
     auto outIndexBufRowView    = outIndexBufs[curStream]->narrow(0, 0, curQuerySize);
 
     // Tile over the centroids
-    for (int j = 0; j < m; j += tileCols) {
+    for (std::size_t j = 0; j < m; j += tileCols) {
       if (faiss::InterruptCallback::is_interrupted()) {
         interrupt = true;
         break;
       }
 
-      int curCentroidSize = std::min((size_t)tileCols, m - j);
+      int curCentroidSize = std::min(static_cast<std::size_t>(tileCols), m - j);
       int curColTile      = j / tileCols;
 
       auto centroidsView = sliceCentroids(x_tensor, true, j, curCentroidSize);
@@ -352,7 +357,7 @@ void mutual_reachability_knn_l2(const raft::handle_t& handle,
                     gpu_res->getBlasHandleCurrentDevice(),
                     streams[curStream]);
 
-      if (tileCols == m) {
+      if (static_cast<std::size_t>(tileCols) == m) {
         // Write into the final output
         runL2SelectMin<value_t>(distanceBufView,
                                 norms_tensor,
@@ -364,7 +369,7 @@ void mutual_reachability_knn_l2(const raft::handle_t& handle,
                                 alpha,
                                 streams[curStream]);
       } else {
-        auto centroidNormsView = norms_tensor.narrow(0, j, curCentroidSize);
+        norms_tensor.narrow(0, j, curCentroidSize);
 
         // Write into our intermediate output
         runL2SelectMin<value_t>(distanceBufView,
@@ -381,7 +386,7 @@ void mutual_reachability_knn_l2(const raft::handle_t& handle,
 
     // As we're finished with processing a full set of centroids, perform
     // the final k-selection
-    if (tileCols != m) {
+    if (static_cast<std::size_t>(tileCols) != m) {
       // The indices are tile-relative; for each tile of k, we need to add
       // tileCols to the index
       faiss::gpu::runIncrementIndex(outIndexBufRowView, k, tileCols, streams[curStream]);
@@ -407,4 +412,4 @@ void mutual_reachability_knn_l2(const raft::handle_t& handle,
 };  // end namespace Reachability
 };  // end namespace detail
 };  // end namespace HDBSCAN
-};  // end namespace ML
\ No newline at end of file
+};  // end namespace ML
diff --git a/cpp/src/holtwinters/internal/hw_eval.cuh b/cpp/src/holtwinters/internal/hw_eval.cuh
index c0867b9ace..82ebe593fb 100644
--- a/cpp/src/holtwinters/internal/hw_eval.cuh
+++ b/cpp/src/holtwinters/internal/hw_eval.cuh
@@ -15,9 +15,13 @@
  */
 
 #pragma once
-#include <raft/cudart_utils.h>
+
 #include "hw_utils.cuh"
 
+#include <raft/cudart_utils.h>
+#include <raft/handle.hpp>
+#include <raft/mr/device/buffer.hpp>
+
 template <typename Dtype>
 __device__ Dtype holtwinters_eval_device(int tid,
                                          const Dtype* ts,
@@ -252,11 +256,11 @@ void holtwinters_eval_gpu(const raft::handle_t& handle,
   int threads_per_block = GET_THREADS_PER_BLOCK(batch_size);
 
   // How much sm needed for shared kernel
-  size_t sm_needed = sizeof(Dtype) * threads_per_block * frequency;
+  int sm_needed    = sizeof(Dtype) * threads_per_block * frequency;
   bool is_additive = seasonal == ML::SeasonalType::ADDITIVE;
 
   if (sm_needed > raft::getSharedMemPerBlock()) {
-    MLCommon::device_buffer<Dtype> pseason(dev_allocator, stream, batch_size * frequency);
+    raft::mr::device::buffer<Dtype> pseason(dev_allocator, stream, batch_size * frequency);
     holtwinters_eval_gpu_global_kernel<Dtype>
       <<<total_blocks, threads_per_block, 0, stream>>>(ts,
                                                        n,
diff --git a/cpp/src/holtwinters/internal/hw_optim.cuh b/cpp/src/holtwinters/internal/hw_optim.cuh
index 9458e03f88..4483df2543 100644
--- a/cpp/src/holtwinters/internal/hw_optim.cuh
+++ b/cpp/src/holtwinters/internal/hw_optim.cuh
@@ -15,10 +15,13 @@
  */
 
 #pragma once
-#include <raft/cudart_utils.h>
+
 #include "hw_eval.cuh"
 #include "hw_utils.cuh"
 
+#include <raft/cudart_utils.h>
+#include <raft/mr/device/buffer.hpp>
+
 template <typename Dtype>
 __device__ Dtype golden_step(Dtype a, Dtype b, Dtype c)
 {
@@ -871,12 +874,12 @@ void holtwinters_optim_gpu(const raft::handle_t& handle,
   int threads_per_block = 128;
 
   // How much sm needed for shared kernel
-  size_t sm_needed  = sizeof(Dtype) * threads_per_block * frequency;
+  int sm_needed     = sizeof(Dtype) * threads_per_block * frequency;
   bool is_additive  = seasonal == ML::SeasonalType::ADDITIVE;
   bool single_param = (optim_alpha + optim_beta + optim_gamma > 1) ? false : true;
 
   if (sm_needed > raft::getSharedMemPerBlock()) {  // Global memory //
-    MLCommon::device_buffer<Dtype> pseason(dev_allocator, stream, batch_size * frequency);
+    raft::mr::device::buffer<Dtype> pseason(dev_allocator, stream, batch_size * frequency);
     holtwinters_optim_gpu_global_kernel<Dtype>
       <<<total_blocks, threads_per_block, 0, stream>>>(ts,
                                                        n,
diff --git a/cpp/src/kmeans/common.cuh b/cpp/src/kmeans/common.cuh
index e1bcf23cc9..9a275de18c 100644
--- a/cpp/src/kmeans/common.cuh
+++ b/cpp/src/kmeans/common.cuh
@@ -15,41 +15,40 @@
  */
 #pragma once
 
+#include <cuml/cluster/kmeans_mg.hpp>
+#include <cuml/common/device_buffer.hpp>
+#include <cuml/common/host_buffer.hpp>
+#include <cuml/common/logger.hpp>
 #include <cuml/metrics/metrics.hpp>
+
+#include <ml_cuda_utils.h>
+
+#include <common/allocatorAdapter.hpp>
+#include <common/tensor.hpp>
+
 #include <linalg/reduce_cols_by_key.cuh>
 #include <linalg/reduce_rows_by_key.cuh>
 #include <matrix/gather.cuh>
+#include <random/permute.cuh>
 
+#include <raft/cudart_utils.h>
+#include <raft/comms/comms.hpp>
 #include <raft/distance/fused_l2_nn.cuh>
 #include <raft/linalg/binary_op.cuh>
 #include <raft/linalg/matrix_vector_op.cuh>
 #include <raft/linalg/mean_squared_error.cuh>
 #include <raft/linalg/reduce.cuh>
 #include <raft/random/rng.cuh>
-#include <random/permute.cuh>
-#include <random>
 
 #include <thrust/equal.h>
 #include <thrust/execution_policy.h>
 #include <thrust/fill.h>
 #include <thrust/for_each.h>
 #include <thrust/scan.h>
-#include <numeric>
 
-#include <ml_cuda_utils.h>
-
-#include <common/allocatorAdapter.hpp>
-#include <common/tensor.hpp>
-#include <cuml/common/device_buffer.hpp>
-#include <cuml/common/host_buffer.hpp>
-#include <raft/comms/comms.hpp>
-
-#include <cuml/common/logger.hpp>
-
-#include <cuml/cluster/kmeans_mg.hpp>
-
-#include <raft/cudart_utils.h>
 #include <fstream>
+#include <numeric>
+#include <random>
 
 namespace ML {
 
@@ -732,7 +731,6 @@ void kmeansPlusPlus(const raft::handle_t& handle,
   std::uniform_int_distribution<> dis(0, n_samples - 1);
 
   ML::thrustAllocatorAdapter alloc(handle.get_device_allocator(), stream);
-  auto thrust_exec_policy = thrust::cuda::par(alloc).on(stream);
 
   // <<< Step-1 >>>: C <-- sample a point uniformly at random from X
   auto initialCentroid  = X.template view<2>({1, n_features}, {dis(gen), 0});
diff --git a/cpp/src/knn/knn.cu b/cpp/src/knn/knn.cu
index 56e1057a47..4c486f984d 100644
--- a/cpp/src/knn/knn.cu
+++ b/cpp/src/knn/knn.cu
@@ -24,9 +24,9 @@
 #include <raft/spatial/knn/knn.hpp>
 #include <selection/knn.cuh>
 
-#include <cuda_runtime.h>
 #include <raft/cuda_utils.cuh>
 
+#include <cstddef>
 #include <sstream>
 #include <vector>
 
@@ -102,7 +102,7 @@ void knn_classify(raft::handle_t& handle,
   std::vector<int*> uniq_labels(y.size());
   std::vector<int> n_unique(y.size());
 
-  for (int i = 0; i < y.size(); i++) {
+  for (std::size_t i = 0; i < y.size(); i++) {
     MLCommon::Label::getUniqueLabels(
       y[i], n_index_rows, &(uniq_labels[i]), &(n_unique[i]), stream, d_alloc);
   }
@@ -137,7 +137,7 @@ void knn_class_proba(raft::handle_t& handle,
   std::vector<int*> uniq_labels(y.size());
   std::vector<int> n_unique(y.size());
 
-  for (int i = 0; i < y.size(); i++) {
+  for (std::size_t i = 0; i < y.size(); i++) {
     MLCommon::Label::getUniqueLabels(
       y[i], n_index_rows, &(uniq_labels[i]), &(n_unique[i]), stream, d_alloc);
   }
diff --git a/cpp/src/knn/knn_opg_common.cuh b/cpp/src/knn/knn_opg_common.cuh
index 2c9eddd5fa..1f41d3db30 100644
--- a/cpp/src/knn/knn_opg_common.cuh
+++ b/cpp/src/knn/knn_opg_common.cuh
@@ -15,22 +15,25 @@
  */
 
 #pragma once
-#include <cuml/neighbors/knn_mg.hpp>
-#include <selection/knn.cuh>
 
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/common/logger.hpp>
-#include <raft/comms/comms.hpp>
-#include <raft/mr/device/allocator.hpp>
+#include <cuml/neighbors/knn_mg.hpp>
 
-#include <memory>
-#include <set>
+#include <selection/knn.cuh>
+
+#include <cumlprims/opg/matrix/data.hpp>
+#include <cumlprims/opg/matrix/part_descriptor.hpp>
 
 #include <raft/cudart_utils.h>
+#include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
+#include <raft/mr/device/allocator.hpp>
 #include <raft/spatial/knn/knn.hpp>
 
-#include "knn_opg_common.cuh"
+#include <cstddef>
+#include <memory>
+#include <set>
 
 namespace ML {
 namespace KNN {
@@ -89,7 +92,7 @@ struct opg_knn_param {
   size_t batch_size = 0;                                  /**< Batch size */
   bool verbose;                                           /**< verbose */
 
-  int n_outputs = 0;                      /**< Number of outputs per query (cl&re) */
+  std::size_t n_outputs = 0;              /**< Number of outputs per query (cl&re) */
   std::vector<std::vector<out_t*>>* y;    /**< Labels input array (cl&re) */
   std::vector<Matrix::Data<out_t>*>* out; /**< KNN outputs output array (cl&re) */
 
@@ -141,7 +144,7 @@ struct KNN_RE_params : public opg_knn_param<in_t, ind_t, dist_t, out_t> {
                 size_t k,
                 size_t batch_size,
                 bool verbose,
-                int n_outputs,
+                std::size_t n_outputs,
                 std::vector<std::vector<out_t*>>* y,
                 std::vector<Matrix::Data<out_t>*>* out)
     : opg_knn_param<in_t, ind_t, dist_t, out_t>(knn_op,
@@ -173,7 +176,7 @@ struct KNN_CL_params : public opg_knn_param<in_t, ind_t, dist_t, out_t> {
                 size_t k,
                 size_t batch_size,
                 bool verbose,
-                int n_outputs,
+                std::size_t n_outputs,
                 std::vector<std::vector<out_t*>>* y,
                 std::vector<int>* n_unique,
                 std::vector<out_t*>* uniq_labels,
@@ -255,9 +258,8 @@ void opg_knn(opg_knn_param<in_t, ind_t, dist_t, out_t>& params, raft::handle_t&
     size_t total_batches     = raft::ceildiv(part_n_rows, params.batch_size);
     size_t total_n_processed = 0;
 
-    // Loop through batches for each query part
-    for (int cur_batch = 0; cur_batch < total_batches;
-         cur_batch++) {  // For each batch in a query partition
+    // For each batch in a query partition
+    for (std::size_t cur_batch = 0; cur_batch < total_batches; cur_batch++) {
       size_t cur_batch_size = params.batch_size;
 
       if (cur_batch == total_batches - 1)
@@ -275,7 +277,7 @@ void opg_knn(opg_knn_param<in_t, ind_t, dist_t, out_t>& params, raft::handle_t&
       size_t batch_input_elms   = cur_batch_size * params.query_desc->N;
       size_t batch_input_offset = batch_input_elms * cur_batch;
 
-      in_t* cur_query_ptr;
+      in_t* cur_query_ptr{nullptr};
 
       device_buffer<in_t> tmp_batch_buf(handle.get_device_allocator(), handle.get_stream(), 0);
       // current partition's owner rank broadcasts
@@ -286,7 +288,7 @@ void opg_knn(opg_knn_param<in_t, ind_t, dist_t, out_t>& params, raft::handle_t&
         // temporary buffer for the batch so that we can stack rows.
         if (!params.rowMajorQuery && total_batches > 1) {
           tmp_batch_buf.resize(batch_input_elms, handle.get_stream());
-          for (int col_data = 0; col_data < params.query_desc->N; col_data++) {
+          for (std::size_t col_data = 0; col_data < params.query_desc->N; col_data++) {
             raft::copy(tmp_batch_buf.data() + (col_data * cur_batch_size),
                        data->ptr + ((col_data * part_n_rows) + total_n_processed),
                        cur_batch_size,
@@ -430,7 +432,7 @@ void perform_local_knn(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
   std::vector<in_t*> ptrs(params.idx_data->size());
   std::vector<int> sizes(params.idx_data->size());
 
-  for (int cur_idx = 0; cur_idx < params.idx_data->size(); cur_idx++) {
+  for (std::size_t cur_idx = 0; cur_idx < params.idx_data->size(); cur_idx++) {
     ptrs[cur_idx]  = params.idx_data->at(cur_idx)->ptr;
     sizes[cur_idx] = work.local_idx_parts[cur_idx]->size;
   }
@@ -487,9 +489,8 @@ __global__ void copy_label_outputs_from_index_parts_kernel(out_t* out,
   if (i >= n_labels) return;
   uint64_t nn_idx = knn_indices[i];
   int part_idx    = 0;
-  for (; part_idx < n_parts && nn_idx >= offsets[part_idx]; part_idx++)
-    ;
-  part_idx        = min(max((int)0, part_idx - 1), n_parts - 1);
+  for (; part_idx < n_parts && nn_idx >= offsets[part_idx]; part_idx++) {}
+  part_idx        = std::min(std::max(0, part_idx - 1), n_parts - 1);
   uint64_t offset = nn_idx - offsets[part_idx];
   out[i]          = parts[part_idx][offset];
 }
@@ -518,14 +519,14 @@ void copy_label_outputs_from_index_parts(opg_knn_param<in_t, ind_t, dist_t, out_
     if (rsp->rank == work.my_rank) { offsets_h.push_back(offset); }
     offset += rsp->size;
   }
-  uint64_t n_parts = offsets_h.size();
+  std::size_t n_parts = offsets_h.size();
   device_buffer<uint64_t> offsets_d(handle.get_device_allocator(), handle.get_stream(), n_parts);
   raft::update_device(offsets_d.data(), offsets_h.data(), n_parts, handle.get_stream());
 
   std::vector<out_t*> parts_h(n_parts);
   device_buffer<out_t*> parts_d(handle.get_device_allocator(), handle.get_stream(), n_parts);
-  for (int o = 0; o < params.n_outputs; o++) {
-    for (int p = 0; p < n_parts; p++) {
+  for (std::size_t o = 0; o < params.n_outputs; o++) {
+    for (std::size_t p = 0; p < n_parts; p++) {
       parts_h[p] = params.y->at(p)[o];
     }
     raft::update_device(parts_d.data(), parts_h.data(), n_parts, handle.get_stream());
@@ -576,7 +577,7 @@ void exchange_results(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
 
     if (params.knn_op != knn_operation::knn) {
       requests.resize(2 + params.n_outputs);
-      for (size_t o = 0; o < params.n_outputs; o++) {
+      for (std::size_t o = 0; o < params.n_outputs; o++) {
         handle.get_comms().isend(work.res.data() + (o * batch_elms),
                                  batch_elms,
                                  part_rank,
@@ -622,7 +623,7 @@ void exchange_results(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
               handle.get_device_allocator(), handle.get_stream(), params.n_outputs * batch_elms);
             raft::copy_async(tmp_res.data(), work.res.data(), tmp_res.size(), handle.get_stream());
 
-            for (int o = 0; o < params.n_outputs; ++o) {
+            for (std::size_t o = 0; o < params.n_outputs; ++o) {
               // Outputs are stored in target order and then in rank order
               raft::copy_async(
                 work.res.data() + (o * work.idxRanks.size() * batch_elms) + batch_offset,
@@ -656,7 +657,7 @@ void exchange_results(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
         ++request_idx;
 
         if (params.knn_op != knn_operation::knn) {
-          for (size_t o = 0; o < params.n_outputs; o++) {
+          for (std::size_t o = 0; o < params.n_outputs; o++) {
             // Outputs are stored in target order and then in rank order
             out_t* r = work.res.data() + (o * work.idxRanks.size() * batch_elms) + batch_offset;
             handle.get_comms().irecv(r, batch_elms, rank, 0, requests.data() + request_idx);
@@ -759,7 +760,7 @@ void reduce(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
       outputs = params.out->at(part_idx)->ptr + (processed_in_part * params.n_outputs);
     } else {
       std::vector<float*>& probas_part = params.probas->at(part_idx);
-      for (int i = 0; i < params.n_outputs; i++) {
+      for (std::size_t i = 0; i < params.n_outputs; i++) {
         float* ptr           = probas_part[i];
         int n_unique_classes = params.n_unique->at(i);
         probas_with_offsets.push_back(ptr + (processed_in_part * n_unique_classes));
@@ -810,9 +811,8 @@ __global__ void merge_labels_kernel(out_t* outputs,
   if (i >= n_labels) return;
   uint64_t nn_idx = knn_indices[i];
   int part_idx    = 0;
-  for (; part_idx < n_parts && nn_idx >= offsets[part_idx]; part_idx++)
-    ;
-  part_idx         = min(max((int)0, part_idx - 1), n_parts - 1);
+  for (; part_idx < n_parts && nn_idx >= offsets[part_idx]; part_idx++) {}
+  part_idx         = std::min(std::max(0, part_idx - 1), n_parts - 1);
   int rank_idx     = parts_to_ranks[part_idx];
   int inbatch_idx  = i / nearest_neighbors;
   uint64_t elm_idx = (rank_idx * n_labels) + inbatch_idx * nearest_neighbors;
@@ -913,7 +913,7 @@ void perform_local_operation(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
 {
   size_t n_labels = batch_size * params.k;
   std::vector<out_t*> y(params.n_outputs);
-  for (int o = 0; o < params.n_outputs; o++) {
+  for (std::size_t o = 0; o < params.n_outputs; o++) {
     y[o] = reinterpret_cast<out_t*>(labels) + (o * n_labels);
   }
 
@@ -953,7 +953,7 @@ void perform_local_operation(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
 {
   size_t n_labels = batch_size * params.k;
   std::vector<out_t*> y(params.n_outputs);
-  for (int o = 0; o < params.n_outputs; o++) {
+  for (std::size_t o = 0; o < params.n_outputs; o++) {
     y[o] = reinterpret_cast<out_t*>(labels) + (o * n_labels);
   }
 
@@ -993,4 +993,4 @@ void perform_local_operation(opg_knn_param<in_t, ind_t, dist_t, out_t>& params,
 };  // namespace knn_common
 };  // namespace opg
 };  // namespace KNN
-};  // namespace ML
\ No newline at end of file
+};  // namespace ML
diff --git a/cpp/src/pca/pca_mg.cu b/cpp/src/pca/pca_mg.cu
index 243b0131f0..4d1a8d364a 100644
--- a/cpp/src/pca/pca_mg.cu
+++ b/cpp/src/pca/pca_mg.cu
@@ -14,23 +14,28 @@
  * limitations under the License.
  */
 
-#include <raft/cudart_utils.h>
-#include <raft/linalg/transpose.h>
+#include "pca.cuh"
+
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/decomposition/pca.hpp>
 #include <cuml/decomposition/pca_mg.hpp>
 #include <cuml/decomposition/sign_flip_mg.hpp>
+
 #include <opg/linalg/qr_based_svd.hpp>
 #include <opg/matrix/matrix_utils.hpp>
 #include <opg/stats/cov.hpp>
 #include <opg/stats/mean.hpp>
 #include <opg/stats/mean_center.hpp>
+
+#include <raft/cudart_utils.h>
+#include <raft/linalg/transpose.h>
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/matrix/math.cuh>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/stats/mean_center.cuh>
-#include "pca.cuh"
+
+#include <cstddef>
 
 using namespace MLCommon;
 
@@ -53,9 +58,8 @@ void fit_impl(raft::handle_t& handle,
               int n_streams,
               bool verbose)
 {
-  const auto& comm             = handle.get_comms();
-  cublasHandle_t cublas_handle = handle.get_cublas_handle();
-  const auto allocator         = handle.get_device_allocator();
+  const auto& comm     = handle.get_comms();
+  const auto allocator = handle.get_device_allocator();
 
   Matrix::Data<T> mu_data{mu, size_t(prms.n_cols)};
 
@@ -222,7 +226,6 @@ void transform_impl(raft::handle_t& handle,
                     int n_streams,
                     bool verbose)
 {
-  cublasHandle_t cublas_h                         = handle.get_cublas_handle();
   const auto allocator                            = handle.get_device_allocator();
   std::vector<Matrix::RankSizePair*> local_blocks = input_desc.partsToRanks;
 
@@ -234,7 +237,7 @@ void transform_impl(raft::handle_t& handle,
       components, singular_vals, prms.n_cols, prms.n_components, true, true, streams[0]);
   }
 
-  for (int i = 0; i < input.size(); i++) {
+  for (std::size_t i = 0; i < input.size(); i++) {
     int si = i % n_streams;
 
     raft::stats::meanCenter(input[i]->ptr,
@@ -361,7 +364,6 @@ void inverse_transform_impl(raft::handle_t& handle,
                             int n_streams,
                             bool verbose)
 {
-  cublasHandle_t cublas_h                         = handle.get_cublas_handle();
   const auto allocator                            = handle.get_device_allocator();
   std::vector<Matrix::RankSizePair*> local_blocks = trans_input_desc.partsToRanks;
 
@@ -373,7 +375,7 @@ void inverse_transform_impl(raft::handle_t& handle,
       components, singular_vals, prms.n_rows, prms.n_components, true, true, streams[0]);
   }
 
-  for (int i = 0; i < local_blocks.size(); i++) {
+  for (std::size_t i = 0; i < local_blocks.size(); i++) {
     int si  = i % n_streams;
     T alpha = T(1);
     T beta  = T(0);
diff --git a/cpp/src/pca/sign_flip_mg.cu b/cpp/src/pca/sign_flip_mg.cu
index bba6dadb66..4f2d65c1bb 100644
--- a/cpp/src/pca/sign_flip_mg.cu
+++ b/cpp/src/pca/sign_flip_mg.cu
@@ -14,17 +14,22 @@
  * limitations under the License.
  */
 
-#include <thrust/device_vector.h>
-#include <thrust/execution_policy.h>
-#include <common/allocatorAdapter.hpp>
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/decomposition/sign_flip_mg.hpp>
+
+#include <common/allocatorAdapter.hpp>
+
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/matrix/math.cuh>
 #include <raft/matrix/matrix.cuh>
 #include <raft/mr/device/allocator.hpp>
 
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+
+#include <cstddef>
+
 using namespace MLCommon;
 
 namespace ML {
@@ -141,7 +146,7 @@ void sign_flip_imp(raft::handle_t& handle,
   device_buffer<T> max_vals(
     allocator, streams[0], std::max(size_t(comm.get_size()), local_blocks.size()) * n_components);
 
-  for (int i = 0; i < input.size(); i++) {
+  for (std::size_t i = 0; i < input.size(); i++) {
     T* mv_loc = max_vals.data() + (i * n_components);
     findMaxAbsOfColumns(
       input[i]->ptr, local_blocks[i]->size, n_components, mv_loc, allocator, streams[i % n_stream]);
@@ -165,7 +170,7 @@ void sign_flip_imp(raft::handle_t& handle,
   findMaxAbsOfColumns(
     max_vals.data(), n_components, comm.get_size(), max_vals.data(), allocator, streams[0], true);
 
-  for (int i = 0; i < local_blocks.size(); i++) {
+  for (std::size_t i = 0; i < local_blocks.size(); i++) {
     flip(input[i]->ptr,
          local_blocks[i]->size,
          n_components,
diff --git a/cpp/src/random_projection/rproj.cuh b/cpp/src/random_projection/rproj.cuh
index 87d2fc7ad7..1ec44fcf9a 100644
--- a/cpp/src/random_projection/rproj.cuh
+++ b/cpp/src/random_projection/rproj.cuh
@@ -16,16 +16,19 @@
 
 #pragma once
 
-#include <random>
-#include <unordered_set>
-#include <vector>
+#include "rproj_utils.cuh"
 
 #include <cuml/random_projection/rproj_c.h>
+
 #include <raft/cudart_utils.h>
 #include <raft/linalg/cublas_wrappers.h>
 #include <raft/sparse/cusparse_wrappers.h>
 #include <raft/cuda_utils.cuh>
-#include "rproj_utils.cuh"
+
+#include <cstddef>
+#include <random>
+#include <unordered_set>
+#include <vector>
 
 namespace ML {
 
@@ -74,16 +77,16 @@ void sparse_random_matrix(const raft::handle_t& h,
   } else {
     auto alloc = h.get_host_allocator();
 
-    size_t indices_alloc = params.n_features * params.n_components * sizeof(int);
-    size_t indptr_alloc  = (params.n_components + 1) * sizeof(int);
-    int* indices         = (int*)alloc->allocate(indices_alloc, stream);
-    int* indptr          = (int*)alloc->allocate(indptr_alloc, stream);
+    std::size_t indices_alloc = params.n_features * params.n_components * sizeof(int);
+    std::size_t indptr_alloc  = (params.n_components + 1) * sizeof(int);
+    int* indices              = (int*)alloc->allocate(indices_alloc, stream);
+    int* indptr               = (int*)alloc->allocate(indptr_alloc, stream);
 
-    size_t offset      = 0;
-    size_t indices_idx = 0;
-    size_t indptr_idx  = 0;
+    std::size_t offset      = 0;
+    std::size_t indices_idx = 0;
+    std::size_t indptr_idx  = 0;
 
-    for (size_t i = 0; i < params.n_components; i++) {
+    for (int i = 0; i < params.n_components; i++) {
       int n_nonzero = binomial(h, params.n_features, params.density, params.random_state);
       sample_without_replacement(params.n_features, n_nonzero, indices, indices_idx);
       indptr[indptr_idx] = offset;
@@ -93,7 +96,7 @@ void sparse_random_matrix(const raft::handle_t& h,
 
     indptr[indptr_idx] = offset;
 
-    size_t len = offset;
+    auto len = offset;
     random_matrix->indices.resize(len, stream);
     raft::update_device(random_matrix->indices.data(), indices, len, stream);
     alloc->deallocate(indices, indices_alloc, stream);
@@ -159,18 +162,18 @@ void RPROJtransform(const raft::handle_t& handle,
     const math_t alfa = 1;
     const math_t beta = 0;
 
-    int& m = params->n_samples;
-    int& n = params->n_components;
-    int& k = params->n_features;
+    auto& m = params->n_samples;
+    auto& n = params->n_components;
+    auto& k = params->n_features;
 
-    int& lda = m;
-    int& ldb = k;
-    int& ldc = m;
+    auto& lda = m;
+    auto& ldb = k;
+    auto& ldc = m;
 
     CUBLAS_CHECK(raft::linalg::cublasgemm(cublas_handle,
                                           CUBLAS_OP_N,
                                           CUBLAS_OP_N,
-                                          m,
+                                          params->n_samples,
                                           n,
                                           k,
                                           &alfa,
@@ -189,13 +192,13 @@ void RPROJtransform(const raft::handle_t& handle,
     const math_t alfa = 1;
     const math_t beta = 0;
 
-    int& m     = params->n_samples;
-    int& n     = params->n_components;
-    int& k     = params->n_features;
-    size_t nnz = random_matrix->sparse_data.size();
+    auto& m         = params->n_samples;
+    auto& n         = params->n_components;
+    auto& k         = params->n_features;
+    std::size_t nnz = random_matrix->sparse_data.size();
 
-    int& lda = m;
-    int& ldc = m;
+    auto& lda = m;
+    auto& ldc = m;
 
     CUSPARSE_CHECK(raft::sparse::cusparsegemmi(cusparse_handle,
                                                m,
diff --git a/cpp/src/random_projection/rproj_utils.cuh b/cpp/src/random_projection/rproj_utils.cuh
index 90e21ec020..0ed3ebcd31 100644
--- a/cpp/src/random_projection/rproj_utils.cuh
+++ b/cpp/src/random_projection/rproj_utils.cuh
@@ -17,11 +17,15 @@
 #pragma once
 
 #include <cuml/random_projection/rproj_c.h>
+
 #include <raft/cudart_utils.h>
-#include <sys/time.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/random/rng.cuh>
 
+#include <sys/time.h>
+
+#include <unordered_set>
+
 const int TPB_X = 256;
 
 inline void sample_without_replacement(size_t n_population,
diff --git a/cpp/src/randomforest/randomforest.cu b/cpp/src/randomforest/randomforest.cu
index 38a13c936f..51caa8f714 100644
--- a/cpp/src/randomforest/randomforest.cu
+++ b/cpp/src/randomforest/randomforest.cu
@@ -13,23 +13,19 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
-#ifdef _OPENMP
-#include <omp.h>
-#else
-#define omp_get_max_threads() 1
-#endif
 
-#include <cuml/ensemble/randomforest.hpp>
+#include "randomforest.cuh"
 
 #include <cuml/tree/flatnode.h>
+#include <cuml/common/logger.hpp>
+#include <cuml/ensemble/randomforest.hpp>
+
 #include <treelite/c_api.h>
 #include <treelite/tree.h>
 
-#include <cuml/common/logger.hpp>
 #include <raft/error.hpp>
 
-#include "randomforest.cuh"
-
+#include <cstddef>
 #include <cstdio>
 #include <cstring>
 #include <fstream>
@@ -368,7 +364,7 @@ void compare_concat_forest_to_subforests(ModelHandle concat_tree_handle,
 {
   size_t concat_forest;
   size_t total_num_trees = 0;
-  for (int forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
+  for (std::size_t forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
     size_t num_trees_each_forest;
     TREELITE_CHECK(TreeliteQueryNumTree(treelite_handles[forest_idx], &num_trees_each_forest));
     total_num_trees = total_num_trees + num_trees_each_forest;
@@ -382,7 +378,7 @@ void compare_concat_forest_to_subforests(ModelHandle concat_tree_handle,
 
   int concat_mod_tree_num = 0;
   tl::Model& concat_model = *(tl::Model*)(concat_tree_handle);
-  for (int forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
+  for (std::size_t forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
     tl::Model& model = *(tl::Model*)(treelite_handles[forest_idx]);
 
     ASSERT(concat_model.GetThresholdType() == model.GetThresholdType(),
@@ -405,7 +401,7 @@ void compare_concat_forest_to_subforests(ModelHandle concat_tree_handle,
       // model_inner is of the concrete type tl::ModelImpl<T, L>
       using model_type         = std::remove_reference_t<decltype(model_inner)>;
       auto& concat_model_inner = dynamic_cast<model_type&>(concat_model);
-      for (int indiv_trees = 0; indiv_trees < model_inner.trees.size(); indiv_trees++) {
+      for (std::size_t indiv_trees = 0; indiv_trees < model_inner.trees.size(); indiv_trees++) {
         compare_trees(concat_model_inner.trees[concat_mod_tree_num + indiv_trees],
                       model_inner.trees[indiv_trees]);
       }
@@ -431,7 +427,7 @@ ModelHandle concatenate_trees(std::vector<ModelHandle> treelite_handles)
     auto* concat_model = dynamic_cast<model_type*>(
       tl::Model::Create(first_model_inner.GetThresholdType(), first_model_inner.GetLeafOutputType())
         .release());
-    for (int forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
+    for (std::size_t forest_idx = 0; forest_idx < treelite_handles.size(); forest_idx++) {
       tl::Model& model  = *(tl::Model*)treelite_handles[forest_idx];
       auto& model_inner = dynamic_cast<model_type&>(model);
       for (const auto& tree : model_inner.trees) {
diff --git a/cpp/src/randomforest/randomforest.cuh b/cpp/src/randomforest/randomforest.cuh
index 1551abd205..1b1e458d2c 100644
--- a/cpp/src/randomforest/randomforest.cuh
+++ b/cpp/src/randomforest/randomforest.cuh
@@ -15,25 +15,33 @@
  */
 
 #pragma once
+
+#include <common/nvtx.hpp>
+
 #include <decisiontree/treelite_util.h>
-#include <raft/cudart_utils.h>
+#include <decisiontree/decisiontree.cuh>
+#include <decisiontree/quantile/quantile.cuh>
+
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/common/logger.hpp>
 #include <cuml/ensemble/randomforest.hpp>
-#include <decisiontree/decisiontree.cuh>
-#include <decisiontree/quantile/quantile.cuh>
-#include <map>
+
 #include <metrics/scores.cuh>
-#include <raft/mr/device/allocator.hpp>
-#include <raft/random/rng.cuh>
 #include <random/permute.cuh>
 
-#include <common/nvtx.hpp>
+#include <raft/cudart_utils.h>
+#include <raft/mr/device/allocator.hpp>
+#include <raft/random/rng.cuh>
 
-#ifndef _OPENMP
-#define omp_get_thread_num() 0
+#ifdef _OPENMP
+#include <omp.h>
+#else
+#define omp_get_thread_num()  0
+#define omp_get_max_threads() 1
 #endif
 
+#include <map>
+
 namespace ML {
 
 template <class T, class L>
diff --git a/cpp/src/solver/cd_mg.cu b/cpp/src/solver/cd_mg.cu
index 9184ad187d..f8fbb713f8 100644
--- a/cpp/src/solver/cd_mg.cu
+++ b/cpp/src/solver/cd_mg.cu
@@ -14,13 +14,18 @@
  * limitations under the License.
  */
 
-#include <raft/cudart_utils.h>
+#include "shuffle.h"
+
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/linear_model/preprocess_mg.hpp>
 #include <cuml/solvers/cd_mg.hpp>
+
 #include <functions/softThres.cuh>
+
 #include <opg/linalg/mv_aTb.hpp>
 #include <opg/linalg/norm.hpp>
+
+#include <raft/cudart_utils.h>
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/linalg/add.cuh>
@@ -31,7 +36,8 @@
 #include <raft/matrix/math.cuh>
 #include <raft/matrix/matrix.cuh>
 #include <raft/mr/device/allocator.hpp>
-#include "shuffle.h"
+
+#include <cstddef>
 
 using namespace MLCommon;
 
@@ -57,14 +63,13 @@ void fit_impl(raft::handle_t& handle,
               int n_streams,
               bool verbose)
 {
-  const auto& comm             = handle.get_comms();
-  cublasHandle_t cublas_handle = handle.get_cublas_handle();
-  const auto allocator         = handle.get_device_allocator();
+  const auto& comm     = handle.get_comms();
+  const auto allocator = handle.get_device_allocator();
 
   std::vector<Matrix::RankSizePair*> partsToRanks = input_desc.blocksOwnedBy(comm.get_rank());
 
   size_t total_M = 0.0;
-  for (int i = 0; i < partsToRanks.size(); i++) {
+  for (std::size_t i = 0; i < partsToRanks.size(); i++) {
     total_M += partsToRanks[i]->size;
   }
 
@@ -105,7 +110,7 @@ void fit_impl(raft::handle_t& handle,
 
   if (comm.get_rank() == 0) {
     ML::Solver::initShuffle(ri, g);
-    for (int i = 0; i < input_desc.N; i++) {
+    for (std::size_t i = 0; i < input_desc.N; i++) {
       ri_h[i] = ri[i];
     }
   }
@@ -132,7 +137,7 @@ void fit_impl(raft::handle_t& handle,
   Matrix::Data<T> coef_loc_data;
 
   T* rs = residual.data();
-  for (int i = 0; i < partsToRanks.size(); i++) {
+  for (std::size_t i = 0; i < partsToRanks.size(); i++) {
     raft::copy(rs, labels[i]->ptr, partsToRanks[i]->size, streams[0]);
 
     Matrix::Data<T>* rs_data = new Matrix::Data<T>();
@@ -151,7 +156,7 @@ void fit_impl(raft::handle_t& handle,
     if (i > 0 && shuffle) {
       if (comm.get_rank() == 0) {
         Solver::shuffle(ri, g);
-        for (int k = 0; k < input_desc.N; k++) {
+        for (std::size_t k = 0; k < input_desc.N; k++) {
           ri_h[k] = ri[k];
         }
       }
@@ -164,7 +169,7 @@ void fit_impl(raft::handle_t& handle,
     T d_coef_max = 0.0;
     T coef_prev  = 0.0;
 
-    for (int j = 0; j < input_desc.N; j++) {
+    for (std::size_t j = 0; j < input_desc.N; j++) {
       int ci         = ri_h[j];
       T* coef_loc    = coef + ci;
       T* squared_loc = squared.data() + ci;
@@ -172,7 +177,7 @@ void fit_impl(raft::handle_t& handle,
       T* pred_loc     = pred.data();
       T* residual_loc = residual.data();
 
-      for (int k = 0; k < input_data.size(); k++) {
+      for (std::size_t k = 0; k < input_data.size(); k++) {
         input_col_loc = input_data[k]->ptr + (ci * partsToRanks[k]->size);
 
         input_data_temp[k]->ptr       = input_col_loc;
@@ -214,7 +219,7 @@ void fit_impl(raft::handle_t& handle,
       pred_loc     = pred.data();
       residual_loc = residual.data();
 
-      for (int k = 0; k < input_data.size(); k++) {
+      for (std::size_t k = 0; k < input_data.size(); k++) {
         input_col_loc = input_data[k]->ptr + (ci * partsToRanks[k]->size);
 
         raft::linalg::multiplyScalar(
@@ -243,7 +248,7 @@ void fit_impl(raft::handle_t& handle,
   CUDA_CHECK(cudaHostUnregister(ri_h));
   free(ri_h);
 
-  for (int i = 0; i < partsToRanks.size(); i++) {
+  for (std::size_t i = 0; i < partsToRanks.size(); i++) {
     delete residual_temp[i];
     delete input_data_temp[i];
   }
@@ -350,7 +355,7 @@ void predict_impl(raft::handle_t& handle,
   T alpha                                         = T(1);
   T beta                                          = T(0);
 
-  for (int i = 0; i < input_data.size(); i++) {
+  for (std::size_t i = 0; i < input_data.size(); i++) {
     int si = i % n_streams;
     raft::linalg::gemm(handle,
                        input_data[i]->ptr,
diff --git a/cpp/src/solver/lars_impl.cuh b/cpp/src/solver/lars_impl.cuh
index c1a2b00616..92b6d99f21 100644
--- a/cpp/src/solver/lars_impl.cuh
+++ b/cpp/src/solver/lars_impl.cuh
@@ -703,10 +703,10 @@ void larsInit(const raft::handle_t& handle,
     *ld_U                   = raft::alignTo<idx_t>(*max_iter, align_bytes);
     try {
       U_buffer.resize((*ld_U) * (*max_iter), stream);
-    } catch (std::bad_alloc) {
+    } catch (std::bad_alloc const&) {
       THROW(
         "Not enough GPU memory! The memory usage depends quadraticaly on the "
-        "n_nonzero_coefs parameter, try to decrease it!");
+        "n_nonzero_coefs parameter, try to decrease it.");
     }
     *U = U_buffer.data();
   } else {
diff --git a/cpp/src/solver/shuffle.h b/cpp/src/solver/shuffle.h
index efd4603a55..c5b70bb3de 100644
--- a/cpp/src/solver/shuffle.h
+++ b/cpp/src/solver/shuffle.h
@@ -16,12 +16,9 @@
 
 #pragma once
 
-#include <stdio.h>
-#include <stdlib.h>
 #include <algorithm>
-#include <iostream>
+#include <cstddef>
 #include <random>
-#include <sstream>
 
 namespace ML {
 namespace Solver {
@@ -30,7 +27,7 @@ template <typename math_t>
 void initShuffle(std::vector<math_t>& rand_indices, std::mt19937& g, math_t random_state = 0)
 {
   g.seed((int)random_state);
-  for (int i = 0; i < rand_indices.size(); ++i)
+  for (std::size_t i = 0; i < rand_indices.size(); ++i)
     rand_indices[i] = i;
 }
 
diff --git a/cpp/src/svm/kernelcache.cuh b/cpp/src/svm/kernelcache.cuh
index 33600afdc7..45aa8ad435 100644
--- a/cpp/src/svm/kernelcache.cuh
+++ b/cpp/src/svm/kernelcache.cuh
@@ -17,16 +17,22 @@
 #pragma once
 
 #include <cuml/svm/svm_parameter.h>
+
 #include <linalg/init.h>
-#include <raft/cudart_utils.h>
 #include <cache/cache.cuh>
 #include <cache/cache_util.cuh>
-#include <cub/cub.cuh>
 #include <matrix/grammatrix.cuh>
+
+#include <raft/cudart_utils.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/linalg/gemm.cuh>
 #include <raft/matrix/matrix.cuh>
 
+#include <cub/cub.cuh>
+
+#include <algorithm>
+#include <cstddef>
+
 namespace ML {
 namespace SVM {
 
@@ -132,12 +138,13 @@ class KernelCache {
     MLCommon::LinAlg::range(k_col_idx.data(), n_ws, stream);
 
     // Init cub buffers
-    size_t bytes1, bytes2;
+    std::size_t bytes1{};
+    std::size_t bytes2{};
     cub::DeviceRadixSort::SortKeys(
       NULL, bytes1, unique_idx.data(), unique_idx.data(), n_ws, 0, sizeof(int) * 8, stream);
     cub::DeviceSelect::Unique(
       NULL, bytes2, unique_idx.data(), unique_idx.data(), d_num_selected_out.data(), n_ws, stream);
-    d_temp_storage_size = max(bytes1, bytes2);
+    d_temp_storage_size = std::max(bytes1, bytes2);
     d_temp_storage.resize(d_temp_storage_size, stream);
   }
 
diff --git a/cpp/src/svm/smosolver.cuh b/cpp/src/svm/smosolver.cuh
index b5231ea050..ddf861c210 100644
--- a/cpp/src/svm/smosolver.cuh
+++ b/cpp/src/svm/smosolver.cuh
@@ -16,7 +16,24 @@
 
 #pragma once
 
-#include <math.h>
+#include "kernelcache.cuh"
+#include "results.cuh"
+#include "smo_sets.cuh"
+#include "smoblocksolve.cuh"
+#include "workingset.cuh"
+#include "ws_util.cuh"
+
+#include <cuml/matrix/kernelparams.h>
+#include <cuml/common/device_buffer.hpp>
+#include <cuml/common/logger.hpp>
+
+#include <matrix/grammatrix.cuh>
+#include <matrix/kernelfactory.cuh>
+
+#include <raft/linalg/cublas_wrappers.h>
+#include <raft/linalg/gemv.h>
+#include <raft/linalg/unary_op.cuh>
+
 #include <raft/cudart_utils.h>
 #include <thrust/device_ptr.h>
 #include <thrust/fill.h>
@@ -26,22 +43,6 @@
 #include <string>
 #include <type_traits>
 
-#include <cuml/matrix/kernelparams.h>
-#include <raft/linalg/cublas_wrappers.h>
-#include <raft/linalg/gemv.h>
-#include <cuml/common/logger.hpp>
-#include <matrix/grammatrix.cuh>
-#include <matrix/kernelfactory.cuh>
-#include <raft/linalg/unary_op.cuh>
-#include "kernelcache.cuh"
-#include "smo_sets.cuh"
-#include "smoblocksolve.cuh"
-#include "workingset.cuh"
-#include "ws_util.cuh"
-
-#include <cuml/common/device_buffer.hpp>
-#include "results.cuh"
-
 namespace ML {
 namespace SVM {
 
@@ -77,7 +78,6 @@ class SmoSolver {
             svmParameter param,
             MLCommon::Matrix::GramMatrixBase<math_t>* kernel)
     : handle(handle),
-      n_rows(n_rows),
       C(param.C),
       tol(param.tol),
       kernel(kernel),
diff --git a/cpp/src/svm/workingset.cuh b/cpp/src/svm/workingset.cuh
index 985bf5c109..cc0147fe36 100644
--- a/cpp/src/svm/workingset.cuh
+++ b/cpp/src/svm/workingset.cuh
@@ -16,26 +16,36 @@
 
 #pragma once
 
+#include "smo_sets.cuh"
+#include "ws_util.cuh"
+
 #include <cuml/svm/svm_parameter.h>
-#include <limits.h>
+#include <cuml/common/device_buffer.hpp>
+#include <cuml/common/logger.hpp>
+
 #include <linalg/init.h>
+
 #include <raft/cudart_utils.h>
-#include <thrust/device_ptr.h>
-#include <thrust/iterator/permutation_iterator.h>
-#include <cub/cub.cuh>
-#include <cuml/common/device_buffer.hpp>
 #include <raft/cuda_utils.cuh>
+#include <raft/handle.hpp>
 #include <raft/linalg/add.cuh>
 #include <raft/linalg/unary_op.cuh>
-#include "smo_sets.cuh"
-#include "ws_util.cuh"
+
+#include <thrust/device_ptr.h>
+#include <thrust/iterator/permutation_iterator.h>
+
+#include <cub/cub.cuh>
+
+#include <algorithm>
+#include <cstddef>
+#include <limits>
 
 namespace ML {
 namespace SVM {
 
 namespace {
-// Unnamed namespace to avoid multiple definition error
-__device__ bool dummy_select_op(int idx) { return true; }
+//  placeholder function passed to configuration call to Cub::DeviceSelect
+__device__ bool always_true(int) { return true; }
 }  // end unnamed namespace
 
 /**
@@ -104,7 +114,7 @@ class WorkingSet {
   void SetSize(int n_train, int n_ws = 0)
   {
     if (n_ws == 0 || n_ws > n_train) { n_ws = n_train; }
-    n_ws       = min(1024, n_ws);
+    n_ws       = std::min(1024, n_ws);
     this->n_ws = n_ws;
     CUML_LOG_DEBUG("Creating working set with %d elements", n_ws);
     AllocateBuffers();
@@ -328,8 +338,8 @@ class WorkingSet {
   MLCommon::device_buffer<int> ws_priority;
   MLCommon::device_buffer<int> ws_priority_sorted;
 
-  int* d_num_selected = nullptr;
-  size_t cub_bytes    = 0;
+  int* d_num_selected   = nullptr;
+  std::size_t cub_bytes = 0;
   MLCommon::device_buffer<char> cub_storage;
 
   void AllocateBuffers()
@@ -352,7 +362,7 @@ class WorkingSet {
       d_num_selected = (int*)handle.get_device_allocator()->allocate(1 * sizeof(int), stream);
 
       // Determine temporary device storage requirements for cub
-      size_t cub_bytes2 = 0;
+      std::size_t cub_bytes2 = 0;
       cub::DeviceRadixSort::SortPairs(NULL,
                                       cub_bytes,
                                       f_sorted.data(),
@@ -363,15 +373,9 @@ class WorkingSet {
                                       0,
                                       8 * sizeof(math_t),
                                       stream);
-      cub::DeviceSelect::If(NULL,
-                            cub_bytes2,
-                            f_idx.data(),
-                            f_idx.data(),
-                            d_num_selected,
-                            n_train,
-                            dummy_select_op,
-                            stream);
-      cub_bytes = max(cub_bytes, cub_bytes2);
+      cub::DeviceSelect::If(
+        NULL, cub_bytes2, f_idx.data(), f_idx.data(), d_num_selected, n_train, always_true, stream);
+      cub_bytes = std::max(cub_bytes, cub_bytes2);
       cub_storage.resize(cub_bytes, stream);
       Initialize();
     }
diff --git a/cpp/src/tsvd/tsvd_mg.cu b/cpp/src/tsvd/tsvd_mg.cu
index 0d8e1a7948..93fe41de5d 100644
--- a/cpp/src/tsvd/tsvd_mg.cu
+++ b/cpp/src/tsvd/tsvd_mg.cu
@@ -14,22 +14,27 @@
  * limitations under the License.
  */
 
-#include <raft/cudart_utils.h>
+#include "tsvd.cuh"
+
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/decomposition/sign_flip_mg.hpp>
 #include <cuml/decomposition/tsvd.hpp>
 #include <cuml/decomposition/tsvd_mg.hpp>
+
 #include <opg/linalg/mm_aTa.hpp>
 #include <opg/stats/mean.hpp>
 #include <opg/stats/mean_center.hpp>
 #include <opg/stats/stddev.hpp>
+
+#include <raft/cudart_utils.h>
 #include <raft/comms/comms.hpp>
 #include <raft/cuda_utils.cuh>
 #include <raft/linalg/eltwise.cuh>
 #include <raft/matrix/math.cuh>
 #include <raft/mr/device/allocator.hpp>
 #include <raft/stats/mean_center.cuh>
-#include "tsvd.cuh"
+
+#include <cstddef>
 
 using namespace MLCommon;
 
@@ -135,12 +140,11 @@ void transform_impl(raft::handle_t& handle,
 {
   int rank = handle.get_comms().get_rank();
 
-  cublasHandle_t cublas_h = handle.get_cublas_handle();
-  const auto allocator    = handle.get_device_allocator();
+  const auto allocator = handle.get_device_allocator();
 
   std::vector<Matrix::RankSizePair*> local_blocks = input_desc.blocksOwnedBy(rank);
 
-  for (int i = 0; i < input.size(); i++) {
+  for (std::size_t i = 0; i < input.size(); i++) {
     int si = i % n_streams;
 
     T alpha = T(1);
@@ -223,11 +227,10 @@ void inverse_transform_impl(raft::handle_t& handle,
                             int n_streams,
                             bool verbose)
 {
-  cublasHandle_t cublas_h                         = handle.get_cublas_handle();
   const auto allocator                            = handle.get_device_allocator();
   std::vector<Matrix::RankSizePair*> local_blocks = trans_input_desc.partsToRanks;
 
-  for (int i = 0; i < local_blocks.size(); i++) {
+  for (std::size_t i = 0; i < local_blocks.size(); i++) {
     int si  = i % n_streams;
     T alpha = T(1);
     T beta  = T(0);
diff --git a/cpp/src/umap/simpl_set_embed/optimize_batch_kernel.cuh b/cpp/src/umap/simpl_set_embed/optimize_batch_kernel.cuh
index c20f8d8dc5..bb330a7436 100644
--- a/cpp/src/umap/simpl_set_embed/optimize_batch_kernel.cuh
+++ b/cpp/src/umap/simpl_set_embed/optimize_batch_kernel.cuh
@@ -16,12 +16,16 @@
 
 #pragma once
 
+#include <common/fast_int_div.cuh>
+
 #include <cuml/manifold/umapparams.h>
+
 #include <raft/cudart_utils.h>
-#include <common/fast_int_div.cuh>
 #include <raft/cuda_utils.cuh>
 #include <raft/random/rng.cuh>
 
+#include <cstddef>
+
 namespace UMAPAlgo {
 namespace SimplSetEmbed {
 namespace Algo {
@@ -368,9 +372,9 @@ void call_optimize_batch_kernel(T const* head_embedding,
                                 cudaStream_t& stream,
                                 T rounding)
 {
-  size_t requiredSize = TPB_X * params->n_components;
+  std::size_t requiredSize = TPB_X * params->n_components;
   requiredSize *= sizeof(T);
-  bool use_shared_mem = requiredSize < raft::getSharedMemPerBlock();
+  bool use_shared_mem = requiredSize < static_cast<std::size_t>(raft::getSharedMemPerBlock());
   T nsr_inv           = T(1.0) / params->negative_sample_rate;
   if (params->n_components == 2) {
     // multicore implementation with registers
diff --git a/cpp/src_prims/cache/cache.cuh b/cpp/src_prims/cache/cache.cuh
index cf815f5d4c..076b246737 100644
--- a/cpp/src_prims/cache/cache.cuh
+++ b/cpp/src_prims/cache/cache.cuh
@@ -16,13 +16,18 @@
 
 #pragma once
 
-#include <raft/cudart_utils.h>
-#include <cub/cub.cuh>
+#include "cache_util.cuh"
+
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/common/logger.hpp>
+
+#include <raft/cudart_utils.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/mr/device/allocator.hpp>
-#include "cache_util.cuh"
+
+#include <cub/cub.cuh>
+
+#include <cstddef>
 
 namespace MLCommon {
 namespace Cache {
@@ -390,7 +395,7 @@ class Cache {
 
   void ResizeTmpBuffers(int n, cudaStream_t stream)
   {
-    if (ws_tmp.size() < n) {
+    if (ws_tmp.size() < static_cast<std::size_t>(n)) {
       ws_tmp.resize(n, stream);
       is_cached.resize(n, stream);
       idx_tmp.resize(n, stream);
diff --git a/cpp/src_prims/cache/cache_util.cuh b/cpp/src_prims/cache/cache_util.cuh
index 2ad4ed0bd7..7c18f6f0d5 100644
--- a/cpp/src_prims/cache/cache_util.cuh
+++ b/cpp/src_prims/cache/cache_util.cuh
@@ -16,10 +16,12 @@
 
 #pragma once
 
-#include <cub/cub.cuh>
-#include <raft/cuda_utils.cuh>
 #include <selection/kselection.cuh>
 
+#include <raft/cuda_utils.cuh>
+
+#include <cub/cub.cuh>
+
 namespace MLCommon {
 namespace Cache {
 
diff --git a/cpp/src_prims/linalg/batched/matrix.cuh b/cpp/src_prims/linalg/batched/matrix.cuh
index 74877f754d..5ed44b7445 100644
--- a/cpp/src_prims/linalg/batched/matrix.cuh
+++ b/cpp/src_prims/linalg/batched/matrix.cuh
@@ -16,6 +16,23 @@
 
 #pragma once
 
+#include <cuml/common/device_buffer.hpp>
+#include <cuml/common/utils.hpp>
+
+#include <common/fast_int_div.cuh>
+
+#include <raft/cudart_utils.h>
+#include <raft/linalg/cublas_wrappers.h>
+#include <raft/cuda_utils.cuh>
+#include <raft/linalg/binary_op.cuh>
+#include <raft/linalg/unary_op.cuh>
+#include <raft/mr/device/allocator.hpp>
+
+#include <thrust/execution_policy.h>
+#include <thrust/for_each.h>
+#include <thrust/iterator/counting_iterator.h>
+
+#include <cstddef>
 #include <functional>
 #include <memory>
 #include <stdexcept>
@@ -24,21 +41,6 @@
 #include <unordered_map>
 #include <vector>
 
-#include <thrust/execution_policy.h>
-#include <thrust/for_each.h>
-#include <thrust/iterator/counting_iterator.h>
-
-#include <raft/cudart_utils.h>
-#include <common/fast_int_div.cuh>
-#include <cuml/common/device_buffer.hpp>
-#include <cuml/common/utils.hpp>
-#include <raft/cuda_utils.cuh>
-#include <raft/mr/device/allocator.hpp>
-
-#include <raft/linalg/cublas_wrappers.h>
-#include <raft/linalg/binary_op.cuh>
-#include <raft/linalg/unary_op.cuh>
-
 namespace MLCommon {
 namespace LinAlg {
 namespace Batched {
@@ -57,6 +59,7 @@ __global__ void identity_matrix_kernel(T* I, int m)
 {
   T* I_b     = I + blockIdx.x * m * m;
   int stride = (m + 1);
+
   for (int idx = threadIdx.x; idx < m; idx += blockDim.x) {
     I_b[idx * stride] = 1;
   }
@@ -155,6 +158,8 @@ class Matrix {
   }
 
  public:
+  using shape_type = std::pair<std::size_t, std::size_t>;
+
   /**
    * @brief Constructor that allocates memory using the memory pool.
    *
@@ -166,9 +171,9 @@ class Matrix {
    * @param[in]  stream       CUDA stream
    * @param[in]  setZero      Should matrix be zeroed on allocation?
    */
-  Matrix(int m,
-         int n,
-         int batch_size,
+  Matrix(std::size_t m,
+         std::size_t n,
+         std::size_t batch_size,
          cublasHandle_t cublasHandle,
          std::shared_ptr<raft::mr::device::allocator> allocator,
          cudaStream_t stream,
@@ -202,9 +207,9 @@ class Matrix {
    * @param[in]  stream       CUDA stream
    * @param[in]  setZero      Should matrix be zeroed on allocation?
    */
-  Matrix(int m,
-         int n,
-         int batch_size,
+  Matrix(std::size_t m,
+         std::size_t n,
+         std::size_t batch_size,
          cublasHandle_t cublasHandle,
          T** d_batches,
          T* d_dense,
@@ -268,7 +273,7 @@ class Matrix {
   }
 
   //! Return batches
-  int batches() const { return m_batch_size; }
+  std::size_t batches() const { return m_batch_size; }
 
   //! Return cublas handle
   cublasHandle_t cublasHandle() const { return m_cublasHandle; }
@@ -280,7 +285,7 @@ class Matrix {
   cudaStream_t stream() const { return m_stream; }
 
   //! Return shape
-  const std::pair<int, int>& shape() const { return m_shape; }
+  const shape_type& shape() const { return m_shape; }
 
   //! Return array of pointers to the offsets in the data buffer
   const T** data() const { return d_batches; }
@@ -305,19 +310,17 @@ class Matrix {
    * @param[in]  m  Number of desired rows
    * @param[in]  n  Number of desired columns
    */
-  void reshape(int m, int n)
+  void reshape(std::size_t m, std::size_t n)
   {
-    const int r = m_shape.first * m_shape.second;
+    const auto r = m_shape.first * m_shape.second;
     ASSERT(r == m * n, "ERROR: Size mismatch - Cannot reshape matrix into desired shape");
-    m_shape = std::pair<int, int>(m, n);
+    m_shape = shape_type(m, n);
   }
 
   //! Stack the matrix by columns creating a long vector
   Matrix<T> vec() const
   {
-    int m = m_shape.first;
-    int n = m_shape.second;
-    int r = m * n;
+    const auto r = m_shape.first * m_shape.second;
     Matrix<T> toVec(r, 1, m_batch_size, m_cublasHandle, m_allocator, m_stream, false);
     raft::copy(toVec[0], raw_data(), m_batch_size * r, m_stream);
     return toVec;
@@ -332,7 +335,7 @@ class Matrix {
    */
   Matrix<T> mat(int m, int n) const
   {
-    const int r = m_shape.first * m_shape.second;
+    const auto r = m_shape.first * m_shape.second;
     ASSERT(r == m * n, "ERROR: Size mismatch - Cannot reshape array into desired size");
     Matrix<T> toMat(m, n, m_batch_size, m_cublasHandle, m_allocator, m_stream, false);
     raft::copy(toMat[0], raw_data(), m_batch_size * r, m_stream);
@@ -343,7 +346,7 @@ class Matrix {
   //! Visualize the first matrix.
   void print(std::string name) const
   {
-    size_t len = m_shape.first * m_shape.second * m_batch_size;
+    std::size_t len = m_shape.first * m_shape.second * m_batch_size;
     std::vector<T> A(len);
     raft::update_host(A.data(), raw_data(), len, m_stream);
     std::cout << name << "=\n";
@@ -440,8 +443,8 @@ class Matrix {
    */
   Matrix<T> transpose() const
   {
-    int m = m_shape.first;
-    int n = m_shape.second;
+    auto m = m_shape.first;
+    auto n = m_shape.second;
 
     Matrix<T> At(n, m, m_batch_size, m_cublasHandle, m_allocator, m_stream);
 
@@ -476,22 +479,23 @@ class Matrix {
    *
    * @return A batched identity matrix
    */
-  static Matrix<T> Identity(int m,
-                            int batch_size,
+  static Matrix<T> Identity(std::size_t m,
+                            std::size_t batch_size,
                             cublasHandle_t cublasHandle,
                             std::shared_ptr<raft::mr::device::allocator> allocator,
                             cudaStream_t stream)
   {
     Matrix<T> I(m, m, batch_size, cublasHandle, allocator, stream, true);
 
-    identity_matrix_kernel<T><<<batch_size, std::min(256, m), 0, stream>>>(I.raw_data(), m);
+    identity_matrix_kernel<T>
+      <<<batch_size, std::min(std::size_t{256}, m), 0, stream>>>(I.raw_data(), m);
     CUDA_CHECK(cudaPeekAtLastError());
     return I;
   }
 
  protected:
   //! Shape (rows, cols) of matrices. We assume all matrices in batch have same shape.
-  std::pair<int, int> m_shape;
+  shape_type m_shape;
 
   //! Pointers to each matrix in the contiguous data buffer (strided offsets)
   device_buffer<T*> m_batches;
@@ -502,7 +506,7 @@ class Matrix {
   T* d_dense;  // When pre-allocated
 
   //! Number of matrices in batch
-  int m_batch_size;
+  std::size_t m_batch_size;
 
   std::shared_ptr<raft::mr::device::allocator> m_allocator;
   cublasHandle_t m_cublasHandle;
@@ -567,9 +571,9 @@ __global__ void kronecker_product_kernel(
 template <typename T>
 void b_gemm(bool aT,
             bool bT,
-            int m,
-            int n,
-            int k,
+            std::size_t m,
+            std::size_t n,
+            std::size_t k,
             T alpha,
             const Matrix<T>& A,
             const Matrix<T>& B,
@@ -580,10 +584,10 @@ void b_gemm(bool aT,
   {
     ASSERT(A.batches() == B.batches(), "A and B must have the same number of batches");
     ASSERT(A.batches() == C.batches(), "A and C must have the same number of batches");
-    int Arows = !aT ? A.shape().first : A.shape().second;
-    int Acols = !aT ? A.shape().second : A.shape().first;
-    int Brows = !bT ? B.shape().first : B.shape().second;
-    int Bcols = !bT ? B.shape().second : B.shape().first;
+    auto Arows = !aT ? A.shape().first : A.shape().second;
+    auto Acols = !aT ? A.shape().second : A.shape().first;
+    auto Brows = !bT ? B.shape().first : B.shape().second;
+    auto Bcols = !bT ? B.shape().second : B.shape().first;
     ASSERT(m <= Arows, "m should be <= number of rows of A");
     ASSERT(k <= Acols, "k should be <= number of columns of A");
     ASSERT(k <= Brows, "k should be <= number of rows of B");
@@ -665,11 +669,11 @@ template <typename T>
 void b_gels(const Matrix<T>& A, Matrix<T>& C)
 {
   ASSERT(A.batches() == C.batches(), "A and C must have the same number of batches");
-  int m = A.shape().first;
+  auto m = A.shape().first;
   ASSERT(C.shape().first == m, "Dimension mismatch: A rows, C rows");
-  int n = A.shape().second;
+  auto n = A.shape().second;
   ASSERT(m > n, "Only overdetermined systems (m > n) are supported");
-  int nrhs = C.shape().second;
+  auto nrhs = C.shape().second;
 
   Matrix<T> Acopy(A);
 
@@ -700,8 +704,8 @@ template <typename T, typename F>
 Matrix<T> b_op_A(const Matrix<T>& A, F unary_op)
 {
   auto batch_size = A.batches();
-  int m           = A.shape().first;
-  int n           = A.shape().second;
+  auto m          = A.shape().first;
+  auto n          = A.shape().second;
 
   Matrix<T> C(m, n, batch_size, A.cublasHandle(), A.allocator(), A.stream());
 
@@ -728,8 +732,8 @@ Matrix<T> b_aA_op_B(const Matrix<T>& A, const Matrix<T>& B, F binary_op)
   ASSERT(A.batches() == B.batches(), "A & B must have same number of batches");
 
   auto batch_size = A.batches();
-  int m           = A.shape().first;
-  int n           = A.shape().second;
+  auto m          = A.shape().first;
+  auto n          = A.shape().second;
 
   Matrix<T> C(m, n, batch_size, A.cublasHandle(), A.allocator(), A.stream());
 
@@ -818,20 +822,20 @@ Matrix<T> b_solve(const Matrix<T>& A, const Matrix<T>& b)
 template <typename T>
 void b_kron(const Matrix<T>& A, const Matrix<T>& B, Matrix<T>& AkB, T alpha = (T)1)
 {
-  int m = A.shape().first;
-  int n = A.shape().second;
+  auto m = A.shape().first;
+  auto n = A.shape().second;
 
-  int p = B.shape().first;
-  int q = B.shape().second;
+  auto p = B.shape().first;
+  auto q = B.shape().second;
 
   // Resulting shape
-  int k_m = m * p;
-  int k_n = n * q;
+  auto k_m = m * p;
+  auto k_n = n * q;
   ASSERT(AkB.shape().first == k_m, "Kronecker product output dimensions mismatch");
   ASSERT(AkB.shape().second == k_n, "Kronecker product output dimensions mismatch");
 
   // Run kronecker
-  dim3 threads(std::min(p, 32), std::min(q, 32));
+  dim3 threads(std::min(p, std::size_t{32}), std::min(q, std::size_t{32}));
   kronecker_product_kernel<T><<<A.batches(), threads, 0, A.stream()>>>(
     A.raw_data(), m, n, B.raw_data(), p, q, AkB.raw_data(), k_m, k_n, alpha);
   CUDA_CHECK(cudaPeekAtLastError());
@@ -848,15 +852,15 @@ void b_kron(const Matrix<T>& A, const Matrix<T>& B, Matrix<T>& AkB, T alpha = (T
 template <typename T>
 Matrix<T> b_kron(const Matrix<T>& A, const Matrix<T>& B)
 {
-  int m = A.shape().first;
-  int n = A.shape().second;
+  auto m = A.shape().first;
+  auto n = A.shape().second;
 
-  int p = B.shape().first;
-  int q = B.shape().second;
+  auto p = B.shape().first;
+  auto q = B.shape().second;
 
   // Resulting shape
-  int k_m = m * p;
-  int k_n = n * q;
+  auto k_m = m * p;
+  auto k_n = n * q;
 
   Matrix<T> AkB(k_m, k_n, A.batches(), A.cublasHandle(), A.allocator(), A.stream());
 
@@ -1040,12 +1044,12 @@ static __global__ void batched_2dcopy_kernel(const T* in,
 template <typename T>
 void b_2dcopy(const Matrix<T>& in,
               Matrix<T>& out,
-              int in_starting_row,
-              int in_starting_col,
-              int copy_rows,
-              int copy_cols,
-              int out_starting_row = 0,
-              int out_starting_col = 0)
+              std::size_t in_starting_row,
+              std::size_t in_starting_col,
+              std::size_t copy_rows,
+              std::size_t copy_cols,
+              std::size_t out_starting_row = 0,
+              std::size_t out_starting_col = 0)
 {
   ASSERT(in_starting_row + copy_rows <= in.shape().first,
          "[2D copy] Dimension mismatch: rows for input matrix");
@@ -1057,7 +1061,7 @@ void b_2dcopy(const Matrix<T>& in,
          "[2D copy] Dimension mismatch: columns for output matrix");
 
   // Execute the kernel
-  const int TPB = copy_rows * copy_cols > 512 ? 256 : 128;  // quick heuristics
+  const int TPB = copy_rows * copy_cols > std::size_t{512} ? 256 : 128;  // quick heuristics
   batched_2dcopy_kernel<<<in.batches(), TPB, 0, in.stream()>>>(in.raw_data(),
                                                                out.raw_data(),
                                                                in_starting_row,
diff --git a/cpp/src_prims/linalg/init.h b/cpp/src_prims/linalg/init.h
index 763bcb665a..842edce5b1 100644
--- a/cpp/src_prims/linalg/init.h
+++ b/cpp/src_prims/linalg/init.h
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <raft/cudart_utils.h>
+
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
 #include <thrust/iterator/counting_iterator.h>
diff --git a/cpp/src_prims/linalg/reduce_rows_by_key.cuh b/cpp/src_prims/linalg/reduce_rows_by_key.cuh
index b883532950..0fe50894f3 100644
--- a/cpp/src_prims/linalg/reduce_rows_by_key.cuh
+++ b/cpp/src_prims/linalg/reduce_rows_by_key.cuh
@@ -16,13 +16,13 @@
 
 #pragma once
 
-#define MAX_BLOCKS 65535u
+#include <raft/cuda_utils.cuh>
 
-#include <stdlib.h>
 #include <cub/cub.cuh>
+
 #include <limits>
-#include <raft/cuda_utils.cuh>
 
+#define MAX_BLOCKS 65535u
 namespace MLCommon {
 namespace LinAlg {
 
@@ -100,8 +100,8 @@ __launch_bounds__(SUM_ROWS_SMALL_K_DIMX, 4) __global__
   typedef cub::BlockReduce<quad<DataType>, SUM_ROWS_SMALL_K_DIMX> BlockReduce;
   __shared__ typename BlockReduce::TempStorage temp_storage;
 
-  for (int idim = blockIdx.y; idim < ncols; idim += gridDim.y) {
-    if (idim != blockIdx.y) __syncthreads();  // we're reusing temp_storage
+  for (int idim = static_cast<int>(blockIdx.y); idim < ncols; idim += gridDim.y) {
+    if (idim != static_cast<int>(blockIdx.y)) __syncthreads();  // we're reusing temp_storage
 
     // threadIdx.x stores partial sum for current dim and key=threadIdx.x in this reg
     quad<DataType> thread_sums;
@@ -139,7 +139,7 @@ __launch_bounds__(SUM_ROWS_SMALL_K_DIMX, 4) __global__
     if (threadIdx.x < 32) {
       // We only need 4
       thread_sums = cub::ShuffleIndex<32>(thread_sums, 0, 0xffffffff);
-      if (threadIdx.x < nkeys) {
+      if (static_cast<int>(threadIdx.x) < nkeys) {
         if (threadIdx.x == 0) raft::myAtomicAdd(&d_sums[threadIdx.x * ncols + idim], thread_sums.x);
         if (threadIdx.x == 1) raft::myAtomicAdd(&d_sums[threadIdx.x * ncols + idim], thread_sums.y);
         if (threadIdx.x == 2) raft::myAtomicAdd(&d_sums[threadIdx.x * ncols + idim], thread_sums.z);
@@ -382,7 +382,8 @@ void reduce_rows_by_key(const DataIteratorT d_A,
     sum_rows_by_key_small_nkeys(
       d_A, lda, d_keys_char, d_weights, nrows, ncols, nkeys, d_sums, stream);
   } else {
-    for (KeyType key_offset = 0; key_offset < nkeys; key_offset += SUM_ROWS_BY_KEY_LARGE_K_MAX_K) {
+    for (KeyType key_offset = 0; key_offset < static_cast<KeyType>(nkeys);
+         key_offset += SUM_ROWS_BY_KEY_LARGE_K_MAX_K) {
       KeyType this_call_nkeys = std::min(SUM_ROWS_BY_KEY_LARGE_K_MAX_K, nkeys);
       sum_rows_by_key_large_nkeys_rowmajor(
         d_A, lda, d_keys, d_weights, nrows, ncols, key_offset, this_call_nkeys, d_sums, stream);
diff --git a/cpp/src_prims/metrics/batched/information_criterion.cuh b/cpp/src_prims/metrics/batched/information_criterion.cuh
index 87e9a42302..a4a61ce7a1 100644
--- a/cpp/src_prims/metrics/batched/information_criterion.cuh
+++ b/cpp/src_prims/metrics/batched/information_criterion.cuh
@@ -25,11 +25,10 @@
  *  - BIC: https://en.wikipedia.org/wiki/Bayesian_information_criterion
  */
 
-#include <cuda_runtime.h>
-#include <cmath>
-
 #include <raft/linalg/unary_op.cuh>
 
+#include <cmath>
+
 namespace MLCommon {
 namespace Metrics {
 
@@ -62,7 +61,7 @@ void information_criterion(ScalarT* d_ic,
                            IdxT n_samples,
                            cudaStream_t stream)
 {
-  ScalarT ic_base;
+  ScalarT ic_base{};
   ScalarT N = static_cast<ScalarT>(n_params);
   ScalarT T = static_cast<ScalarT>(n_samples);
   switch (ic_type) {
diff --git a/cpp/src_prims/metrics/contingencyMatrix.cuh b/cpp/src_prims/metrics/contingencyMatrix.cuh
index e1080ecc6f..0091234611 100644
--- a/cpp/src_prims/metrics/contingencyMatrix.cuh
+++ b/cpp/src_prims/metrics/contingencyMatrix.cuh
@@ -16,12 +16,15 @@
 
 #pragma once
 
-#include <math.h>
 #include <raft/cudart_utils.h>
+#include <raft/cuda_utils.cuh>
+
 #include <thrust/device_ptr.h>
 #include <thrust/reduce.h>
+
 #include <cub/cub.cuh>
-#include <raft/cuda_utils.cuh>
+
+#include <math.h>
 
 namespace MLCommon {
 namespace Metrics {
@@ -41,7 +44,7 @@ __global__ void devConstructContingencyMatrix(const T* groundTruth,
                                               int outIdxOffset,
                                               int outMatWidth)
 {
-  auto elementId = threadIdx.x + blockDim.x * blockIdx.x;
+  int elementId = threadIdx.x + blockDim.x * blockIdx.x;
   if (elementId < nSamples) {
     T gt           = groundTruth[elementId];
     T pd           = predicted[elementId];
@@ -77,11 +80,11 @@ __global__ void devConstructContingencyMatrixSmem(const T* groundTruth,
 {
   extern __shared__ char smem[];
   auto* sMemMatrix = reinterpret_cast<OutT*>(smem);
-  for (auto smemIdx = threadIdx.x; smemIdx < outMatWidth * outMatWidth; smemIdx += blockDim.x) {
+  for (int smemIdx = threadIdx.x; smemIdx < outMatWidth * outMatWidth; smemIdx += blockDim.x) {
     sMemMatrix[smemIdx] = 0;
   }
   __syncthreads();
-  auto elementId = threadIdx.x + blockDim.x * blockIdx.x;
+  int elementId = threadIdx.x + blockDim.x * blockIdx.x;
   if (elementId < nSamples) {
     T gt           = groundTruth[elementId];
     T pd           = predicted[elementId];
@@ -89,7 +92,7 @@ __global__ void devConstructContingencyMatrixSmem(const T* groundTruth,
     raft::myAtomicAdd(sMemMatrix + outputIdx, OutT(1));
   }
   __syncthreads();
-  for (auto smemIdx = threadIdx.x; smemIdx < outMatWidth * outMatWidth; smemIdx += blockDim.x) {
+  for (int smemIdx = threadIdx.x; smemIdx < outMatWidth * outMatWidth; smemIdx += blockDim.x) {
     raft::myAtomicAdd(outMat + smemIdx, sMemMatrix[smemIdx]);
   }
 }
@@ -158,9 +161,9 @@ ContingencyMatrixImplType getImplVersion(OutT outDimN)
   ContingencyMatrixImplType implVersion = IMPL_NONE;
   // keeping 8 block per SM to get good utilization
   // can go higher but reduced L1 size degrades perf
-  int upperLimitSmemAtomics =
+  OutT upperLimitSmemAtomics =
     std::floor(std::sqrt(maxSmemPerBlock / (sizeof(OutT) * (maxBlocksResidentPerSM / 2))));
-  int upperLimitL2Atomics = std::floor(std::sqrt(l2CacheSize / sizeof(OutT)));
+  OutT upperLimitL2Atomics = std::floor(std::sqrt(l2CacheSize / sizeof(OutT)));
   if (outDimN <= upperLimitSmemAtomics)
     implVersion = SMEM_ATOMICS;
   else if (outDimN <= upperLimitL2Atomics)
@@ -215,12 +218,12 @@ size_t getContingencyMatrixWorkspaceSize(int nSamples,
   auto outDimN                          = OutT(maxLabel - minLabel + 1);
   ContingencyMatrixImplType implVersion = getImplVersion<OutT>(outDimN);
   if (implVersion == SORT_AND_GATOMICS) {
-    void* pWorkspaceCub    = nullptr;
+    void* pWorkspaceCub{};
     size_t tmpStorageBytes = 0;
-    // bunch of no-op pointers to get workspace size
-    T *pTmpKey, *pTmpValue, *pTmpKeyOut, *pTmpValueOut;
+    // no-op pointers to get workspace size
+    T* pTmpUnused{};
     CUDA_CHECK(cub::DeviceRadixSort::SortPairs(
-      pWorkspaceCub, tmpStorageBytes, pTmpKey, pTmpValue, pTmpKeyOut, pTmpValueOut, nSamples));
+      pWorkspaceCub, tmpStorageBytes, pTmpUnused, pTmpUnused, pTmpUnused, pTmpUnused, nSamples));
     auto tmpStagingMemorySize = raft::alignTo<size_t>(nSamples * sizeof(T), 256);
     tmpStagingMemorySize *= 2;
     workspaceSize = tmpStagingMemorySize + tmpStorageBytes;
@@ -300,6 +303,7 @@ void contingencyMatrix(const T* groundTruth,
                                       workspaceSize,
                                       stream);
       break;
+    case IMPL_NONE: break;
   }
 }
 
diff --git a/cpp/src_prims/selection/haversine_knn.cuh b/cpp/src_prims/selection/haversine_knn.cuh
index 58df0ec99d..c531c3577b 100644
--- a/cpp/src_prims/selection/haversine_knn.cuh
+++ b/cpp/src_prims/selection/haversine_knn.cuh
@@ -16,6 +16,8 @@
 
 #pragma once
 
+#include <raft/cuda_utils.cuh>
+
 #include <faiss/gpu/utils/Limits.cuh>
 #include <faiss/gpu/utils/Select.cuh>
 
@@ -131,4 +133,4 @@ void haversine_knn(value_idx* out_inds,
 }
 
 };  // namespace selection
-};  // namespace raft
\ No newline at end of file
+};  // namespace raft
diff --git a/cpp/src_prims/selection/knn.cuh b/cpp/src_prims/selection/knn.cuh
index c1eb9e2b21..66d788dfc1 100644
--- a/cpp/src_prims/selection/knn.cuh
+++ b/cpp/src_prims/selection/knn.cuh
@@ -16,11 +16,19 @@
 
 #pragma once
 
-#include <raft/cudart_utils.h>
-#include <raft/cuda_utils.cuh>
+#include "haversine_knn.cuh"
+#include "processing.cuh"
 
 #include <label/classlabels.cuh>
+
+#include <cuml/common/device_buffer.hpp>
+#include <cuml/neighbors/knn.hpp>
+
+#include <raft/cudart_utils.h>
+#include <raft/linalg/distance_type.h>
+#include <raft/cuda_utils.cuh>
 #include <raft/distance/distance.cuh>
+#include <raft/mr/device/allocator.hpp>
 
 #include <faiss/gpu/GpuDistance.h>
 #include <faiss/gpu/GpuIndexFlat.h>
@@ -37,15 +45,7 @@
 #include <thrust/device_vector.h>
 #include <thrust/iterator/transform_iterator.h>
 
-#include <raft/linalg/distance_type.h>
-#include "processing.cuh"
-
-#include "haversine_knn.cuh"
-
-#include <cuml/common/device_buffer.hpp>
-#include <cuml/neighbors/knn.hpp>
-#include <raft/mr/device/allocator.hpp>
-
+#include <cstddef>
 #include <iostream>
 #include <set>
 
@@ -68,7 +68,7 @@ __global__ void class_probs_kernel(OutType* out,
                                    const int64_t* knn_indices,
                                    const int* labels,
                                    int n_uniq_labels,
-                                   size_t n_samples,
+                                   std::size_t n_samples,
                                    int n_neighbors)
 {
   int row = (blockIdx.x * blockDim.x) + threadIdx.x;
@@ -90,7 +90,7 @@ __global__ void class_vote_kernel(OutType* out,
                                   const float* class_proba,
                                   int* unique_labels,
                                   int n_uniq_labels,
-                                  size_t n_samples,
+                                  std::size_t n_samples,
                                   int n_outputs,
                                   int output_offset,
                                   bool use_shared_mem)
@@ -127,7 +127,7 @@ template <typename LabelType, bool precomp_lbls = false>
 __global__ void regress_avg_kernel(LabelType* out,
                                    const int64_t* knn_indices,
                                    const LabelType* labels,
-                                   size_t n_samples,
+                                   std::size_t n_samples,
                                    int n_neighbors,
                                    int n_outputs,
                                    int output_offset)
@@ -174,8 +174,8 @@ template <int TPB_X = 32, bool precomp_lbls = false>
 void class_probs(std::vector<float*>& out,
                  const int64_t* knn_indices,
                  std::vector<int*>& y,
-                 size_t n_index_rows,
-                 size_t n_query_rows,
+                 std::size_t n_index_rows,
+                 std::size_t n_query_rows,
                  int k,
                  std::vector<int*>& uniq_labels,
                  std::vector<int>& n_unique,
@@ -184,7 +184,7 @@ void class_probs(std::vector<float*>& out,
                  cudaStream_t* int_streams = nullptr,
                  int n_int_streams         = 0)
 {
-  for (int i = 0; i < y.size(); i++) {
+  for (std::size_t i = 0; i < y.size(); i++) {
     cudaStream_t stream = raft::select_stream(user_stream, int_streams, n_int_streams, i);
 
     int n_unique_labels = n_unique[i];
@@ -192,7 +192,7 @@ void class_probs(std::vector<float*>& out,
 
     CUDA_CHECK(cudaMemsetAsync(out[i], 0, cur_size * sizeof(float), stream));
 
-    dim3 grid(raft::ceildiv(n_query_rows, (size_t)TPB_X), 1, 1);
+    dim3 grid(raft::ceildiv(n_query_rows, static_cast<std::size_t>(TPB_X)), 1, 1);
     dim3 blk(TPB_X, 1, 1);
 
     /**
@@ -254,8 +254,8 @@ template <int TPB_X = 32, bool precomp_lbls = false>
 void knn_classify(int* out,
                   const int64_t* knn_indices,
                   std::vector<int*>& y,
-                  size_t n_index_rows,
-                  size_t n_query_rows,
+                  std::size_t n_index_rows,
+                  std::size_t n_query_rows,
                   int k,
                   std::vector<int*>& uniq_labels,
                   std::vector<int>& n_unique,
@@ -268,7 +268,7 @@ void knn_classify(int* out,
   std::vector<device_buffer<float>*> tmp_probs;
 
   // allocate temporary memory
-  for (int i = 0; i < n_unique.size(); i++) {
+  for (std::size_t i = 0; i < n_unique.size(); i++) {
     int size = n_unique[i];
 
     cudaStream_t stream = raft::select_stream(user_stream, int_streams, n_int_streams, i);
@@ -299,10 +299,10 @@ void knn_classify(int* out,
                                 int_streams,
                                 n_int_streams);
 
-  dim3 grid(raft::ceildiv(n_query_rows, (size_t)TPB_X), 1, 1);
+  dim3 grid(raft::ceildiv(n_query_rows, static_cast<std::size_t>(TPB_X)), 1, 1);
   dim3 blk(TPB_X, 1, 1);
 
-  for (int i = 0; i < y.size(); i++) {
+  for (std::size_t i = 0; i < y.size(); i++) {
     cudaStream_t stream = raft::select_stream(user_stream, int_streams, n_int_streams, i);
 
     int n_unique_labels = n_unique[i];
@@ -359,11 +359,11 @@ void knn_regress(ValType* out,
   /**
    * Vote average regression value
    */
-  for (int i = 0; i < y.size(); i++) {
+  for (std::size_t i = 0; i < y.size(); i++) {
     cudaStream_t stream = raft::select_stream(user_stream, int_streams, n_int_streams, i);
 
     regress_avg_kernel<ValType, precomp_lbls>
-      <<<raft::ceildiv(n_query_rows, (size_t)TPB_X), TPB_X, 0, stream>>>(
+      <<<raft::ceildiv(n_query_rows, static_cast<std::size_t>(TPB_X)), TPB_X, 0, stream>>>(
         out, knn_indices, y[i], n_query_rows, k, y.size(), i);
 
     CUDA_CHECK(cudaStreamSynchronize(stream));
diff --git a/cpp/src_prims/sparse/batched/csr.cuh b/cpp/src_prims/sparse/batched/csr.cuh
index 798d45fd5d..202a77e9f8 100644
--- a/cpp/src_prims/sparse/batched/csr.cuh
+++ b/cpp/src_prims/sparse/batched/csr.cuh
@@ -26,22 +26,24 @@
 
 #pragma once
 
-#include <algorithm>
-#include <memory>
-#include <vector>
-
-#include <thrust/execution_policy.h>
-#include <thrust/for_each.h>
-#include <thrust/iterator/counting_iterator.h>
-
+#include <cuml/common/device_buffer.hpp>
 #include <cuml/common/utils.hpp>
-#include <raft/mr/device/allocator.hpp>
+
+#include <linalg/batched/matrix.cuh>
 
 #include <raft/cudart_utils.h>
 #include <raft/linalg/cusolver_wrappers.h>
-#include <cuml/common/device_buffer.hpp>
-#include <linalg/batched/matrix.cuh>
 #include <raft/matrix/matrix.cuh>
+#include <raft/mr/device/allocator.hpp>
+
+#include <thrust/execution_policy.h>
+#include <thrust/for_each.h>
+#include <thrust/iterator/counting_iterator.h>
+
+#include <algorithm>
+#include <cstddef>
+#include <memory>
+#include <vector>
 
 namespace MLCommon {
 namespace Sparse {
@@ -137,6 +139,8 @@ static __global__ void csr_to_dense_kernel(T* dense,
 template <typename T>
 class CSR {
  public:
+  using shape_type = std::pair<std::size_t, std::size_t>;
+
   /**
    * @brief Constructor that leaves the matrix uninitialized
    *
@@ -149,10 +153,10 @@ class CSR {
    * @param[in] allocator        Device memory allocator
    * @param[in] stream           CUDA stream
    */
-  CSR(int m,
-      int n,
-      int nnz,
-      int batch_size,
+  CSR(std::size_t m,
+      std::size_t n,
+      std::size_t nnz,
+      std::size_t batch_size,
       cublasHandle_t cublasHandle,
       cusolverSpHandle_t cusolverSpHandle,
       std::shared_ptr<raft::mr::device::allocator> allocator,
@@ -188,10 +192,10 @@ class CSR {
    * @param[in] allocator        Device memory allocator
    * @param[in] stream           CUDA stream
    */
-  CSR(int m,
-      int n,
-      int nnz,
-      int batch_size,
+  CSR(std::size_t m,
+      std::size_t n,
+      std::size_t nnz,
+      std::size_t batch_size,
       cublasHandle_t cublasHandle,
       cusolverSpHandle_t cusolverSpHandle,
       T* d_values,
@@ -285,15 +289,15 @@ class CSR {
                            int* d_col_index = nullptr,
                            int* d_row_index = nullptr)
   {
-    std::pair<int, int> shape = dense.shape();
+    auto shape = dense.shape();
 
     // Create the index arrays from the mask
     std::vector<int> h_col_index;
     std::vector<int> h_row_index = std::vector<int>(shape.first + 1);
     int nnz                      = 0;
-    for (int i = 0; i < shape.first; i++) {
+    for (std::size_t i = 0; i < shape.first; i++) {
       h_row_index[i] = nnz;
-      for (int j = 0; j < shape.second; j++) {
+      for (std::size_t j = 0; j < shape.second; j++) {
         if (mask[j * shape.first + i]) {
           h_col_index.push_back(j);
           nnz++;
@@ -369,10 +373,10 @@ class CSR {
   }
 
   //! Return batch size
-  size_t batches() const { return m_batch_size; }
+  std::size_t batches() const { return m_batch_size; }
 
   //! Return number of non-zero elements
-  size_t nnz() const { return m_nnz; }
+  std::size_t nnz() const { return m_nnz; }
 
   //! Return cublas handle
   cublasHandle_t cublasHandle() const { return m_cublasHandle; }
@@ -387,7 +391,7 @@ class CSR {
   cudaStream_t stream() const { return m_stream; }
 
   //! Return shape
-  const std::pair<int, int>& shape() const { return m_shape; }
+  const shape_type& shape() const { return m_shape; }
 
   //! Return values array
   T* get_values() { return d_values; }
@@ -403,10 +407,10 @@ class CSR {
 
  protected:
   //! Shape (rows, cols) of matrices.
-  std::pair<int, int> m_shape;
+  shape_type m_shape;
 
   //! Number of non-zero values per matrix
-  int m_nnz;
+  std::size_t m_nnz;
 
   //! Array(pointer) to the values in all the batched matrices.
   device_buffer<T> m_values;
@@ -421,7 +425,7 @@ class CSR {
   int* d_row_index;
 
   //! Number of matrices in batch
-  size_t m_batch_size;
+  std::size_t m_batch_size;
 
   std::shared_ptr<raft::mr::device::allocator> m_allocator;
   cublasHandle_t m_cublasHandle;
@@ -496,8 +500,8 @@ void b_spmv(T alpha,
             T beta,
             LinAlg::Batched::Matrix<T>& y)
 {
-  int m = A.shape().first;
-  int n = A.shape().second;
+  auto m = A.shape().first;
+  auto n = A.shape().second;
   // A few checks
   ASSERT(std::min(x.shape().first, x.shape().second) == 1 &&
            std::max(x.shape().first, x.shape().second) == n,
@@ -673,11 +677,11 @@ void b_spmm(T alpha,
             LinAlg::Batched::Matrix<T>& C,
             bool use_shared_mem = true)
 {
-  int m   = A.shape().first;
-  int n   = B.shape().second;
-  int k   = A.shape().second;
-  int nb  = A.batches();
-  int nnz = A.nnz();
+  auto m   = A.shape().first;
+  auto n   = B.shape().second;
+  auto k   = A.shape().second;
+  auto nb  = A.batches();
+  auto nnz = A.nnz();
   // Check the parameters
   ASSERT(B.batches() == nb, "SpMM: A and B must have the same batch size");
   ASSERT(C.batches() == nb, "SpMM: A and C must have the same batch size");
diff --git a/cpp/src_prims/stats/minmax.cuh b/cpp/src_prims/stats/minmax.cuh
index f7d05ea280..d1d1376912 100644
--- a/cpp/src_prims/stats/minmax.cuh
+++ b/cpp/src_prims/stats/minmax.cuh
@@ -17,12 +17,27 @@
 #pragma once
 
 #include <raft/cudart_utils.h>
-#include <limits>
 #include <raft/cuda_utils.cuh>
 
+#include <limits>
+
 namespace MLCommon {
 namespace Stats {
 
+namespace detail {
+
+// TODO: replace with `std::bitcast` once we adopt C++20 or libcu++ adds it
+template <class To, class From>
+constexpr To bit_cast(const From& from) noexcept
+{
+  To to{};
+  static_assert(sizeof(To) == sizeof(From));
+  memcpy(&to, &from, sizeof(To));
+  return to;
+}
+
+}  // namespace detail
+
 template <typename T>
 struct encode_traits {
 };
@@ -39,26 +54,26 @@ struct encode_traits<double> {
 
 HDI int encode(float val)
 {
-  int i = *(int*)&val;
+  int i = detail::bit_cast<int>(val);
   return i >= 0 ? i : (1 << 31) | ~i;
 }
 
 HDI long long encode(double val)
 {
-  long long i = *(long long*)&val;
+  std::int64_t i = detail::bit_cast<std::int64_t>(val);
   return i >= 0 ? i : (1ULL << 63) | ~i;
 }
 
 HDI float decode(int val)
 {
   if (val < 0) val = (1 << 31) | ~val;
-  return *(float*)&val;
+  return detail::bit_cast<float>(val);
 }
 
 HDI double decode(long long val)
 {
   if (val < 0) val = (1ULL << 63) | ~val;
-  return *(double*)&val;
+  return detail::bit_cast<double>(val);
 }
 
 template <typename T, typename E>
diff --git a/cpp/test/prims/batched/csr.cu b/cpp/test/prims/batched/csr.cu
index af2be6eefd..e3e2e413e5 100644
--- a/cpp/test/prims/batched/csr.cu
+++ b/cpp/test/prims/batched/csr.cu
@@ -14,17 +14,20 @@
  * limitations under the License.
  */
 
-#include <gtest/gtest.h>
-#include <raft/mr/device/allocator.hpp>
-#include <random>
-#include <vector>
-
 #include <linalg_naive.h>
-#include <raft/cudart_utils.h>
 #include <test_utils.h>
+
 #include <linalg/batched/matrix.cuh>
 #include <sparse/batched/csr.cuh>
-#include "../test_utils.h"
+
+#include <raft/cudart_utils.h>
+#include <raft/mr/device/allocator.hpp>
+
+#include <gtest/gtest.h>
+
+#include <cstddef>
+#include <random>
+#include <vector>
 
 namespace MLCommon {
 namespace Sparse {
@@ -55,7 +58,9 @@ class CSRTest : public ::testing::TestWithParam<CSRInputs<T>> {
     params = ::testing::TestWithParam<CSRInputs<T>>::GetParam();
 
     // Check if the dimensions are valid and compute the output dimensions
-    int m_r, n_r;
+    int m_r{};
+    int n_r{};
+
     switch (params.operation) {
       case SpMV_op:
         ASSERT_TRUE(params.n == params.p);
@@ -97,10 +102,10 @@ class CSRTest : public ::testing::TestWithParam<CSRInputs<T>> {
     }
 
     // Generate random dense matrices/vectors
-    for (int i = 0; i < Bx.size(); i++)
+    for (std::size_t i = 0; i < Bx.size(); i++)
       Bx[i] = udis(gen);
     res_h.resize(params.batch_size * m_r * n_r);
-    for (int i = 0; i < res_h.size(); i++)
+    for (std::size_t i = 0; i < res_h.size(); i++)
       res_h[i] = udis(gen);
 
     // Create handles, stream, allocator
diff --git a/cpp/test/prims/batched/information_criterion.cu b/cpp/test/prims/batched/information_criterion.cu
index 6237135a3a..37e2331887 100644
--- a/cpp/test/prims/batched/information_criterion.cu
+++ b/cpp/test/prims/batched/information_criterion.cu
@@ -14,17 +14,19 @@
  * limitations under the License.
  */
 
+#include <test_utils.h>
+
+#include <metrics/batched/information_criterion.cuh>
+
+#include <raft/cudart_utils.h>
+#include <raft/mr/device/allocator.hpp>
+
 #include <gtest/gtest.h>
+
 #include <cmath>
-#include <raft/mr/device/allocator.hpp>
 #include <random>
 #include <vector>
 
-#include <raft/cudart_utils.h>
-#include <test_utils.h>
-#include <metrics/batched/information_criterion.cuh>
-#include "../test_utils.h"
-
 namespace MLCommon {
 namespace Metrics {
 namespace Batched {
@@ -33,7 +35,7 @@ template <typename T>
 void naive_ic(
   T* h_ic, const T* h_loglike, IC_Type ic_type, int n_params, int batch_size, int n_samples)
 {
-  T ic_base;
+  T ic_base{};
   T N = static_cast<T>(n_params);
   T M = static_cast<T>(n_samples);
   switch (ic_type) {
diff --git a/cpp/test/prims/batched/matrix.cu b/cpp/test/prims/batched/matrix.cu
index 8509e5bd8d..6657552165 100644
--- a/cpp/test/prims/batched/matrix.cu
+++ b/cpp/test/prims/batched/matrix.cu
@@ -14,21 +14,23 @@
  * limitations under the License.
  */
 
+#include <linalg_naive.h>
+#include <test_utils.h>
+
+#include <linalg/batched/matrix.cuh>
+
+#include <raft/cudart_utils.h>
+#include <raft/linalg/add.cuh>
+#include <raft/mr/device/allocator.hpp>
+
 #include <gtest/gtest.h>
+
 #include <algorithm>
 #include <cmath>
+#include <cstddef>
 #include <random>
 #include <vector>
 
-#include <raft/mr/device/allocator.hpp>
-
-#include <raft/cudart_utils.h>
-#include <test_utils.h>
-#include <linalg/batched/matrix.cuh>
-#include <raft/linalg/add.cuh>
-#include "../linalg_naive.h"
-#include "../test_utils.h"
-
 namespace MLCommon {
 namespace LinAlg {
 namespace Batched {
@@ -81,7 +83,9 @@ class MatrixTest : public ::testing::TestWithParam<MatrixInputs<T>> {
     int r      = params.operation == AZT_op ? params.n : params.m;
 
     // Check if the dimensions are valid and compute the output dimensions
-    int m_r, n_r;
+    int m_r{};
+    int n_r{};
+
     switch (params.operation) {
       case AB_op:
         ASSERT_TRUE(params.n == params.p);
@@ -149,11 +153,11 @@ class MatrixTest : public ::testing::TestWithParam<MatrixInputs<T>> {
     std::random_device rd;
     std::mt19937 gen(rd());
     std::uniform_real_distribution<T> udis(-1.0, 3.0);
-    for (int i = 0; i < A.size(); i++)
+    for (std::size_t i = 0; i < A.size(); i++)
       A[i] = udis(gen);
-    for (int i = 0; i < B.size(); i++)
+    for (std::size_t i = 0; i < B.size(); i++)
       B[i] = udis(gen);
-    for (int i = 0; i < Z.size(); i++)
+    for (std::size_t i = 0; i < Z.size(); i++)
       Z[i] = udis(gen);
 
     // Create handles, stream, allocator
diff --git a/cpp/test/prims/rand_index.cu b/cpp/test/prims/rand_index.cu
index 3948903faa..101e89e6c9 100644
--- a/cpp/test/prims/rand_index.cu
+++ b/cpp/test/prims/rand_index.cu
@@ -14,14 +14,18 @@
  * limitations under the License.
  */
 
-#include <gtest/gtest.h>
+#include "test_utils.h"
+
+#include <metrics/rand_index.cuh>
+
 #include <raft/cudart_utils.h>
+#include <raft/mr/device/allocator.hpp>
+
+#include <gtest/gtest.h>
+
 #include <algorithm>
 #include <iostream>
-#include <metrics/rand_index.cuh>
-#include <raft/mr/device/allocator.hpp>
 #include <random>
-#include "test_utils.h"
 
 namespace MLCommon {
 namespace Metrics {
@@ -59,9 +63,11 @@ class randIndexTest : public ::testing::TestWithParam<randIndexParam> {
     std::generate(arr2.begin(), arr2.end(), [&]() { return intGenerator(dre); });
 
     // generating the golden output
-    int64_t a_truth = 0, b_truth = 0, iter = 0, jiter;
-    for (; iter < size; ++iter) {
-      for (jiter = 0; jiter < iter; ++jiter) {
+    int64_t a_truth = 0;
+    int64_t b_truth = 0;
+
+    for (uint64_t iter = 0; iter < size; ++iter) {
+      for (uint64_t jiter = 0; jiter < iter; ++jiter) {
         if (arr1[iter] == arr1[jiter] && arr2[iter] == arr2[jiter]) {
           ++a_truth;
         } else if (arr1[iter] != arr1[jiter] && arr2[iter] != arr2[jiter]) {
diff --git a/cpp/test/sg/fil_test.cu b/cpp/test/sg/fil_test.cu
index 666b94ed55..c0fb4b15e5 100644
--- a/cpp/test/sg/fil_test.cu
+++ b/cpp/test/sg/fil_test.cu
@@ -14,25 +14,30 @@
  * limitations under the License.
  */
 
+#include "../../src/fil/internal.cuh"
+
+#include <test_utils.h>
+
 #include <cuml/fil/fil.h>
-#include <gtest/gtest.h>
+
 #include <raft/cudart_utils.h>
-#include <test_utils.h>
+#include <raft/cuda_utils.cuh>
+#include <raft/random/rng.cuh>
+
 #include <treelite/c_api.h>
 #include <treelite/frontend.h>
 #include <treelite/tree.h>
+
+#include <gtest/gtest.h>
+
 #include <cmath>
 #include <cstdio>
 #include <limits>
 #include <memory>
 #include <numeric>
 #include <ostream>
-#include <raft/cuda_utils.cuh>
-#include <raft/random/rng.cuh>
 #include <utility>
 
-#include "../../src/fil/internal.cuh"
-
 #define TL_CPP_CHECK(call) ASSERT(int(call) >= 0, "treelite call error")
 
 namespace ML {
@@ -206,7 +211,7 @@ class BaseFilTest : public testing::TestWithParam<FilTestParams> {
     CUDA_CHECK(cudaStreamSynchronize(stream));
 
     // mark leaves
-    for (size_t i = 0; i < ps.num_trees; ++i) {
+    for (int i = 0; i < ps.num_trees; ++i) {
       int num_tree_nodes = tree_num_nodes();
       size_t leaf_start  = num_tree_nodes * i + num_tree_nodes / 2;
       size_t leaf_end    = num_tree_nodes * (i + 1);
@@ -370,6 +375,8 @@ class BaseFilTest : public testing::TestWithParam<FilTestParams> {
             class_probabilities.begin();
         }
         break;
+      case fil::leaf_algo_t::GROVE_PER_CLASS_FEW_CLASSES:
+      case fil::leaf_algo_t::GROVE_PER_CLASS_MANY_CLASSES: break;
     }
 
     // copy to GPU
@@ -585,6 +592,8 @@ class TreeliteFilTest : public BaseFilTest {
           builder->SetLeafVectorNode(key, vec);
           break;
         }
+        case fil::leaf_algo_t::GROVE_PER_CLASS_FEW_CLASSES:
+        case fil::leaf_algo_t::GROVE_PER_CLASS_MANY_CLASSES: break;
       }
     } else {
       int left          = root + 2 * (node - root) + 1;
diff --git a/cpp/test/sg/multi_sum_test.cu b/cpp/test/sg/multi_sum_test.cu
index 2b4b77e984..ac7796565c 100644
--- a/cpp/test/sg/multi_sum_test.cu
+++ b/cpp/test/sg/multi_sum_test.cu
@@ -14,15 +14,21 @@
  * limitations under the License.
  */
 
-#include <gtest/gtest.h>
-#include <raft/cudart_utils.h>
 #include <test_utils.h>
-#include <thrust/device_vector.h>
-#include <thrust/host_vector.h>
+
 #include <cuml/fil/multi_sum.cuh>
+
+#include <raft/cudart_utils.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/random/rng.cuh>
 
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+
+#include <gtest/gtest.h>
+
+#include <cstddef>
+
 template <typename T>
 __device__ void serial_multi_sum(const T* in, T* out, int n_groups, int n_values)
 {
@@ -108,9 +114,9 @@ class MultiSumTest : public testing::TestWithParam<int> {
     CUDA_CHECK(cudaPeekAtLastError());
     error = error_d;
     CUDA_CHECK(cudaDeviceSynchronize());
-    for (int i = 0; i < params_h.size(); ++i) {
+    for (std::size_t i = 0; i < params_h.size(); ++i) {
       ASSERT(error[i] == 0,
-             "test # %d: block_dim_x %d multi_sum<%d>(on %d sets sized"
+             "test # %lu: block_dim_x %d multi_sum<%d>(on %d sets sized"
              " %d) gave wrong result",
              i,
              block_dim_x,
diff --git a/cpp/test/sg/rf_test.cu b/cpp/test/sg/rf_test.cu
index 6d3ba544ac..e055ed61d6 100644
--- a/cpp/test/sg/rf_test.cu
+++ b/cpp/test/sg/rf_test.cu
@@ -14,21 +14,29 @@
  * limitations under the License.
  */
 
-#include <cuml/tree/algo_helper.h>
+#include <test_utils.h>
+
 #include <decisiontree/quantile/quantile.h>
-#include <gtest/gtest.h>
+#include <decisiontree/batched-levelalgo/kernels.cuh>
+
+#include <cuml/tree/algo_helper.h>
+#include <cuml/datasets/make_blobs.hpp>
+#include <cuml/ensemble/randomforest.hpp>
+
+#include <random/make_blobs.cuh>
+
 #include <raft/cudart_utils.h>
 #include <raft/linalg/transpose.h>
-#include <test_utils.h>
+#include <raft/cuda_utils.cuh>
+#include <raft/handle.hpp>
+
 #include <thrust/device_vector.h>
 #include <thrust/host_vector.h>
-#include <cuml/datasets/make_blobs.hpp>
-#include <cuml/ensemble/randomforest.hpp>
-#include <decisiontree/batched-levelalgo/kernels.cuh>
+
+#include <gtest/gtest.h>
+
+#include <cstddef>
 #include <memory>
-#include <raft/cuda_utils.cuh>
-#include <raft/handle.hpp>
-#include <random/make_blobs.cuh>
 #include <random>
 #include <tuple>
 
@@ -81,7 +89,7 @@ std::vector<ParamT> SampleParameters(int num_samples, size_t seed, Args... args)
   std::default_random_engine gen(seed);
   AddParameters<0>(gen, tuple_sample, args...);
   std::vector<ParamT> sample(num_samples);
-  for (size_t i = 0; i < num_samples; i++) {
+  for (int i = 0; i < num_samples; i++) {
     sample[i] = make_struct<ParamT>(tuple_sample[i]);
   }
   return sample;
@@ -411,7 +419,7 @@ class RFQuantileBinsLowerBoundTest : public ::testing::TestWithParam<QuantileTes
                          nullptr);
     h_quantiles = quantiles;
     h_data      = data;
-    for (int i = 0; i < h_data.size(); ++i) {
+    for (std::size_t i = 0; i < h_data.size(); ++i) {
       auto d = h_data[i];
       // golden lower bound from thrust
       auto golden_lb = thrust::lower_bound(
diff --git a/cpp/test/sg/rf_treelite_test.cu b/cpp/test/sg/rf_treelite_test.cu
index f613a0fe03..ad1a156d02 100644
--- a/cpp/test/sg/rf_treelite_test.cu
+++ b/cpp/test/sg/rf_treelite_test.cu
@@ -14,23 +14,29 @@
  * limitations under the License.
  */
 
+#include <test_utils.h>
+
 #include <decisiontree/treelite_util.h>
-#include <gtest/gtest.h>
+#include <decisiontree/decisiontree.cuh>
+
+#include <cuml/ensemble/randomforest.hpp>
+
 #include <raft/cudart_utils.h>
 #include <raft/linalg/gemv.h>
 #include <raft/linalg/transpose.h>
-#include <sys/stat.h>
-#include <test_utils.h>
+#include <raft/cuda_utils.cuh>
+#include <raft/random/rng.cuh>
+
 #include <treelite/c_api.h>
 #include <treelite/c_api_runtime.h>
+
+#include <gtest/gtest.h>
+
+#include <sys/stat.h>
 #include <cstdlib>
-#include <cuml/ensemble/randomforest.hpp>
-#include <decisiontree/decisiontree.cuh>
 #include <fstream>
 #include <iostream>
 #include <limits>
-#include <raft/cuda_utils.cuh>
-#include <raft/random/rng.cuh>
 #include <string>
 
 namespace ML {
@@ -159,7 +165,7 @@ class RfTreeliteTestCommon : public ::testing::TestWithParam<RfInputs<T>> {
             params.n_inference_rows,
             params.n_cols,
             predicted_labels_d);
-    RF_metrics tmp = score(*handle, forest, labels_d, params.n_inference_rows, predicted_labels_d);
+    score(*handle, forest, labels_d, params.n_inference_rows, predicted_labels_d);
 
     CUDA_CHECK(cudaStreamSynchronize(stream));
 
diff --git a/cpp/test/sg/shap_kernel.cu b/cpp/test/sg/shap_kernel.cu
index d99c8a7996..212e737081 100644
--- a/cpp/test/sg/shap_kernel.cu
+++ b/cpp/test/sg/shap_kernel.cu
@@ -14,22 +14,19 @@
  * limitations under the License.
  */
 
-#include <gtest/gtest.h>
 #include <cuml/explainer/kernel_shap.hpp>
-#include <raft/mr/device/allocator.hpp>
+
+#include <test_utils.h>
 
 #include <raft/cudart_utils.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/handle.hpp>
+#include <raft/mr/device/allocator.hpp>
 
 #include <thrust/device_ptr.h>
 #include <thrust/device_vector.h>
 
-#include <test_utils.h>
-
-namespace MLCommon {
-class raft::mr::device::allocator;
-}
+#include <gtest/gtest.h>
 
 namespace ML {
 
diff --git a/cpp/test/sg/svc_test.cu b/cpp/test/sg/svc_test.cu
index 15d4c18019..8c3a069b94 100644
--- a/cpp/test/sg/svc_test.cu
+++ b/cpp/test/sg/svc_test.cu
@@ -1224,41 +1224,44 @@ TYPED_TEST(SmoSolverTest, MemoryLeak)
 
 TYPED_TEST(SmoSolverTest, DISABLED_MillionRows)
 {
-  // Stress test the kernel matrix calculation by calculating a kernel tile
-  // with more the 2.8B elemnts. This would fail with int32 adressing. The test
-  // is currently disabled because the memory usage might be prohibitive on CI
-  // The test will be enabled once https://github.com/rapidsai/cuml/pull/2449
-  // is merged, that PR would reduce the kernel tile memory size.
-  std::vector<std::pair<blobInput, TypeParam>> data{
-    {blobInput{1, 0.001, KernelParams{RBF, 3, 1, 0}, 2800000, 4}, 98},
-    {blobInput{1, 0.001, KernelParams{LINEAR, 3, 1, 0}, 2800000, 4}, 98},
-    {blobInput{1, 0.001, KernelParams{POLYNOMIAL, 3, 1, 0}, 2800000, 4}, 98},
-    {blobInput{1, 0.001, KernelParams{TANH, 3, 1, 0}, 2800000, 4}, 98}};
-  auto allocator = this->handle.get_device_allocator();
-
   if (sizeof(TypeParam) == 8) {
     GTEST_SKIP();  // Skip the test for double imput
-  }
-  for (auto d : data) {
-    auto p = d.first;
-    SCOPED_TRACE(p);
-    // explicit centers for the blobs
-    device_buffer<float> centers(allocator, this->stream, 2 * p.n_cols);
-    thrust::device_ptr<float> thrust_ptr(centers.data());
-    thrust::fill(thrust::cuda::par.on(this->stream), thrust_ptr, thrust_ptr + p.n_cols, -5.0f);
-    thrust::fill(
-      thrust::cuda::par.on(this->stream), thrust_ptr + p.n_cols, thrust_ptr + 2 * p.n_cols, +5.0f);
+  } else {
+    // Stress test the kernel matrix calculation by calculating a kernel tile
+    // with more the 2.8B elemnts. This would fail with int32 adressing. The test
+    // is currently disabled because the memory usage might be prohibitive on CI
+    // The test will be enabled once https://github.com/rapidsai/cuml/pull/2449
+    // is merged, that PR would reduce the kernel tile memory size.
+    std::vector<std::pair<blobInput, TypeParam>> data{
+      {blobInput{1, 0.001, KernelParams{RBF, 3, 1, 0}, 2800000, 4}, 98},
+      {blobInput{1, 0.001, KernelParams{LINEAR, 3, 1, 0}, 2800000, 4}, 98},
+      {blobInput{1, 0.001, KernelParams{POLYNOMIAL, 3, 1, 0}, 2800000, 4}, 98},
+      {blobInput{1, 0.001, KernelParams{TANH, 3, 1, 0}, 2800000, 4}, 98}};
+    auto allocator = this->handle.get_device_allocator();
 
-    device_buffer<TypeParam> x(allocator, this->stream, p.n_rows * p.n_cols);
-    device_buffer<TypeParam> y(allocator, this->stream, p.n_rows);
-    device_buffer<TypeParam> y_pred(allocator, this->stream, p.n_rows);
-    make_blobs(this->handle, x.data(), y.data(), p.n_rows, p.n_cols, 2, centers.data());
-    const int max_iter = 2;
-    SVC<TypeParam> svc(
-      this->handle, p.C, p.tol, p.kernel_params, 0, max_iter, 50, CUML_LEVEL_DEBUG);
-    svc.fit(x.data(), p.n_rows, p.n_cols, y.data());
-    // predict on the same dataset
-    svc.predict(x.data(), p.n_rows, p.n_cols, y_pred.data());
+    for (auto d : data) {
+      auto p = d.first;
+      SCOPED_TRACE(p);
+      // explicit centers for the blobs
+      device_buffer<float> centers(allocator, this->stream, 2 * p.n_cols);
+      thrust::device_ptr<float> thrust_ptr(centers.data());
+      thrust::fill(thrust::cuda::par.on(this->stream), thrust_ptr, thrust_ptr + p.n_cols, -5.0f);
+      thrust::fill(thrust::cuda::par.on(this->stream),
+                   thrust_ptr + p.n_cols,
+                   thrust_ptr + 2 * p.n_cols,
+                   +5.0f);
+
+      device_buffer<TypeParam> x(allocator, this->stream, p.n_rows * p.n_cols);
+      device_buffer<TypeParam> y(allocator, this->stream, p.n_rows);
+      device_buffer<TypeParam> y_pred(allocator, this->stream, p.n_rows);
+      make_blobs(this->handle, x.data(), y.data(), p.n_rows, p.n_cols, 2, centers.data());
+      const int max_iter = 2;
+      SVC<TypeParam> svc(
+        this->handle, p.C, p.tol, p.kernel_params, 0, max_iter, 50, CUML_LEVEL_DEBUG);
+      svc.fit(x.data(), p.n_rows, p.n_cols, y.data());
+      // predict on the same dataset
+      svc.predict(x.data(), p.n_rows, p.n_cols, y_pred.data());
+    }
   }
 }
 
diff --git a/cpp/test/sg/umap_parametrizable_test.cu b/cpp/test/sg/umap_parametrizable_test.cu
index f3baf5a1cf..066d66e469 100644
--- a/cpp/test/sg/umap_parametrizable_test.cu
+++ b/cpp/test/sg/umap_parametrizable_test.cu
@@ -14,27 +14,33 @@
  * limitations under the License.
  */
 
-#include <cuml/metrics/metrics.hpp>
+#include <test_utils.h>
 
-#include <gtest/gtest.h>
-#include <iostream>
-#include <vector>
+#include <umap/runner.cuh>
 
 #include <cuml/manifold/umapparams.h>
-#include <datasets/digits.h>
-#include <raft/cudart_utils.h>
-#include <test_utils.h>
 #include <cuml/common/device_buffer.hpp>
 #include <cuml/datasets/make_blobs.hpp>
 #include <cuml/manifold/umap.hpp>
+#include <cuml/metrics/metrics.hpp>
 #include <cuml/neighbors/knn.hpp>
+
+#include <datasets/digits.h>
 #include <linalg/reduce_rows_by_key.cuh>
+#include <selection/knn.cuh>
+
+#include <raft/cudart_utils.h>
 #include <raft/cuda_utils.cuh>
 #include <raft/distance/distance.cuh>
 #include <raft/handle.hpp>
 #include <raft/mr/device/allocator.hpp>
-#include <selection/knn.cuh>
-#include <umap/runner.cuh>
+
+#include <gtest/gtest.h>
+
+#include <cstddef>
+#include <iostream>
+#include <type_traits>
+#include <vector>
 
 using namespace ML;
 using namespace ML::Metrics;
@@ -47,10 +53,9 @@ using namespace MLCommon::Datasets::Digits;
 template <typename T>
 __global__ void has_nan_kernel(T* data, size_t len, bool* answer)
 {
-  int tid = threadIdx.x + blockIdx.x * blockDim.x;
-  if (tid >= len) return;
-  bool val = data[tid];
-  if (val != val) { *answer = true; }
+  static_assert(std::is_floating_point<T>());
+  std::size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
+  if ((tid < len) && isnan(data[tid])) { *answer = true; }
 }
 
 template <typename T>
@@ -128,10 +133,10 @@ class UMAPParametrizableTest : public ::testing::Test {
     int& n_samples      = test_params.n_samples;
     int& n_features     = test_params.n_features;
 
-    device_buffer<int64_t>* knn_indices_b;
-    device_buffer<float>* knn_dists_b;
-    int64_t* knn_indices = nullptr;
-    float* knn_dists     = nullptr;
+    device_buffer<int64_t>* knn_indices_b{};
+    device_buffer<float>* knn_dists_b{};
+    int64_t* knn_indices{};
+    float* knn_dists{};
     if (test_params.knn_params) {
       knn_indices_b =
         new device_buffer<int64_t>(alloc, stream, n_samples * umap_params.n_neighbors);
@@ -158,7 +163,8 @@ class UMAPParametrizableTest : public ::testing::Test {
     }
 
     float* model_embedding = nullptr;
-    device_buffer<float>* model_embedding_b;
+    device_buffer<float>* model_embedding_b{};
+
     if (test_params.fit_transform) {
       model_embedding = embedding_ptr;
     } else {