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 {