From cca81c4970233b8f8dc14f8ab7bc952cb703cad2 Mon Sep 17 00:00:00 2001
From: Ben Frederickson <github@benfrederickson.com>
Date: Mon, 13 Feb 2023 08:38:49 -0800
Subject: [PATCH] Add innerproduct to the pairwise distance api (#1226)

Authors:
  - Ben Frederickson (https://github.com/benfred)
  - Corey J. Nolet (https://github.com/cjnolet)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: https://github.com/rapidsai/raft/pull/1226
---
 cpp/CMakeLists.txt                            |   2 +
 cpp/bench/distance/distance_common.cuh        |   4 +-
 cpp/include/raft/distance/detail/distance.cuh | 202 ++++++++++++------
 .../distance/detail/kernels/gram_matrix.cuh   |   6 +-
 .../detail/kernels/kernel_matrices.cuh        |   6 +-
 cpp/include/raft/distance/distance.cuh        |  79 +++----
 .../specializations/detail/canberra.cuh       |   5 +-
 .../specializations/detail/chebyshev.cuh      |   4 +-
 .../specializations/detail/correlation.cuh    |   5 +-
 .../specializations/detail/cosine.cuh         |   5 +-
 .../detail/hamming_unexpanded.cuh             |   5 +-
 .../detail/hellinger_expanded.cuh             |   4 +-
 .../specializations/detail/inner_product.cuh  |  52 +++++
 .../specializations/detail/jensen_shannon.cuh |  25 +--
 .../specializations/detail/kl_divergence.cuh  |   4 +-
 .../distance/specializations/detail/l1.cuh    |   4 +-
 .../specializations/detail/l2_expanded.cuh    |  25 +--
 .../detail/l2_sqrt_expanded.cuh               |   4 +-
 .../detail/l2_sqrt_unexpanded.cuh             |   4 +-
 .../specializations/detail/l2_unexpanded.cuh  |   4 +-
 .../specializations/detail/lp_unexpanded.cuh  |   4 +-
 .../specializations/detail/russel_rao.cuh     |   4 +-
 .../distance/specializations/distance.cuh     |   3 +-
 cpp/include/raft/linalg/detail/gemm.hpp       |  17 +-
 cpp/include/raft/linalg/gemm.cuh              |   8 +-
 .../canberra_double_double_double_int.cu      |  16 +-
 .../detail/canberra_float_float_float_int.cu  |   2 +-
 .../chebyshev_double_double_double_int.cu     |   2 +-
 .../detail/chebyshev_float_float_float_int.cu |   2 +-
 .../correlation_double_double_double_int.cu   |   2 +-
 .../correlation_float_float_float_int.cu      |   2 +-
 .../detail/cosine_double_double_double_int.cu |   2 +-
 .../detail/cosine_float_float_float_int.cu    |   3 +-
 ...ing_unexpanded_double_double_double_int.cu |   2 +-
 ...amming_unexpanded_float_float_float_int.cu |   2 +-
 ...inger_expanded_double_double_double_int.cu |   2 +-
 ...ellinger_expanded_float_float_float_int.cu |   3 +-
 .../inner_product_double_double_double_int.cu |  37 ++++
 .../inner_product_float_float_float_int.cu    |  36 ++++
 ...jensen_shannon_double_double_double_int.cu |   4 +-
 .../jensen_shannon_float_float_float_int.cu   |  17 +-
 .../kl_divergence_double_double_double_int.cu |   4 +-
 .../kl_divergence_float_float_float_int.cu    |   4 +-
 .../detail/l1_double_double_double_int.cu     |   4 +-
 .../detail/l1_float_float_float_int.cu        |   4 +-
 .../l2_expanded_double_double_double_int.cu   |   4 +-
 .../l2_expanded_float_float_float_int.cu      |   4 +-
 ..._sqrt_expanded_double_double_double_int.cu |   4 +-
 .../l2_sqrt_expanded_float_float_float_int.cu |   4 +-
 ...qrt_unexpanded_double_double_double_int.cu |   4 +-
 ...2_sqrt_unexpanded_float_float_float_int.cu |   4 +-
 .../l2_unexpanded_double_double_double_int.cu |   4 +-
 .../l2_unexpanded_float_float_float_int.cu    |   4 +-
 .../lp_unexpanded_double_double_double_int.cu |   4 +-
 .../lp_unexpanded_float_float_float_int.cu    |   4 +-
 .../russel_rao_double_double_double_int.cu    |   4 +-
 .../russel_rao_float_float_float_int.cu       |   4 +-
 cpp/test/CMakeLists.txt                       |   1 +
 cpp/test/distance/dist_adj.cu                 |   4 +-
 cpp/test/distance/dist_inner_product.cu       |  76 +++++++
 cpp/test/distance/distance_base.cuh           |  55 +++--
 .../pylibraft/distance/pairwise_distance.pyx  |   2 +-
 .../pylibraft/pylibraft/test/test_distance.py |  10 +-
 63 files changed, 555 insertions(+), 276 deletions(-)
 create mode 100644 cpp/include/raft/distance/specializations/detail/inner_product.cuh
 create mode 100644 cpp/src/distance/distance/specializations/detail/inner_product_double_double_double_int.cu
 create mode 100644 cpp/src/distance/distance/specializations/detail/inner_product_float_float_float_int.cu
 create mode 100644 cpp/test/distance/dist_inner_product.cu

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 6ec8dada43..350306f523 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -321,6 +321,8 @@ if(RAFT_COMPILE_DIST_LIBRARY)
     src/distance/distance/specializations/detail/hamming_unexpanded_float_float_float_int.cu
     src/distance/distance/specializations/detail/hellinger_expanded_float_float_float_int.cu
     src/distance/distance/specializations/detail/hellinger_expanded_double_double_double_int.cu
+    src/distance/distance/specializations/detail/inner_product_float_float_float_int.cu
+    src/distance/distance/specializations/detail/inner_product_double_double_double_int.cu
     src/distance/distance/specializations/detail/jensen_shannon_float_float_float_int.cu
     src/distance/distance/specializations/detail/jensen_shannon_double_double_double_int.cu
     src/distance/distance/specializations/detail/kernels/gram_matrix_base_double.cu
diff --git a/cpp/bench/distance/distance_common.cuh b/cpp/bench/distance/distance_common.cuh
index 7d51bfa608..906271bf5a 100644
--- a/cpp/bench/distance/distance_common.cuh
+++ b/cpp/bench/distance/distance_common.cuh
@@ -49,7 +49,8 @@ struct distance : public fixture {
   void run_benchmark(::benchmark::State& state) override
   {
     loop_on_state(state, [this]() {
-      raft::distance::distance<DType, T, T, T>(x.data(),
+      raft::distance::distance<DType, T, T, T>(handle,
+                                               x.data(),
                                                y.data(),
                                                out.data(),
                                                params.m,
@@ -57,7 +58,6 @@ struct distance : public fixture {
                                                params.k,
                                                (void*)workspace.data(),
                                                worksize,
-                                               stream,
                                                params.isRowMajor);
     });
   }
diff --git a/cpp/include/raft/distance/detail/distance.cuh b/cpp/include/raft/distance/detail/distance.cuh
index b459c73bee..f9a85ea627 100644
--- a/cpp/include/raft/distance/detail/distance.cuh
+++ b/cpp/include/raft/distance/detail/distance.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuda_runtime_api.h>
+#include <raft/core/resource/cuda_stream.hpp>
 #include <raft/distance/detail/canberra.cuh>
 #include <raft/distance/detail/chebyshev.cuh>
 #include <raft/distance/detail/correlation.cuh>
@@ -30,6 +31,7 @@
 #include <raft/distance/detail/minkowski.cuh>
 #include <raft/distance/detail/russell_rao.cuh>
 #include <raft/distance/distance_types.hpp>
+#include <raft/linalg/gemm.cuh>
 #include <raft/util/cuda_utils.cuh>
 #include <rmm/device_uvector.hpp>
 
@@ -92,7 +94,8 @@ template <raft::distance::DistanceType distanceType,
           typename FinalLambda,
           typename Index_>
 struct DistanceImpl {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -101,7 +104,6 @@ struct DistanceImpl {
            void* workspace,
            size_t worksize,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType metric_arg = 2.0f)
   {
@@ -119,7 +121,8 @@ struct DistanceImpl<raft::distance::DistanceType::L2Expanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -128,12 +131,22 @@ struct DistanceImpl<raft::distance::DistanceType::L2Expanded,
            void* workspace,
            size_t worksize,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::euclideanAlgo1<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, false, (AccType*)workspace, worksize, fin_op, stream, isRowMajor);
+      m,
+      n,
+      k,
+      x,
+      y,
+      dist,
+      false,
+      (AccType*)workspace,
+      worksize,
+      fin_op,
+      raft::resource::get_cuda_stream(handle),
+      isRowMajor);
   }
 };
 
@@ -148,7 +161,8 @@ struct DistanceImpl<raft::distance::DistanceType::L2SqrtExpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -157,12 +171,22 @@ struct DistanceImpl<raft::distance::DistanceType::L2SqrtExpanded,
            void* workspace,
            size_t worksize,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::euclideanAlgo1<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, true, (AccType*)workspace, worksize, fin_op, stream, isRowMajor);
+      m,
+      n,
+      k,
+      x,
+      y,
+      dist,
+      true,
+      (AccType*)workspace,
+      worksize,
+      fin_op,
+      raft::resource::get_cuda_stream(handle),
+      isRowMajor);
   }
 };
 
