Test rtcx refactoring

Author: arhag23

cpp/cmake/thirdparty/get_rtcx.cmake

@@ -22,15 +22,15 @@ function(find_and_configure_rtcx VERSION)
     BUILD_EXPORT_SET    cuvs-static-exports
     INSTALL_EXPORT_SET  cuvs-static-exports
     CPM_ARGS
-    GIT_REPOSITORY https://github.com/rapidsai/librtcx
-    GIT_TAG b5e753e9427e331c02f8c370143cb0d5f40da0fd
+    GIT_REPOSITORY https://github.com/arhag23/librtcx
+    GIT_TAG 67b5e05f5282fca9ad04d44b7dc93d75d89f1ad1
     GIT_SHALLOW FALSE
   )
 
   # When CPM fetches from source (add_subdirectory), generate_jit_lto_kernels.cmake is not
   # auto-included. Include it explicitly so necessary functions are available.
   if(rtcx_ADDED OR DEFINED CPM_rtcx_SOURCE)
-    include("${rtcx_SOURCE_DIR}/generate_jit_lto_kernels.cmake")
+    include("${rtcx_SOURCE_DIR}/cmake/modules/generate_jit_lto_kernels.cmake")
   endif()
 endfunction()
 

cpp/src/distance/detail/pairwise_matrix/jit_lto_kernels/pairwise_matrix_planner.hpp

@@ -5,8 +5,8 @@
 
 #pragma once
 
-#include <AlgorithmPlanner.hpp>
 #include <cuvs/detail/jit_lto/pairwise_matrix/pairwise_matrix_fragments.hpp>
+#include <rtcx/algorithm_planner.hpp>
 
 namespace cuvs::distance::detail {
 
@@ -20,7 +20,7 @@ template <typename DistanceTag_,
           typename FinOpTag_,
           typename LayoutTag_,
           int Veclen_>
-struct PairwiseMatrixPlanner : AlgorithmPlanner {
+struct PairwiseMatrixPlanner : rtcx::algorithm_planner {
   using DistanceTag = DistanceTag_;
   using DataTag     = DataTag_;
   using AccTag      = AccTag_;
@@ -31,9 +31,12 @@ struct PairwiseMatrixPlanner : AlgorithmPlanner {
 
   static constexpr int Veclen = Veclen_;
 
-  inline static LauncherJitCache launcher_jit_cache{};
+  inline static rtcx::launcher_jit_cache launcher_jit_cache{};
 
-  PairwiseMatrixPlanner() : AlgorithmPlanner(kPairwiseMatrixJitEntrypoint, launcher_jit_cache) {}
+  PairwiseMatrixPlanner()
+    : rtcx::algorithm_planner(kPairwiseMatrixJitEntrypoint, launcher_jit_cache)
+  {
+  }
 
   void add_entrypoint()
   {

cpp/src/neighbors/detail/cagra/jit_lto_kernels/cagra_jit_launcher_factory.hpp

@@ -12,9 +12,9 @@
 #include "search_multi_kernel_planner.hpp"
 #include "search_single_cta_planner.hpp"
 
-#include <AlgorithmLauncher.hpp>
 #include <cuvs/detail/jit_lto/cagra/cagra_fragments.hpp>
 #include <cuvs/distance/distance.hpp>
+#include <rtcx/algorithm_launcher.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -34,12 +34,12 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_single_cta_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> build_single_cta_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   bool topk_by_bitonic_sort,
   bool bitonic_sort_and_merge_multi_warps,
   bool persistent,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment)
 {
   single_cta_search::CagraSingleCtaSearchPlanner<DataTag,
                                                  IndexTag,
@@ -90,9 +90,9 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_multi_cta_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> build_multi_cta_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment)
 {
   multi_cta_search::CagraMultiCtaSearchPlanner<DataTag,
                                                IndexTag,
@@ -139,7 +139,7 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_single_cta_mp_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> build_single_cta_mp_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   bool topk_by_bitonic_sort,
   bool bitonic_sort_and_merge_multi_warps)
@@ -191,7 +191,7 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_multi_cta_mp_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> build_multi_cta_mp_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc)
 {
   multi_cta_search::CagraMultiCtaMpSearchPlanner<DataTag,
@@ -239,10 +239,10 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_multi_kernel_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> build_multi_kernel_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   const char* linked_kernel_name,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment)
 {
   multi_kernel_search::CagraMultiKernelSearchPlanner<DataTag,
                                                      IndexTag,
@@ -291,8 +291,8 @@ template <typename DataTag,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT>
-std::shared_ptr<AlgorithmLauncher> build_apply_filter_only_launcher(
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment)
+std::shared_ptr<rtcx::algorithm_launcher> build_apply_filter_only_launcher(
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment)
 {
   multi_kernel_search::CagraMultiKernelSearchPlanner<DataTag,
                                                      IndexTag,
@@ -309,20 +309,20 @@ std::shared_ptr<AlgorithmLauncher> build_apply_filter_only_launcher(
 
 }  // namespace cagra_jit_launcher_factory_detail
 
-/// Build a JIT AlgorithmLauncher for single-CTA CAGRA search (runtime VPQ / metric → tag
+/// Build a JIT rtcx::algorithm_launcher for single-CTA CAGRA search (runtime VPQ / metric → tag
 /// dispatch). `SampleFilterJitTag` is `cuvs::neighbors::detail::tag_filter_none`,
 /// `tag_filter_bitset`, or use `sample_filter_jit_tag_t<SAMPLE_FILTER_T>`.
 template <typename DataT,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag>
-std::shared_ptr<AlgorithmLauncher> make_cagra_single_cta_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_single_cta_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   bool topk_by_bitonic_sort,
   bool bitonic_sort_and_merge_multi_warps,
   bool persistent,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment = nullptr)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment = nullptr)
 {
   using DataTag   = decltype(get_data_type_tag<DataT>());
   using IndexTag  = decltype(get_index_type_tag<IndexT>());
@@ -389,15 +389,15 @@ std::shared_ptr<AlgorithmLauncher> make_cagra_single_cta_jit_launcher(
     std::move(sample_filter_udf_fragment));
 }
 
-/// Build a JIT AlgorithmLauncher for multi-CTA CAGRA search.
+/// Build a JIT rtcx::algorithm_launcher for multi-CTA CAGRA search.
 template <typename DataT,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag>
-std::shared_ptr<AlgorithmLauncher> make_cagra_multi_cta_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_multi_cta_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment = nullptr)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment = nullptr)
 {
   using DataTag   = decltype(get_data_type_tag<DataT>());
   using IndexTag  = decltype(get_index_type_tag<IndexT>());
@@ -452,13 +452,13 @@ std::shared_ptr<AlgorithmLauncher> make_cagra_multi_cta_jit_launcher(
     dataset_desc, std::move(sample_filter_udf_fragment));
 }
 
-/// Build a JIT AlgorithmLauncher for the multi-partition single-CTA CAGRA search.
+/// Build a JIT rtcx::algorithm_launcher for the multi-partition single-CTA CAGRA search.
 template <typename DataT,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag>
-std::shared_ptr<AlgorithmLauncher> make_cagra_single_cta_mp_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_single_cta_mp_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   bool topk_by_bitonic_sort,
   bool bitonic_sort_and_merge_multi_warps)