@@ -177,7 +201,8 @@ struct DistanceImpl<raft::distance::DistanceType::CosineExpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -186,12 +211,60 @@ struct DistanceImpl<raft::distance::DistanceType::CosineExpanded,
            void* workspace,
            size_t worksize,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::cosineAlgo1<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, (AccType*)workspace, worksize, fin_op, stream, isRowMajor);
+      m,
+      n,
+      k,
+      x,
+      y,
+      dist,
+      (AccType*)workspace,
+      worksize,
+      fin_op,
+      raft::resource::get_cuda_stream(handle),
+      isRowMajor);
+  }
+};
+
+template <typename InType,
+          typename AccType,
+          typename OutType,
+          typename FinalLambda,
+          typename Index_>
+struct DistanceImpl<raft::distance::DistanceType::InnerProduct,
+                    InType,
+                    AccType,
+                    OutType,
+                    FinalLambda,
+                    Index_> {
+  void run(raft::resources const& handle,
+           const InType* x,
+           const InType* y,
+           OutType* dist,
+           Index_ m,
+           Index_ n,
+           Index_ k,
+           void* workspace,
+           size_t worksize,
+           FinalLambda fin_op,
+           bool isRowMajor,
+           InType)
+  {
+    cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+    raft::linalg::gemm(handle,
+                       dist,
+                       const_cast<InType*>(x),
+                       const_cast<InType*>(y),
+                       m,
+                       n,
+                       k,
+                       !isRowMajor,
+                       !isRowMajor,
+                       isRowMajor,
+                       stream);
   }
 };
 
@@ -206,7 +279,8 @@ struct DistanceImpl<raft::distance::DistanceType::L2Unexpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -215,12 +289,11 @@ struct DistanceImpl<raft::distance::DistanceType::L2Unexpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::euclideanAlgo2<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, false, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, false, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -235,7 +308,8 @@ struct DistanceImpl<raft::distance::DistanceType::L2SqrtUnexpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -244,12 +318,11 @@ struct DistanceImpl<raft::distance::DistanceType::L2SqrtUnexpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::euclideanAlgo2<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, true, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, true, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -264,7 +337,8 @@ struct DistanceImpl<raft::distance::DistanceType::L1,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -273,12 +347,11 @@ struct DistanceImpl<raft::distance::DistanceType::L1,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::l1Impl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -293,7 +366,8 @@ struct DistanceImpl<raft::distance::DistanceType::Linf,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -302,12 +376,11 @@ struct DistanceImpl<raft::distance::DistanceType::Linf,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::chebyshevImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -322,7 +395,8 @@ struct DistanceImpl<raft::distance::DistanceType::HellingerExpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -331,12 +405,11 @@ struct DistanceImpl<raft::distance::DistanceType::HellingerExpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::hellingerImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -351,7 +424,8 @@ struct DistanceImpl<raft::distance::DistanceType::LpUnexpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -360,12 +434,11 @@ struct DistanceImpl<raft::distance::DistanceType::LpUnexpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType metric_arg)
   {
     raft::distance::detail::minkowskiImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor, metric_arg);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor, metric_arg);
   }
 };
 
@@ -380,7 +453,8 @@ struct DistanceImpl<raft::distance::DistanceType::Canberra,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -389,12 +463,11 @@ struct DistanceImpl<raft::distance::DistanceType::Canberra,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::canberraImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -409,7 +482,8 @@ struct DistanceImpl<raft::distance::DistanceType::HammingUnexpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -418,12 +492,11 @@ struct DistanceImpl<raft::distance::DistanceType::HammingUnexpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::hammingUnexpandedImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -438,7 +511,8 @@ struct DistanceImpl<raft::distance::DistanceType::JensenShannon,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -447,12 +521,11 @@ struct DistanceImpl<raft::distance::DistanceType::JensenShannon,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::jensenShannonImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -467,7 +540,8 @@ struct DistanceImpl<raft::distance::DistanceType::RusselRaoExpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -476,12 +550,11 @@ struct DistanceImpl<raft::distance::DistanceType::RusselRaoExpanded,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::russellRaoImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -496,7 +569,8 @@ struct DistanceImpl<raft::distance::DistanceType::KLDivergence,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -505,12 +579,11 @@ struct DistanceImpl<raft::distance::DistanceType::KLDivergence,
            void*,
            size_t,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::klDivergenceImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, fin_op, stream, isRowMajor);
+      m, n, k, x, y, dist, fin_op, raft::resource::get_cuda_stream(handle), isRowMajor);
   }
 };
 
@@ -525,7 +598,8 @@ struct DistanceImpl<raft::distance::DistanceType::CorrelationExpanded,
                     OutType,
                     FinalLambda,
                     Index_> {
-  void run(const InType* x,
+  void run(raft::resources const& handle,
+           const InType* x,
            const InType* y,
            OutType* dist,
            Index_ m,
@@ -534,12 +608,21 @@ struct DistanceImpl<raft::distance::DistanceType::CorrelationExpanded,
            void* workspace,
            size_t worksize,
            FinalLambda fin_op,
-           cudaStream_t stream,
            bool isRowMajor,
            InType)
   {
     raft::distance::detail::correlationImpl<InType, AccType, OutType, FinalLambda, Index_>(
-      m, n, k, x, y, dist, (AccType*)workspace, worksize, fin_op, stream, isRowMajor);
+      m,
+      n,
+      k,
+      x,
+      y,
+      dist,
+      (AccType*)workspace,
+      worksize,
+      fin_op,
+      raft::resource::get_cuda_stream(handle),
+      isRowMajor);
   }
 };
 
@@ -562,7 +645,6 @@ struct DistanceImpl<raft::distance::DistanceType::CorrelationExpanded,
  * @param workspace temporary workspace needed for computations
  * @param worksize number of bytes of the workspace
  * @param fin_op the final gemm epilogue lambda
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  *
  * @note fin_op: This is a device lambda which is supposed to operate upon the
@@ -576,7 +658,8 @@ template <raft::distance::DistanceType distanceType,
           typename OutType,
           typename FinalLambda,
           typename Index_ = int>
-void distance(const InType* x,
+void distance(raft::resources const& handle,
+              const InType* x,
               const InType* y,
               OutType* dist,
               Index_ m,
@@ -585,12 +668,11 @@ void distance(const InType* x,
               void* workspace,
               size_t worksize,
               FinalLambda fin_op,
-              cudaStream_t stream,
               bool isRowMajor   = true,
               InType metric_arg = 2.0f)
 {
   DistanceImpl<distanceType, InType, AccType, OutType, FinalLambda, Index_> distImpl;
-  distImpl.run(x, y, dist, m, n, k, workspace, worksize, fin_op, stream, isRowMajor, metric_arg);
+  distImpl.run(handle, x, y, dist, m, n, k, workspace, worksize, fin_op, isRowMajor, metric_arg);
   RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
@@ -609,7 +691,6 @@ void distance(const InType* x,
  * @param k dimensionality
  * @param workspace temporary workspace needed for computations
  * @param worksize number of bytes of the workspace
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  *
  * @note if workspace is passed as nullptr, this will return in
@@ -633,7 +714,8 @@ template <raft::distance::DistanceType distanceType,
           typename AccType,
           typename OutType,
           typename Index_ = int>
-void distance(const InType* x,
+void distance(raft::resources const& handle,
+              const InType* x,
               const InType* y,
               OutType* dist,
               Index_ m,
@@ -641,7 +723,6 @@ void distance(const InType* x,
               Index_ k,
               void* workspace,
               size_t worksize,
-              cudaStream_t stream,
               bool isRowMajor   = true,
               InType metric_arg = 2.0f)
 {
@@ -653,7 +734,7 @@ void distance(const InType* x,
                 "OutType can be uint8_t, float, double,"
                 "if sizeof(OutType) > 1 then sizeof(AccType) == sizeof(OutType).");
   distance<distanceType, InType, AccType, OutType, final_op_type, Index_>(
-    x, y, dist, m, n, k, workspace, worksize, fin_op, stream, isRowMajor, metric_arg);
+    handle, x, y, dist, m, n, k, workspace, worksize, fin_op, isRowMajor, metric_arg);
   RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
@@ -710,25 +791,24 @@ size_t getWorkspaceSize(const InType* x, const InType* y, Index_ m, Index_ n, In
  * @param workspace temporary workspace buffer which can get resized as per the
  * needed workspace size
  * @param metric distance metric
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  */
 template <typename Type, typename Index_, raft::distance::DistanceType DistType>
-void pairwise_distance_impl(const Type* x,
+void pairwise_distance_impl(raft::resources const& handle,
+                            const Type* x,
                             const Type* y,
                             Type* dist,
                             Index_ m,
                             Index_ n,
                             Index_ k,
                             rmm::device_uvector<char>& workspace,
-                            cudaStream_t stream,
                             bool isRowMajor,
                             Type metric_arg = 2.0f)
 {
   auto worksize = getWorkspaceSize<DistType, Type, Type, Type, Index_>(x, y, m, n, k);
-  workspace.resize(worksize, stream);
+  workspace.resize(worksize, raft::resource::get_cuda_stream(handle));
   distance<DistType, Type, Type, Type, Index_>(
-    x, y, dist, m, n, k, workspace.data(), worksize, stream, isRowMajor, metric_arg);
+    handle, x, y, dist, m, n, k, workspace.data(), worksize, isRowMajor, metric_arg);
 }
 /** @} */
 };  // namespace detail
diff --git a/cpp/include/raft/distance/detail/kernels/gram_matrix.cuh b/cpp/include/raft/distance/detail/kernels/gram_matrix.cuh
index 344dda693e..aaf3052892 100644
--- a/cpp/include/raft/distance/detail/kernels/gram_matrix.cuh
+++ b/cpp/include/raft/distance/detail/kernels/gram_matrix.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -212,7 +212,7 @@ class GramMatrixBase {
                         int ld_out)
   {
     raft::distance::distance<raft::distance::DistanceType::L2Unexpanded, math_t, math_t, math_t>(
-      x1, x2, out, n1, n2, n_cols, stream, is_row_major);
+      raft::device_resources(stream), x1, x2, out, n1, n2, n_cols, is_row_major);
   }
 };
-};  // end namespace raft::distance::kernels::detail
\ No newline at end of file
+};  // end namespace raft::distance::kernels::detail
diff --git a/cpp/include/raft/distance/detail/kernels/kernel_matrices.cuh b/cpp/include/raft/distance/detail/kernels/kernel_matrices.cuh
index b74de84d80..d1465efdb0 100644
--- a/cpp/include/raft/distance/detail/kernels/kernel_matrices.cuh
+++ b/cpp/include/raft/distance/detail/kernels/kernel_matrices.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -359,7 +359,8 @@ class RBFKernel : public GramMatrixBase<math_t> {
                              math_t,
                              math_t,
                              decltype(fin_op),
-                             index_t>(const_cast<math_t*>(x1),
+                             index_t>(device_resources(stream),
+                                      const_cast<math_t*>(x1),
                                       const_cast<math_t*>(x2),
                                       out,
                                       n1,
@@ -368,7 +369,6 @@ class RBFKernel : public GramMatrixBase<math_t> {
                                       NULL,
                                       0,
                                       fin_op,
-                                      stream,
                                       is_row_major);
   }
 };
diff --git a/cpp/include/raft/distance/distance.cuh b/cpp/include/raft/distance/distance.cuh
index 93a5ce7f1a..5fd5ebe98d 100644
--- a/cpp/include/raft/distance/distance.cuh
+++ b/cpp/include/raft/distance/distance.cuh
@@ -18,7 +18,8 @@
 
 #pragma once
 
-#include <raft/core/device_resources.hpp>
+#include <raft/core/resource/cuda_stream.hpp>
+#include <raft/core/resources.hpp>
 #include <raft/distance/detail/distance.cuh>
 #include <raft/distance/distance_types.hpp>
 #include <rmm/device_uvector.hpp>
@@ -50,7 +51,6 @@ namespace distance {
  * @param workspace temporary workspace needed for computations
  * @param worksize number of bytes of the workspace
  * @param fin_op the final gemm epilogue lambda
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  * @param metric_arg metric argument (used for Minkowski distance)
  *
@@ -65,7 +65,8 @@ template <raft::distance::DistanceType distanceType,
           typename OutType,
           typename FinalLambda,
           typename Index_ = int>
-void distance(const InType* x,
+void distance(raft::resources const& handle,
+              const InType* x,
               const InType* y,
               OutType* dist,
               Index_ m,
@@ -74,12 +75,11 @@ void distance(const InType* x,
               void* workspace,
               size_t worksize,
               FinalLambda fin_op,
-              cudaStream_t stream,
               bool isRowMajor   = true,
               InType metric_arg = 2.0f)
 {
   detail::distance<distanceType, InType, AccType, OutType, FinalLambda, Index_>(
-    x, y, dist, m, n, k, workspace, worksize, fin_op, stream, isRowMajor, metric_arg);
+    handle, x, y, dist, m, n, k, workspace, worksize, fin_op, isRowMajor, metric_arg);
 }
 
 /**
@@ -97,7 +97,6 @@ void distance(const InType* x,
  * @param k dimensionality
  * @param workspace temporary workspace needed for computations
  * @param worksize number of bytes of the workspace
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  * @param metric_arg metric argument (used for Minkowski distance)
  *
@@ -109,7 +108,8 @@ template <raft::distance::DistanceType distanceType,
           typename AccType,
           typename OutType,
           typename Index_ = int>
-void distance(const InType* x,
+void distance(raft::resources const& handle,
+              const InType* x,
               const InType* y,
               OutType* dist,
               Index_ m,
@@ -117,12 +117,11 @@ void distance(const InType* x,
               Index_ k,
               void* workspace,
               size_t worksize,
-              cudaStream_t stream,
               bool isRowMajor   = true,
               InType metric_arg = 2.0f)
 {
   detail::distance<distanceType, InType, AccType, OutType, Index_>(
-    x, y, dist, m, n, k, workspace, worksize, stream, isRowMajor, metric_arg);
+    handle, x, y, dist, m, n, k, workspace, worksize, isRowMajor, metric_arg);
 }
 
 /**
@@ -193,7 +192,6 @@ size_t getWorkspaceSize(const raft::device_matrix_view<InType, layout> x,
  * @param m number of points in x
  * @param n number of points in y
  * @param k dimensionality
- * @param stream cuda stream
  * @param isRowMajor whether the matrices are row-major or col-major
  * @param metric_arg metric argument (used for Minkowski distance)
  */
@@ -202,21 +200,22 @@ template <raft::distance::DistanceType distanceType,
           typename AccType,
           typename OutType,
           typename Index_ = int>
-void distance(const InType* x,
+void distance(raft::resources const& handle,
+              const InType* x,
               const InType* y,
               OutType* dist,
               Index_ m,
               Index_ n,
               Index_ k,
-              cudaStream_t stream,
               bool isRowMajor   = true,
               InType metric_arg = 2.0f)
 {
+  auto stream = raft::resource::get_cuda_stream(handle);
   rmm::device_uvector<char> workspace(0, stream);
   auto worksize = getWorkspaceSize<distanceType, InType, AccType, OutType, Index_>(x, y, m, n, k);
   workspace.resize(worksize, stream);
   detail::distance<distanceType, InType, AccType, OutType, Index_>(
-    x, y, dist, m, n, k, workspace.data(), worksize, stream, isRowMajor, metric_arg);
+    handle, x, y, dist, m, n, k, workspace.data(), worksize, isRowMajor, metric_arg);
 }
 
 /**
@@ -238,7 +237,7 @@ void distance(const InType* x,
  * @param metric_arg metric argument (used for Minkowski distance)
  */
 template <typename Type, typename Index_ = int>
-void pairwise_distance(raft::device_resources const& handle,
+void pairwise_distance(raft::resources const& handle,
                        const Type* x,
                        const Type* y,
                        Type* dist,
@@ -253,64 +252,68 @@ void pairwise_distance(raft::device_resources const& handle,
   switch (metric) {
     case raft::distance::DistanceType::L2Expanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L2Expanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::L2SqrtExpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L2SqrtExpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::CosineExpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::CosineExpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::L1:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L1>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::L2Unexpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L2Unexpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::L2SqrtUnexpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::L2SqrtUnexpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::Linf:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::Linf>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::HellingerExpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::HellingerExpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::LpUnexpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::LpUnexpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor, metric_arg);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor, metric_arg);
       break;
     case raft::distance::DistanceType::Canberra:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::Canberra>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::HammingUnexpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::HammingUnexpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::JensenShannon:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::JensenShannon>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::RusselRaoExpanded:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::RusselRaoExpanded>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::KLDivergence:
       detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::KLDivergence>(
-        x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     case raft::distance::DistanceType::CorrelationExpanded:
       detail::
         pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::CorrelationExpanded>(
-          x, y, dist, m, n, k, workspace, handle.get_stream(), isRowMajor);
+          handle, x, y, dist, m, n, k, workspace, isRowMajor);
+      break;
+    case raft::distance::DistanceType::InnerProduct:
+      detail::pairwise_distance_impl<Type, Index_, raft::distance::DistanceType::InnerProduct>(
+        handle, x, y, dist, m, n, k, workspace, isRowMajor);
       break;
     default: THROW("Unknown or unsupported distance metric '%d'!", (int)metric);
   };
@@ -333,7 +336,7 @@ void pairwise_distance(raft::device_resources const& handle,
  * @param metric_arg metric argument (used for Minkowski distance)
  */
 template <typename Type, typename Index_ = int>
-void pairwise_distance(raft::device_resources const& handle,
+void pairwise_distance(raft::resources const& handle,
                        const Type* x,
                        const Type* y,
                        Type* dist,
@@ -344,7 +347,8 @@ void pairwise_distance(raft::device_resources const& handle,
                        bool isRowMajor = true,
                        Type metric_arg = 2.0f)
 {
-  rmm::device_uvector<char> workspace(0, handle.get_stream());
+  auto stream = raft::resource::get_cuda_stream(handle);
+  rmm::device_uvector<char> workspace(0, stream);
   pairwise_distance<Type, Index_>(
     handle, x, y, dist, m, n, k, workspace, metric, isRowMajor, metric_arg);
 }
@@ -398,7 +402,7 @@ template <raft::distance::DistanceType distanceType,
           typename OutType,
           typename layout = raft::layout_c_contiguous,
           typename Index_ = int>
-void distance(raft::device_resources const& handle,
+void distance(raft::resources const& handle,
               raft::device_matrix_view<InType, Index_, layout> const x,
               raft::device_matrix_view<InType, Index_, layout> const y,
               raft::device_matrix_view<OutType, Index_, layout> dist,
@@ -417,13 +421,13 @@ void distance(raft::device_resources const& handle,
 
   constexpr auto is_rowmajor = std::is_same_v<layout, layout_c_contiguous>;
 
-  distance<distanceType, InType, AccType, OutType, Index_>(x.data_handle(),
+  distance<distanceType, InType, AccType, OutType, Index_>(handle,
+                                                           x.data_handle(),
                                                            y.data_handle(),
                                                            dist.data_handle(),
                                                            x.extent(0),
                                                            y.extent(0),
                                                            x.extent(1),
-                                                           handle.get_stream(),
                                                            is_rowmajor,
                                                            metric_arg);
 }
@@ -441,7 +445,7 @@ void distance(raft::device_resources const& handle,
  * @param metric_arg metric argument (used for Minkowski distance)
  */
 template <typename Type, typename layout = layout_c_contiguous, typename Index_ = int>
-void pairwise_distance(raft::device_resources const& handle,
+void pairwise_distance(raft::resources const& handle,
                        device_matrix_view<Type, Index_, layout> const x,
                        device_matrix_view<Type, Index_, layout> const y,
                        device_matrix_view<Type, Index_, layout> dist,
@@ -462,7 +466,8 @@ void pairwise_distance(raft::device_resources const& handle,
 
   constexpr auto rowmajor = std::is_same_v<layout, layout_c_contiguous>;
 
-  rmm::device_uvector<char> workspace(0, handle.get_stream());
+  auto stream = raft::resource::get_cuda_stream(handle);
+  rmm::device_uvector<char> workspace(0, stream);
 
   pairwise_distance(handle,
                     x.data_handle(),
@@ -481,4 +486,4 @@ void pairwise_distance(raft::device_resources const& handle,
 };  // namespace distance
 };  // namespace raft
 
-#endif
\ No newline at end of file
+#endif
diff --git a/cpp/include/raft/distance/specializations/detail/canberra.cuh b/cpp/include/raft/distance/specializations/detail/canberra.cuh
index 6b6364fb58..badce715a5 100644
--- a/cpp/include/raft/distance/specializations/detail/canberra.cuh
+++ b/cpp/include/raft/distance/specializations/detail/canberra.cuh
@@ -23,6 +23,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::Canberra, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,11 +32,11 @@ extern template void distance<raft::distance::DistanceType::Canberra, float, flo
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void distance<raft::distance::DistanceType::Canberra, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -44,10 +45,8 @@ extern template void distance<raft::distance::DistanceType::Canberra, double, do
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/include/raft/distance/specializations/detail/chebyshev.cuh b/cpp/include/raft/distance/specializations/detail/chebyshev.cuh
index bbf5635ceb..9a46d7b488 100644
--- a/cpp/include/raft/distance/specializations/detail/chebyshev.cuh
+++ b/cpp/include/raft/distance/specializations/detail/chebyshev.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::Linf, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,11 +31,11 @@ extern template void distance<raft::distance::DistanceType::Linf, float, float,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void distance<raft::distance::DistanceType::Linf, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -43,7 +44,6 @@ extern template void distance<raft::distance::DistanceType::Linf, double, double
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/correlation.cuh b/cpp/include/raft/distance/specializations/detail/correlation.cuh
index 3826fb1030..013a0d43a3 100644
--- a/cpp/include/raft/distance/specializations/detail/correlation.cuh
+++ b/cpp/include/raft/distance/specializations/detail/correlation.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::CorrelationExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::CorrelationExpanded, float, float, float,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::CorrelationExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,10 +46,8 @@ distance<raft::distance::DistanceType::CorrelationExpanded, double, double, doub
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/include/raft/distance/specializations/detail/cosine.cuh b/cpp/include/raft/distance/specializations/detail/cosine.cuh
index 40e4a55aae..c88bd1b0f6 100644
--- a/cpp/include/raft/distance/specializations/detail/cosine.cuh
+++ b/cpp/include/raft/distance/specializations/detail/cosine.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::CosineExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::CosineExpanded, float, float, float, int>
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::CosineExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,10 +46,8 @@ distance<raft::distance::DistanceType::CosineExpanded, double, double, double, i
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/include/raft/distance/specializations/detail/hamming_unexpanded.cuh b/cpp/include/raft/distance/specializations/detail/hamming_unexpanded.cuh
index 08868eb7ea..3c5cad3315 100644
--- a/cpp/include/raft/distance/specializations/detail/hamming_unexpanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/hamming_unexpanded.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::HammingUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::HammingUnexpanded, float, float, float, i
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::HammingUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,10 +46,8 @@ distance<raft::distance::DistanceType::HammingUnexpanded, double, double, double
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/include/raft/distance/specializations/detail/hellinger_expanded.cuh b/cpp/include/raft/distance/specializations/detail/hellinger_expanded.cuh
index 9fe5c3d6f9..bf214c046f 100644
--- a/cpp/include/raft/distance/specializations/detail/hellinger_expanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/hellinger_expanded.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::HellingerExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::HellingerExpanded, float, float, float, i
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::HellingerExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,7 +46,6 @@ distance<raft::distance::DistanceType::HellingerExpanded, double, double, double
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/inner_product.cuh b/cpp/include/raft/distance/specializations/detail/inner_product.cuh
new file mode 100644
index 0000000000..d97d678928
--- /dev/null
+++ b/cpp/include/raft/distance/specializations/detail/inner_product.cuh
@@ -0,0 +1,52 @@
+/*
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <raft/distance/detail/distance.cuh>
+
+namespace raft {
+namespace distance {
+namespace detail {
+extern template void distance<raft::distance::DistanceType::InnerProduct, float, float, float, int>(
+  raft::resources const& handle,
+  const float* x,
+  const float* y,
+  float* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  size_t worksize,
+  bool isRowMajor,
+  float metric_arg);
+
+extern template void
+distance<raft::distance::DistanceType::InnerProduct, double, double, double, int>(
+  raft::resources const& handle,
+  const double* x,
+  const double* y,
+  double* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  size_t worksize,
+  bool isRowMajor,
+  double metric_arg);
+}  // namespace detail
+}  // namespace distance
+}  // namespace raft
diff --git a/cpp/include/raft/distance/specializations/detail/jensen_shannon.cuh b/cpp/include/raft/distance/specializations/detail/jensen_shannon.cuh
index 4e86417840..145834fb70 100644
--- a/cpp/include/raft/distance/specializations/detail/jensen_shannon.cuh
+++ b/cpp/include/raft/distance/specializations/detail/jensen_shannon.cuh
@@ -22,20 +22,22 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void
-distance<raft::distance::DistanceType::JensenShannon, float, float, float, int>(const float* x,
-                                                                                const float* y,
-                                                                                float* dist,
-                                                                                int m,
-                                                                                int n,
-                                                                                int k,
-                                                                                void* workspace,
-                                                                                size_t worksize,
-                                                                                cudaStream_t stream,
-                                                                                bool isRowMajor,
-                                                                                float metric_arg);
+distance<raft::distance::DistanceType::JensenShannon, float, float, float, int>(
+  raft::resources const& handle,
+  const float* x,
+  const float* y,
+  float* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  size_t worksize,
+  bool isRowMajor,
+  float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::JensenShannon, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -44,7 +46,6 @@ distance<raft::distance::DistanceType::JensenShannon, double, double, double, in
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/kl_divergence.cuh b/cpp/include/raft/distance/specializations/detail/kl_divergence.cuh
index 828b5b09e3..f0928916cd 100644
--- a/cpp/include/raft/distance/specializations/detail/kl_divergence.cuh
+++ b/cpp/include/raft/distance/specializations/detail/kl_divergence.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::KLDivergence, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,12 +31,12 @@ extern template void distance<raft::distance::DistanceType::KLDivergence, float,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::KLDivergence, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -44,7 +45,6 @@ distance<raft::distance::DistanceType::KLDivergence, double, double, double, int
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/l1.cuh b/cpp/include/raft/distance/specializations/detail/l1.cuh
index b4d9d4343b..23261a2571 100644
--- a/cpp/include/raft/distance/specializations/detail/l1.cuh
+++ b/cpp/include/raft/distance/specializations/detail/l1.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::L1, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,11 +31,11 @@ extern template void distance<raft::distance::DistanceType::L1, float, float, fl
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void distance<raft::distance::DistanceType::L1, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -43,7 +44,6 @@ extern template void distance<raft::distance::DistanceType::L1, double, double,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/l2_expanded.cuh b/cpp/include/raft/distance/specializations/detail/l2_expanded.cuh
index 8eadd6e233..f953018b7d 100644
--- a/cpp/include/raft/distance/specializations/detail/l2_expanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/l2_expanded.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::L2Expanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,22 +31,22 @@ extern template void distance<raft::distance::DistanceType::L2Expanded, float, f
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
-distance<raft::distance::DistanceType::L2Expanded, double, double, double, int>(const double* x,
-                                                                                const double* y,
-                                                                                double* dist,
-                                                                                int m,
-                                                                                int n,
-                                                                                int k,
-                                                                                void* workspace,
-                                                                                size_t worksize,
-                                                                                cudaStream_t stream,
-                                                                                bool isRowMajor,
-                                                                                double metric_arg);
+distance<raft::distance::DistanceType::L2Expanded, double, double, double, int>(
+  raft::resources const& handle,
+  const double* x,
+  const double* y,
+  double* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  size_t worksize,
+  bool isRowMajor,
+  double metric_arg);
 
 }  // namespace detail
 }  // namespace distance
diff --git a/cpp/include/raft/distance/specializations/detail/l2_sqrt_expanded.cuh b/cpp/include/raft/distance/specializations/detail/l2_sqrt_expanded.cuh
index 65e48e8401..9f5f6a3706 100644
--- a/cpp/include/raft/distance/specializations/detail/l2_sqrt_expanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/l2_sqrt_expanded.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::L2SqrtExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::L2SqrtExpanded, float, float, float, int>
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::L2SqrtExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,7 +46,6 @@ distance<raft::distance::DistanceType::L2SqrtExpanded, double, double, double, i
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/l2_sqrt_unexpanded.cuh b/cpp/include/raft/distance/specializations/detail/l2_sqrt_unexpanded.cuh
index 8d0de7711f..94531ddc33 100644
--- a/cpp/include/raft/distance/specializations/detail/l2_sqrt_unexpanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/l2_sqrt_unexpanded.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::L2SqrtUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::L2SqrtUnexpanded, float, float, float, in
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::L2SqrtUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,7 +46,6 @@ distance<raft::distance::DistanceType::L2SqrtUnexpanded, double, double, double,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/l2_unexpanded.cuh b/cpp/include/raft/distance/specializations/detail/l2_unexpanded.cuh
index 23c9849917..224b21fce8 100644
--- a/cpp/include/raft/distance/specializations/detail/l2_unexpanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/l2_unexpanded.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::L2Unexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,12 +31,12 @@ extern template void distance<raft::distance::DistanceType::L2Unexpanded, float,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::L2Unexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -44,7 +45,6 @@ distance<raft::distance::DistanceType::L2Unexpanded, double, double, double, int
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/lp_unexpanded.cuh b/cpp/include/raft/distance/specializations/detail/lp_unexpanded.cuh
index 593ac84ebb..e05ef02c42 100644
--- a/cpp/include/raft/distance/specializations/detail/lp_unexpanded.cuh
+++ b/cpp/include/raft/distance/specializations/detail/lp_unexpanded.cuh
@@ -22,6 +22,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 extern template void distance<raft::distance::DistanceType::LpUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -30,12 +31,12 @@ extern template void distance<raft::distance::DistanceType::LpUnexpanded, float,
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::LpUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -44,7 +45,6 @@ distance<raft::distance::DistanceType::LpUnexpanded, double, double, double, int
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/detail/russel_rao.cuh b/cpp/include/raft/distance/specializations/detail/russel_rao.cuh
index 0147559429..afc87997c0 100644
--- a/cpp/include/raft/distance/specializations/detail/russel_rao.cuh
+++ b/cpp/include/raft/distance/specializations/detail/russel_rao.cuh
@@ -23,6 +23,7 @@ namespace distance {
 namespace detail {
 extern template void
 distance<raft::distance::DistanceType::RusselRaoExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -31,12 +32,12 @@ distance<raft::distance::DistanceType::RusselRaoExpanded, float, float, float, i
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
 extern template void
 distance<raft::distance::DistanceType::RusselRaoExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -45,7 +46,6 @@ distance<raft::distance::DistanceType::RusselRaoExpanded, double, double, double
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/include/raft/distance/specializations/distance.cuh b/cpp/include/raft/distance/specializations/distance.cuh
index 053441d68a..a0c35ca9a8 100644
--- a/cpp/include/raft/distance/specializations/distance.cuh
+++ b/cpp/include/raft/distance/specializations/distance.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 #include <raft/distance/specializations/detail/cosine.cuh>
 #include <raft/distance/specializations/detail/hamming_unexpanded.cuh>
 #include <raft/distance/specializations/detail/hellinger_expanded.cuh>
+#include <raft/distance/specializations/detail/inner_product.cuh>
 #include <raft/distance/specializations/detail/jensen_shannon.cuh>
 #include <raft/distance/specializations/detail/kernels.cuh>
 #include <raft/distance/specializations/detail/kl_divergence.cuh>
diff --git a/cpp/include/raft/linalg/detail/gemm.hpp b/cpp/include/raft/linalg/detail/gemm.hpp
index ba9496c3b9..d82c821148 100644
--- a/cpp/include/raft/linalg/detail/gemm.hpp
+++ b/cpp/include/raft/linalg/detail/gemm.hpp
@@ -20,7 +20,8 @@
 
 #include "cublas_wrappers.hpp"
 
-#include <raft/core/device_resources.hpp>
+#include <raft/core/resource/cublas_handle.hpp>
+#include <raft/core/resources.hpp>
 
 namespace raft {
 namespace linalg {
@@ -49,7 +50,7 @@ namespace detail {
  * @param [in] stream
  */
 template <typename math_t, bool DevicePointerMode = false>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const bool trans_a,
           const bool trans_b,
           const int m,
@@ -65,7 +66,7 @@ void gemm(raft::device_resources const& handle,
           const int ldc,
           cudaStream_t stream)
 {
-  cublasHandle_t cublas_h = handle.get_cublas_handle();
+  auto cublas_h = raft::resource::get_cublas_handle(handle);
   cublas_device_pointer_mode<DevicePointerMode> pmode(cublas_h);
   RAFT_CUBLAS_TRY(cublasgemm(cublas_h,
                              trans_a ? CUBLAS_OP_T : CUBLAS_OP_N,
@@ -103,7 +104,7 @@ void gemm(raft::device_resources const& handle,
  * @param stream cuda stream
  */
 template <typename math_t>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const math_t* a,
           int n_rows_a,
           int n_cols_a,
@@ -117,7 +118,7 @@ void gemm(raft::device_resources const& handle,
           math_t beta,
           cudaStream_t stream)
 {
-  cublasHandle_t cublas_h = handle.get_cublas_handle();
+  auto cublas_h = raft::resource::get_cublas_handle(handle);
 
   int m   = n_rows_c;
   int n   = n_cols_c;
@@ -130,7 +131,7 @@ void gemm(raft::device_resources const& handle,
 }
 
 template <typename math_t>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const math_t* a,
           int n_rows_a,
           int n_cols_a,
@@ -149,7 +150,7 @@ void gemm(raft::device_resources const& handle,
 }
 
 template <typename T, bool DevicePointerMode = false>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           T* z,
           T* x,
           T* y,
@@ -163,7 +164,7 @@ void gemm(raft::device_resources const& handle,
           T* alpha,
           T* beta)
 {
-  cublasHandle_t cublas_h = handle.get_cublas_handle();
+  auto cublas_h = raft::resource::get_cublas_handle(handle);
   cublas_device_pointer_mode<DevicePointerMode> pmode(cublas_h);
 
   cublasOperation_t trans_a, trans_b;
diff --git a/cpp/include/raft/linalg/gemm.cuh b/cpp/include/raft/linalg/gemm.cuh
index d5dc5ffab5..a336f844bf 100644
--- a/cpp/include/raft/linalg/gemm.cuh
+++ b/cpp/include/raft/linalg/gemm.cuh
@@ -52,7 +52,7 @@ namespace linalg {
  * @param [in] stream
  */
 template <typename math_t, bool DevicePointerMode = false>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const bool trans_a,
           const bool trans_b,
           const int m,
@@ -91,7 +91,7 @@ void gemm(raft::device_resources const& handle,
  * @param stream cuda stream
  */
 template <typename math_t>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const math_t* a,
           int n_rows_a,
           int n_cols_a,
@@ -126,7 +126,7 @@ void gemm(raft::device_resources const& handle,
  * @param stream cuda stream
  */
 template <typename math_t>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           const math_t* a,
           int n_rows_a,
           int n_cols_a,
@@ -161,7 +161,7 @@ void gemm(raft::device_resources const& handle,
  * @param beta scalar
  */
 template <typename T>
-void gemm(raft::device_resources const& handle,
+void gemm(raft::resources const& handle,
           T* z,
           T* x,
           T* y,
diff --git a/cpp/src/distance/distance/specializations/detail/canberra_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/canberra_double_double_double_int.cu
index a00861f421..d3a629f5a0 100644
--- a/cpp/src/distance/distance/specializations/detail/canberra_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/canberra_double_double_double_int.cu
@@ -19,20 +19,8 @@
 namespace raft {
 namespace distance {
 namespace detail {
-template void distance<raft::distance::DistanceType::Canberra, float, float, float, int>(
-  const float* x,
-  const float* y,
-  float* dist,
-  int m,
-  int n,
-  int k,
-  void* workspace,
-  size_t worksize,
-  cudaStream_t stream,
-  bool isRowMajor,
-  float metric_arg);
-
 template void distance<raft::distance::DistanceType::Canberra, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -41,10 +29,8 @@ template void distance<raft::distance::DistanceType::Canberra, double, double, d
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/src/distance/distance/specializations/detail/canberra_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/canberra_float_float_float_int.cu
index b5ed1dd629..71899c00a7 100644
--- a/cpp/src/distance/distance/specializations/detail/canberra_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/canberra_float_float_float_int.cu
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::Canberra, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::Canberra, float, float, flo
   int k,
   void* workspace,
   size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/chebyshev_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/chebyshev_double_double_double_int.cu
index 14938769c3..c071ddc902 100644
--- a/cpp/src/distance/distance/specializations/detail/chebyshev_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/chebyshev_double_double_double_int.cu
@@ -21,6 +21,7 @@ namespace distance {
 namespace detail {
 
 template void distance<raft::distance::DistanceType::Linf, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -29,7 +30,6 @@ template void distance<raft::distance::DistanceType::Linf, double, double, doubl
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/chebyshev_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/chebyshev_float_float_float_int.cu
index 3a4e5a92fd..d42d769306 100644
--- a/cpp/src/distance/distance/specializations/detail/chebyshev_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/chebyshev_float_float_float_int.cu
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::Linf, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::Linf, float, float, float,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/correlation_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/correlation_double_double_double_int.cu
index 7f80786b7e..4c626a0c73 100644
--- a/cpp/src/distance/distance/specializations/detail/correlation_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/correlation_double_double_double_int.cu
@@ -22,6 +22,7 @@ namespace detail {
 
 template void
 distance<raft::distance::DistanceType::CorrelationExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -30,7 +31,6 @@ distance<raft::distance::DistanceType::CorrelationExpanded, double, double, doub
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/correlation_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/correlation_float_float_float_int.cu
index f1c83fecc1..5f23a5042e 100644
--- a/cpp/src/distance/distance/specializations/detail/correlation_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/correlation_float_float_float_int.cu
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::CorrelationExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::CorrelationExpanded, float,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/cosine_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/cosine_double_double_double_int.cu
index b10d22716b..b7f2eb569d 100644
--- a/cpp/src/distance/distance/specializations/detail/cosine_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/cosine_double_double_double_int.cu
@@ -21,6 +21,7 @@ namespace distance {
 namespace detail {
 
 template void distance<raft::distance::DistanceType::CosineExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -29,7 +30,6 @@ template void distance<raft::distance::DistanceType::CosineExpanded, double, dou
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/cosine_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/cosine_float_float_float_int.cu
index c597826926..650c889a0d 100644
--- a/cpp/src/distance/distance/specializations/detail/cosine_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/cosine_float_float_float_int.cu
@@ -19,7 +19,9 @@
 namespace raft {
 namespace distance {
 namespace detail {
+
 template void distance<raft::distance::DistanceType::CosineExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +30,6 @@ template void distance<raft::distance::DistanceType::CosineExpanded, float, floa
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_double_double_double_int.cu
index 7cd9d74667..282449f48b 100644
--- a/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_double_double_double_int.cu
@@ -22,6 +22,7 @@ namespace detail {
 
 template void
 distance<raft::distance::DistanceType::HammingUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -30,7 +31,6 @@ distance<raft::distance::DistanceType::HammingUnexpanded, double, double, double
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_float_float_float_int.cu
index f38aa01edd..2bed6c2401 100644
--- a/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/hamming_unexpanded_float_float_float_int.cu
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::HammingUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::HammingUnexpanded, float, f
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/hellinger_expanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/hellinger_expanded_double_double_double_int.cu
index a7a56ddc1e..b14425eec3 100644
--- a/cpp/src/distance/distance/specializations/detail/hellinger_expanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/hellinger_expanded_double_double_double_int.cu
@@ -22,6 +22,7 @@ namespace detail {
 
 template void
 distance<raft::distance::DistanceType::HellingerExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -30,7 +31,6 @@ distance<raft::distance::DistanceType::HellingerExpanded, double, double, double
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/hellinger_expanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/hellinger_expanded_float_float_float_int.cu
index 77403af132..b738b3983b 100644
--- a/cpp/src/distance/distance/specializations/detail/hellinger_expanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/hellinger_expanded_float_float_float_int.cu
@@ -19,7 +19,9 @@
 namespace raft {
 namespace distance {
 namespace detail {
+
 template void distance<raft::distance::DistanceType::HellingerExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +30,6 @@ template void distance<raft::distance::DistanceType::HellingerExpanded, float, f
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/inner_product_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/inner_product_double_double_double_int.cu
new file mode 100644
index 0000000000..41bd6c2c14
--- /dev/null
+++ b/cpp/src/distance/distance/specializations/detail/inner_product_double_double_double_int.cu
@@ -0,0 +1,37 @@
+/*
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <raft/distance/detail/distance.cuh>
+
+namespace raft {
+namespace distance {
+namespace detail {
+template void distance<raft::distance::DistanceType::InnerProduct, double, double, double, int>(
+  raft::resources const& handle,
+  const double* x,
+  const double* y,
+  double* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  std::size_t worksize,
+  bool isRowMajor,
+  double metric_arg);
+
+}  // namespace detail
+}  // namespace distance
+}  // namespace raft
diff --git a/cpp/src/distance/distance/specializations/detail/inner_product_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/inner_product_float_float_float_int.cu
new file mode 100644
index 0000000000..f147c51b52
--- /dev/null
+++ b/cpp/src/distance/distance/specializations/detail/inner_product_float_float_float_int.cu
@@ -0,0 +1,36 @@
+/*
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <raft/distance/detail/distance.cuh>
+
+namespace raft {
+namespace distance {
+namespace detail {
+template void distance<raft::distance::DistanceType::InnerProduct, float, float, float, int>(
+  raft::resources const& handle,
+  const float* x,
+  const float* y,
+  float* dist,
+  int m,
+  int n,
+  int k,
+  void* workspace,
+  std::size_t worksize,
+  bool isRowMajor,
+  float metric_arg);
+}  // namespace detail
+}  // namespace distance
+}  // namespace raft
diff --git a/cpp/src/distance/distance/specializations/detail/jensen_shannon_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/jensen_shannon_double_double_double_int.cu
index 30fbf70322..c490033dc3 100644
--- a/cpp/src/distance/distance/specializations/detail/jensen_shannon_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/jensen_shannon_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::JensenShannon, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::JensenShannon, double, doub
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/jensen_shannon_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/jensen_shannon_float_float_float_int.cu
index 31bc11aa71..ae1d860b3e 100644
--- a/cpp/src/distance/distance/specializations/detail/jensen_shannon_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/jensen_shannon_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::JensenShannon, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,23 +29,9 @@ template void distance<raft::distance::DistanceType::JensenShannon, float, float
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
-template void distance<raft::distance::DistanceType::JensenShannon, double, double, double, int>(
-  const double* x,
-  const double* y,
-  double* dist,
-  int m,
-  int n,
-  int k,
-  void* workspace,
-  std::size_t worksize,
-  cudaStream_t stream,
-  bool isRowMajor,
-  double metric_arg);
-
 }  // namespace detail
 }  // namespace distance
 }  // namespace raft
diff --git a/cpp/src/distance/distance/specializations/detail/kl_divergence_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/kl_divergence_double_double_double_int.cu
index 0ac2b23b29..ed2a29bccf 100644
--- a/cpp/src/distance/distance/specializations/detail/kl_divergence_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/kl_divergence_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::KLDivergence, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::KLDivergence, double, doubl
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/kl_divergence_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/kl_divergence_float_float_float_int.cu
index f56179fcac..511e0fb1a6 100644
--- a/cpp/src/distance/distance/specializations/detail/kl_divergence_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/kl_divergence_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::KLDivergence, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::KLDivergence, float, float,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l1_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/l1_double_double_double_int.cu
index b19603dcb9..7a1c884223 100644
--- a/cpp/src/distance/distance/specializations/detail/l1_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l1_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L1, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L1, double, double, double,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l1_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/l1_float_float_float_int.cu
index 8319239115..f307ade06f 100644
--- a/cpp/src/distance/distance/specializations/detail/l1_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l1_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L1, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L1, float, float, float, in
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_expanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/l2_expanded_double_double_double_int.cu
index c77dee4220..18ddf10b14 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_expanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_expanded_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2Expanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2Expanded, double, double,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 }  // namespace detail
diff --git a/cpp/src/distance/distance/specializations/detail/l2_expanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/l2_expanded_float_float_float_int.cu
index 778a65ad01..28e7b73243 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_expanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_expanded_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2Expanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2Expanded, float, float, f
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_double_double_double_int.cu
index 68e178089c..404dc5f8b4 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2SqrtExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2SqrtExpanded, double, dou
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_float_float_float_int.cu
index bcfa5d99d1..b4f8757d4a 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_sqrt_expanded_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2SqrtExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2SqrtExpanded, float, floa
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_double_double_double_int.cu
index 4cfb058a55..10cafac4a7 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@ namespace distance {
 namespace detail {
 
 template void distance<raft::distance::DistanceType::L2SqrtUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -29,7 +30,6 @@ template void distance<raft::distance::DistanceType::L2SqrtUnexpanded, double, d
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 }  // namespace detail
diff --git a/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_float_float_float_int.cu
index b8331977c4..6a5f43bc46 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_sqrt_unexpanded_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2SqrtUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2SqrtUnexpanded, float, fl
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_unexpanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/l2_unexpanded_double_double_double_int.cu
index fec7c2dce3..174a893d36 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_unexpanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_unexpanded_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2Unexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2Unexpanded, double, doubl
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/l2_unexpanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/l2_unexpanded_float_float_float_int.cu
index 8f83b9cfbb..0cc43b099f 100644
--- a/cpp/src/distance/distance/specializations/detail/l2_unexpanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/l2_unexpanded_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::L2Unexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::L2Unexpanded, float, float,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/lp_unexpanded_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/lp_unexpanded_double_double_double_int.cu
index 082fad62d1..66e765b1c0 100644
--- a/cpp/src/distance/distance/specializations/detail/lp_unexpanded_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/lp_unexpanded_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@ namespace distance {
 namespace detail {
 
 template void distance<raft::distance::DistanceType::LpUnexpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -29,7 +30,6 @@ template void distance<raft::distance::DistanceType::LpUnexpanded, double, doubl
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 }  // namespace detail
diff --git a/cpp/src/distance/distance/specializations/detail/lp_unexpanded_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/lp_unexpanded_float_float_float_int.cu
index 919d069e1d..6fbd52ac89 100644
--- a/cpp/src/distance/distance/specializations/detail/lp_unexpanded_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/lp_unexpanded_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::LpUnexpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::LpUnexpanded, float, float,
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/russel_rao_double_double_double_int.cu b/cpp/src/distance/distance/specializations/detail/russel_rao_double_double_double_int.cu
index 809f6ccb78..77e1fc70d0 100644
--- a/cpp/src/distance/distance/specializations/detail/russel_rao_double_double_double_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/russel_rao_double_double_double_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@ namespace distance {
 namespace detail {
 template void
 distance<raft::distance::DistanceType::RusselRaoExpanded, double, double, double, int>(
+  raft::resources const& handle,
   const double* x,
   const double* y,
   double* dist,
@@ -29,7 +30,6 @@ distance<raft::distance::DistanceType::RusselRaoExpanded, double, double, double
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   double metric_arg);
 
diff --git a/cpp/src/distance/distance/specializations/detail/russel_rao_float_float_float_int.cu b/cpp/src/distance/distance/specializations/detail/russel_rao_float_float_float_int.cu
index 831384fdd7..6f898310c7 100644
--- a/cpp/src/distance/distance/specializations/detail/russel_rao_float_float_float_int.cu
+++ b/cpp/src/distance/distance/specializations/detail/russel_rao_float_float_float_int.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,6 +20,7 @@ namespace raft {
 namespace distance {
 namespace detail {
 template void distance<raft::distance::DistanceType::RusselRaoExpanded, float, float, float, int>(
+  raft::resources const& handle,
   const float* x,
   const float* y,
   float* dist,
@@ -28,7 +29,6 @@ template void distance<raft::distance::DistanceType::RusselRaoExpanded, float, f
   int k,
   void* workspace,
   std::size_t worksize,
-  cudaStream_t stream,
   bool isRowMajor,
   float metric_arg);
 
diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt
index 9974fb6188..073681cd65 100644
--- a/cpp/test/CMakeLists.txt
+++ b/cpp/test/CMakeLists.txt
@@ -117,6 +117,7 @@ if(BUILD_TESTS)
     test/distance/dist_eucsqrt_exp.cu
     test/distance/dist_hamming.cu
     test/distance/dist_hellinger.cu
+    test/distance/dist_inner_product.cu
     test/distance/dist_jensen_shannon.cu
     test/distance/dist_kl_divergence.cu
     test/distance/dist_l1.cu
diff --git a/cpp/test/distance/dist_adj.cu b/cpp/test/distance/dist_adj.cu
index bbd06042c3..ce802e5138 100644
--- a/cpp/test/distance/dist_adj.cu
+++ b/cpp/test/distance/dist_adj.cu
@@ -132,7 +132,8 @@ class DistanceAdjTest : public ::testing::TestWithParam<DistanceAdjInputs<DataTy
                              DataType,
                              DataType,
                              uint8_t,
-                             threshold_final_op_>(x.data(),
+                             threshold_final_op_>(handle,
+                                                  x.data(),
                                                   y.data(),
                                                   dist.data(),
                                                   m,
@@ -141,7 +142,6 @@ class DistanceAdjTest : public ::testing::TestWithParam<DistanceAdjInputs<DataTy
                                                   workspace.data(),
                                                   workspace.size(),
                                                   threshold_op,
-                                                  stream,
                                                   isRowMajor);
     handle.sync_stream(stream);
   }
diff --git a/cpp/test/distance/dist_inner_product.cu b/cpp/test/distance/dist_inner_product.cu
new file mode 100644
index 0000000000..68ce4c841a
--- /dev/null
+++ b/cpp/test/distance/dist_inner_product.cu
@@ -0,0 +1,76 @@
+/*
+ * Copyright (c) 2018-2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "../test_utils.cuh"
+#include "distance_base.cuh"
+
+namespace raft {
+namespace distance {
+
+template <typename DataType>
+class DistanceInnerProduct
+  : public DistanceTest<raft::distance::DistanceType::InnerProduct, DataType> {
+};
+
+const std::vector<DistanceInputs<float>> inputsf = {
+  {0.001f, 10, 5, 32, true, 1234ULL},
+  {0.001f, 1024, 32, 1024, true, 1234ULL},
+  {0.001f, 32, 1024, 1024, true, 1234ULL},
+  {0.003f, 1024, 1024, 1024, true, 1234ULL},
+  {0.001f, 1024, 1024, 32, false, 1234ULL},
+  {0.001f, 1024, 32, 1024, false, 1234ULL},
+  {0.001f, 32, 1024, 1024, false, 1234ULL},
+  {0.003f, 1024, 1024, 1024, false, 1234ULL},
+};
+typedef DistanceInnerProduct<float> DistanceInnerProductF;
+TEST_P(DistanceInnerProductF, Result)
+{
+  int m = params.isRowMajor ? params.m : params.n;
+  int n = params.isRowMajor ? params.n : params.m;
+
+  ASSERT_TRUE(devArrMatch(
+    dist_ref.data(), dist.data(), m, n, raft::CompareApprox<float>(params.tolerance), stream));
+}
+INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceInnerProductF, ::testing::ValuesIn(inputsf));
+
+const std::vector<DistanceInputs<double>> inputsd = {
+  {0.001, 1024, 1024, 32, true, 1234ULL},
+  {0.001, 1024, 32, 1024, true, 1234ULL},
+  {0.001, 32, 1024, 1024, true, 1234ULL},
+  {0.003, 1024, 1024, 1024, true, 1234ULL},
+  {0.001f, 1024, 1024, 32, false, 1234ULL},
+  {0.001f, 1024, 32, 1024, false, 1234ULL},
+  {0.001f, 32, 1024, 1024, false, 1234ULL},
+  {0.003f, 1024, 1024, 1024, false, 1234ULL},
+};
+typedef DistanceInnerProduct<double> DistanceInnerProductD;
+TEST_P(DistanceInnerProductD, Result)
+{
+  int m = params.isRowMajor ? params.m : params.n;
+  int n = params.isRowMajor ? params.n : params.m;
+
+  ASSERT_TRUE(devArrMatch(
+    dist_ref.data(), dist.data(), m, n, raft::CompareApprox<double>(params.tolerance), stream));
+}
+INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceInnerProductD, ::testing::ValuesIn(inputsd));
+
+class BigMatrixInnerProduct
+  : public BigMatrixDistanceTest<raft::distance::DistanceType::InnerProduct> {
+};
+TEST_F(BigMatrixInnerProduct, Result) {}
+
+}  // end namespace distance
+}  // end namespace raft
diff --git a/cpp/test/distance/distance_base.cuh b/cpp/test/distance/distance_base.cuh
index be7b2b1de8..278fde0c45 100644
--- a/cpp/test/distance/distance_base.cuh
+++ b/cpp/test/distance/distance_base.cuh
@@ -122,6 +122,28 @@ __global__ void naiveCosineDistanceKernel(
   dist[outidx] = (DataType)1.0 - acc_ab / (raft::sqrt(acc_a) * raft::sqrt(acc_b));
 }
 
+template <typename DataType>
+__global__ void naiveInnerProductKernel(
+  DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor)
+{
+  int midx = threadIdx.x + blockIdx.x * blockDim.x;
+  int nidx = threadIdx.y + blockIdx.y * blockDim.y;
+  if (midx >= m || nidx >= n) { return; }
+
+  DataType acc_ab = DataType(0);
+
+  for (int i = 0; i < k; ++i) {
+    int xidx = isRowMajor ? i + midx * k : i * m + midx;
+    int yidx = isRowMajor ? i + nidx * k : i * n + nidx;
+    auto a   = x[xidx];
+    auto b   = y[yidx];
+    acc_ab += a * b;
+  }
+
+  int outidx   = isRowMajor ? midx * n + nidx : midx + m * nidx;
+  dist[outidx] = acc_ab;
+}
+
 template <typename DataType>
 __global__ void naiveHellingerDistanceKernel(
   DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor)
@@ -348,6 +370,9 @@ void naiveDistance(DataType* dist,
       naiveHammingDistanceKernel<DataType>
         <<<nblks, TPB, 0, stream>>>(dist, x, y, m, n, k, isRowMajor);
       break;
+    case raft::distance::DistanceType::InnerProduct:
+      naiveInnerProductKernel<DataType><<<nblks, TPB, 0, stream>>>(dist, x, y, m, n, k, isRowMajor);
+      break;
     case raft::distance::DistanceType::JensenShannon:
       naiveJensenShannonDistanceKernel<DataType>
         <<<nblks, TPB, 0, stream>>>(dist, x, y, m, n, k, isRowMajor);
@@ -385,7 +410,8 @@ template <typename DataType>
 }
 
 template <raft::distance::DistanceType distanceType, typename DataType, typename layout>
-void distanceLauncher(DataType* x,
+void distanceLauncher(raft::device_resources const& handle,
+                      DataType* x,
                       DataType* y,
                       DataType* dist,
                       DataType* dist2,
@@ -394,11 +420,8 @@ void distanceLauncher(DataType* x,
                       int k,
                       DistanceInputs<DataType>& params,
                       DataType threshold,
-                      cudaStream_t stream,
                       DataType metric_arg = 2.0f)
 {
-  raft::device_resources handle(stream);
-
   auto x_v    = make_device_matrix_view<DataType, int, layout>(x, m, k);
   auto y_v    = make_device_matrix_view<DataType, int, layout>(y, n, k);
   auto dist_v = make_device_matrix_view<DataType, int, layout>(dist, m, n);
@@ -454,7 +477,8 @@ class DistanceTest : public ::testing::TestWithParam<DistanceInputs<DataType>> {
     DataType threshold = -10000.f;
 
     if (isRowMajor) {
-      distanceLauncher<distanceType, DataType, layout_c_contiguous>(x.data(),
+      distanceLauncher<distanceType, DataType, layout_c_contiguous>(handle,
+                                                                    x.data(),
                                                                     y.data(),
                                                                     dist.data(),
                                                                     dist2.data(),
@@ -463,11 +487,11 @@ class DistanceTest : public ::testing::TestWithParam<DistanceInputs<DataType>> {
                                                                     k,
                                                                     params,
                                                                     threshold,
-                                                                    stream,
                                                                     metric_arg);
 
     } else {
-      distanceLauncher<distanceType, DataType, layout_f_contiguous>(x.data(),
+      distanceLauncher<distanceType, DataType, layout_f_contiguous>(handle,
+                                                                    x.data(),
                                                                     y.data(),
                                                                     dist.data(),
                                                                     dist2.data(),
@@ -476,7 +500,6 @@ class DistanceTest : public ::testing::TestWithParam<DistanceInputs<DataType>> {
                                                                     k,
                                                                     params,
                                                                     threshold,
-                                                                    stream,
                                                                     metric_arg);
     }
     handle.sync_stream(stream);
@@ -500,20 +523,8 @@ class BigMatrixDistanceTest : public ::testing::Test {
     auto testInfo = testing::UnitTest::GetInstance()->current_test_info();
     common::nvtx::range fun_scope("test::%s/%s", testInfo->test_suite_name(), testInfo->name());
 
-    size_t worksize = raft::distance::getWorkspaceSize<distanceType, float, float, float>(
-      x.data(), x.data(), m, n, k);
-    rmm::device_uvector<char> workspace(worksize, handle.get_stream());
-    raft::distance::distance<distanceType, float, float, float>(x.data(),
-                                                                x.data(),
-                                                                dist.data(),
-                                                                m,
-                                                                n,
-                                                                k,
-                                                                workspace.data(),
-                                                                worksize,
-                                                                handle.get_stream(),
-                                                                true,
-                                                                0.0f);
+    raft::distance::distance<distanceType, float, float, float>(
+      handle, x.data(), x.data(), dist.data(), m, n, k, true, 0.0f);
 
     RAFT_CUDA_TRY(cudaStreamSynchronize(handle.get_stream()));
   }
diff --git a/python/pylibraft/pylibraft/distance/pairwise_distance.pyx b/python/pylibraft/pylibraft/distance/pairwise_distance.pyx
index 9649531b61..3037b9a725 100644
--- a/python/pylibraft/pylibraft/distance/pairwise_distance.pyx
+++ b/python/pylibraft/pylibraft/distance/pairwise_distance.pyx
@@ -79,7 +79,7 @@ DISTANCE_TYPES = {
     "kl_divergence": DistanceType.KLDivergence,
     "minkowski": DistanceType.LpUnexpanded,
     "russellrao": DistanceType.RusselRaoExpanded,
-    "dice": DistanceType.DiceExpanded
+    "dice": DistanceType.DiceExpanded,
 }
 
 SUPPORTED_DISTANCES = ["euclidean", "l1", "cityblock", "l2", "inner_product",
diff --git a/python/pylibraft/pylibraft/test/test_distance.py b/python/pylibraft/pylibraft/test/test_distance.py
index dd6050a098..2c0a842fe5 100644
--- a/python/pylibraft/pylibraft/test/test_distance.py
+++ b/python/pylibraft/pylibraft/test/test_distance.py
@@ -21,8 +21,8 @@
 from pylibraft.distance import pairwise_distance
 
 
-@pytest.mark.parametrize("n_rows", [100])
-@pytest.mark.parametrize("n_cols", [100])
+@pytest.mark.parametrize("n_rows", [32, 100])
+@pytest.mark.parametrize("n_cols", [40, 100])
 @pytest.mark.parametrize(
     "metric",
     [
@@ -36,6 +36,7 @@
         "russellrao",
         "cosine",
         "sqeuclidean",
+        "inner_product",
     ],
 )
 @pytest.mark.parametrize("inplace", [True, False])
@@ -57,7 +58,10 @@ def test_distance(n_rows, n_cols, inplace, metric, order, dtype):
 
     output = np.zeros((n_rows, n_rows), dtype=dtype)
 
-    expected = cdist(input1, input1, metric)
+    if metric == "inner_product":
+        expected = np.matmul(input1, input1.T)
+    else:
+        expected = cdist(input1, input1, metric)
 
     expected[expected <= 1e-5] = 0.0