@@ -516,13 +516,13 @@ std::shared_ptr<AlgorithmLauncher> make_cagra_single_cta_mp_jit_launcher(
     dataset_desc, topk_by_bitonic_sort, bitonic_sort_and_merge_multi_warps);
 }
 
-/// Build a JIT AlgorithmLauncher for the multi-partition multi-CTA CAGRA search.
+/// Build a JIT rtcx::algorithm_launcher for the multi-partition multi-CTA CAGRA search.
 template <typename DataT,
           typename IndexT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag>
-std::shared_ptr<AlgorithmLauncher> make_cagra_multi_cta_mp_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_multi_cta_mp_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc)
 {
   using DataTag   = decltype(get_data_type_tag<DataT>());
@@ -577,8 +577,8 @@ std::shared_ptr<AlgorithmLauncher> make_cagra_multi_cta_mp_jit_launcher(
                                                                         SourceIndexT>(dataset_desc);
 }
 
-/// Build a JIT AlgorithmLauncher for multi-kernel CAGRA helpers that need `setup_workspace` and
-/// `compute_distance` linked (e.g. `random_pickup`, `compute_distance_to_child_nodes`). For
+/// Build a JIT rtcx::algorithm_launcher for multi-kernel CAGRA helpers that need `setup_workspace`
+/// and `compute_distance` linked (e.g. `random_pickup`, `compute_distance_to_child_nodes`). For
 /// `apply_filter_kernel` only, use `make_cagra_apply_filter_jit_launcher` instead. Use
 /// `SampleFilterJitTag = tag_cagra_jit_sample_filter_link_absent` (default) when the kernel does
 /// not link `sample_filter`; otherwise `sample_filter_jit_tag_t<SAMPLE_FILTER_T>` or a
@@ -588,10 +588,10 @@ template <typename DataT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag = tag_cagra_jit_sample_filter_link_absent>
-std::shared_ptr<AlgorithmLauncher> make_cagra_multi_kernel_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_multi_kernel_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
   const char* linked_kernel_name,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment = nullptr)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment = nullptr)
 {
   using DataTag   = decltype(get_data_type_tag<DataT>());
   using IndexTag  = decltype(get_index_type_tag<IndexT>());
@@ -655,9 +655,9 @@ template <typename DataT,
           typename DistanceT,
           typename SourceIndexT,
           typename SampleFilterJitTag>
-std::shared_ptr<AlgorithmLauncher> make_cagra_apply_filter_jit_launcher(
+std::shared_ptr<rtcx::algorithm_launcher> make_cagra_apply_filter_jit_launcher(
   const dataset_descriptor_host<DataT, IndexT, DistanceT>& dataset_desc,
-  std::unique_ptr<UDFFatbinFragment> sample_filter_udf_fragment = nullptr)
+  std::unique_ptr<rtcx::udf_fatbin_fragment> sample_filter_udf_fragment = nullptr)
 {
   using DataTag   = decltype(get_data_type_tag<DataT>());
   using IndexTag  = decltype(get_index_type_tag<IndexT>());

cpp/src/neighbors/detail/cagra/jit_lto_kernels/cagra_planner_base.hpp

@@ -5,12 +5,12 @@
 
 #pragma once
 
-#include <AlgorithmPlanner.hpp>
 #include <cuvs/detail/jit_lto/cagra/cagra_fragments.hpp>
 #include <cuvs/detail/jit_lto/common_fragments.hpp>
 #include <cuvs/distance/distance.hpp>
 #include <raft/core/error.hpp>
 #include <raft/core/logger.hpp>
+#include <rtcx/algorithm_planner.hpp>
 
 #include <cstdint>
 #include <memory>
@@ -25,16 +25,16 @@ template <typename DataTag_,
           typename QueryTag_,
           typename CodebookTag_,
           typename SampleFilterJitTag_ = tag_cagra_jit_sample_filter_link_absent>
-struct CagraPlannerBase : AlgorithmPlanner {
+struct CagraPlannerBase : rtcx::algorithm_planner {
   using DataTag            = DataTag_;
   using IndexTag           = IndexTag_;
   using DistanceTag        = DistanceTag_;
   using QueryTag           = QueryTag_;
   using CodebookTag        = CodebookTag_;
   using SampleFilterJitTag = SampleFilterJitTag_;
 
-  explicit CagraPlannerBase(std::string entrypoint, LauncherJitCache& jit_cache)
-    : AlgorithmPlanner(std::move(entrypoint), jit_cache)
+  explicit CagraPlannerBase(std::string entrypoint, rtcx::launcher_jit_cache& jit_cache)
+    : rtcx::algorithm_planner(std::move(entrypoint), jit_cache)
   {
   }
 
@@ -385,7 +385,8 @@ struct CagraPlannerBase : AlgorithmPlanner {
       static_cast<unsigned>(dataset_block_dim));
   }
 
-  void add_sample_filter_device_function(std::unique_ptr<UDFFatbinFragment> udf_fragment = nullptr)
+  void add_sample_filter_device_function(
+    std::unique_ptr<rtcx::udf_fatbin_fragment> udf_fragment = nullptr)
   {
     if constexpr (std::is_same_v<SampleFilterJitTag_, tag_cagra_jit_sample_filter_link_absent>) {
       RAFT_EXPECTS(udf_fragment == nullptr, "Unexpected CAGRA sample-filter UDF fragment");

cpp/src/neighbors/detail/cagra/jit_lto_kernels/kernel_def.hpp

@@ -20,7 +20,7 @@
 namespace cuvs::neighbors::cagra::detail {
 
 // Function types for extern "C" __global__ JIT entry points — must match cudaLibraryGetKernel /
-// AlgorithmLauncher::dispatch signatures exactly (see static_assert in each *_kernel.cu).
+// rtcx::algorithm_launcher::dispatch signatures exactly (see static_assert in each *_kernel.cu).
 
 template <typename DataT, typename IndexT, typename DistanceT, typename SourceIndexT>
 using search_single_cta_kernel_func_t =

cpp/src/neighbors/detail/cagra/jit_lto_kernels/sample_filter_udf.cuh

@@ -7,8 +7,8 @@
 
 #include "../cagra_filter_payload.hpp"
 
-#include <NVRTCLTOFragmentCompiler.hpp>
 #include <raft/core/error.hpp>
+#include <rtcx/nvrtc_lto_fragment_compiler.hpp>
 
 #include <cstdint>
 #include <memory>
@@ -66,7 +66,7 @@ __device__ bool sample_filter<source_index_t>(uint32_t query_id,
 }
 
 template <typename SourceIndexT, typename SampleFilterT>
-std::unique_ptr<UDFFatbinFragment> make_cagra_sample_filter_udf_fragment(
+std::unique_ptr<rtcx::udf_fatbin_fragment> make_cagra_sample_filter_udf_fragment(
   const SampleFilterT& sample_filter)
 {
   const auto* udf = get_cagra_udf_filter(sample_filter);
@@ -83,7 +83,7 @@ std::unique_ptr<UDFFatbinFragment> make_cagra_sample_filter_udf_fragment(
   key += udf->function_name;
   key += ":";
   key += code;
-  return nvrtc_compiler().compile(key, code);
+  return rtcx::nvrtc_compiler().compile(key, code);
 }
 
 }  // namespace cuvs::neighbors::cagra::detail

cpp/src/neighbors/detail/cagra/jit_lto_kernels/search_multi_cta_planner.hpp

@@ -20,7 +20,7 @@ template <typename DataTag,
           typename SampleFilterJitTag>
 struct CagraMultiCtaSearchPlanner
   : CagraPlannerBase<DataTag, IndexTag, DistanceTag, QueryTag, CodebookTag, SampleFilterJitTag> {
-  static inline LauncherJitCache launcher_jit_cache{};
+  static inline rtcx::launcher_jit_cache launcher_jit_cache{};
 
   CagraMultiCtaSearchPlanner(cuvs::distance::DistanceType /*metric*/,
                              uint32_t /*team_size*/,
@@ -49,7 +49,7 @@ template <typename DataTag,
           typename SampleFilterJitTag>
 struct CagraMultiCtaMpSearchPlanner
   : CagraPlannerBase<DataTag, IndexTag, DistanceTag, QueryTag, CodebookTag, SampleFilterJitTag> {
-  static inline LauncherJitCache launcher_jit_cache{};
+  static inline rtcx::launcher_jit_cache launcher_jit_cache{};
 
   CagraMultiCtaMpSearchPlanner(cuvs::distance::DistanceType /*metric*/,
                                uint32_t /*team_size*/,

cpp/src/neighbors/detail/cagra/jit_lto_kernels/search_multi_kernel_planner.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -21,7 +21,7 @@ template <typename DataTag,
           typename SampleFilterJitTag = tag_cagra_jit_sample_filter_link_absent>
 struct CagraMultiKernelSearchPlanner
   : CagraPlannerBase<DataTag, IndexTag, DistanceTag, QueryTag, CodebookTag, SampleFilterJitTag> {
-  static inline LauncherJitCache launcher_jit_cache{};
+  static inline rtcx::launcher_jit_cache launcher_jit_cache{};
 
   /// Kernels that only need `sample_filter` + one linked TU (e.g. `apply_filter_kernel`): no
   /// `setup_workspace` / `compute_distance` fragments. Metric / team / VPQ are unused.

cpp/src/neighbors/detail/cagra/jit_lto_kernels/search_single_cta_planner.hpp

@@ -21,7 +21,7 @@ template <typename DataTag,
           typename SampleFilterJitTag>
 struct CagraSingleCtaSearchPlanner
   : CagraPlannerBase<DataTag, IndexTag, DistanceTag, QueryTag, CodebookTag, SampleFilterJitTag> {
-  static inline LauncherJitCache launcher_jit_cache{};
+  static inline rtcx::launcher_jit_cache launcher_jit_cache{};
 
   CagraSingleCtaSearchPlanner(cuvs::distance::DistanceType /*metric*/,
                               bool /*topk_by_bitonic_sort*/,
@@ -114,7 +114,7 @@ template <typename DataTag,
           typename SampleFilterJitTag>
 struct CagraSingleCtaMpSearchPlanner
   : CagraPlannerBase<DataTag, IndexTag, DistanceTag, QueryTag, CodebookTag, SampleFilterJitTag> {
-  static inline LauncherJitCache launcher_jit_cache{};
+  static inline rtcx::launcher_jit_cache launcher_jit_cache{};
 
   CagraSingleCtaMpSearchPlanner(cuvs::distance::DistanceType /*metric*/,
                                 bool /*topk_by_bitonic_sort*/,