Remove raft dependency

AlgorithmLauncher.cpp

@@ -3,8 +3,9 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
+#include "macros.hpp"
+
 #include <AlgorithmLauncher.hpp>
-#include <raft/util/cuda_rt_essentials.hpp>
 
 AlgorithmLauncher::AlgorithmLauncher(cudaKernel_t k, cudaLibrary_t lib) : kernel{k}, library{lib} {}
 
@@ -43,7 +44,7 @@ void AlgorithmLauncher::call(
   config.numAttrs         = 0;
   config.attrs            = NULL;
 
-  RAFT_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
+  RTCX_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
 }
 
 void AlgorithmLauncher::call_cooperative(
@@ -61,5 +62,5 @@ void AlgorithmLauncher::call_cooperative(
   config.numAttrs         = 1;
   config.attrs            = attribute;
 
-  RAFT_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
+  RTCX_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
 }

AlgorithmPlanner.cpp

@@ -4,12 +4,11 @@
  */
 
 #include "cuda_runtime.h"
+#include "macros.hpp"
 #include "nvJitLink.h"
 
 #include <AlgorithmPlanner.hpp>
 #include <nvjitlink_checker.hpp>
-#include <raft/core/logger.hpp>
-#include <raft/util/cuda_rt_essentials.hpp>
 
 #include <chrono>
 #include <iterator>
@@ -47,13 +46,6 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::get_launcher()
   std::unique_lock<std::shared_mutex> write_lock(jit_cache_.mutex);
   if (auto it = launchers.find(launch_key); it != launchers.end()) { return it->second; }
 
-  std::string log_message =
-    "JIT compiling launcher for kernel: " + this->entrypoint + " and device functions: ";
-  for (auto const& fragment : this->fragments) {
-    log_message += std::string{fragment->get_key()} + ",";
-  }
-  log_message.pop_back();
-  RAFT_LOG_DEBUG("%s", log_message.c_str());
   auto launcher         = this->build();
   launchers[launch_key] = launcher;
   return launcher;
@@ -64,9 +56,9 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
   int device = 0;
   int major  = 0;
   int minor  = 0;
-  RAFT_CUDA_TRY(cudaGetDevice(&device));
-  RAFT_CUDA_TRY(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
-  RAFT_CUDA_TRY(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
+  RTCX_CUDA_TRY(cudaGetDevice(&device));
+  RTCX_CUDA_TRY(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
+  RTCX_CUDA_TRY(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
 
   std::string archs = "-arch=sm_" + std::to_string((major * 10 + minor));
 
@@ -101,15 +93,15 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
   check_nvjitlink_result(handle, result);
 
   result = nvJitLinkDestroy(&handle);
-  RAFT_EXPECTS(result == NVJITLINK_SUCCESS, "nvJitLinkDestroy failed");
+  RTCX_EXPECTS(result == NVJITLINK_SUCCESS, "nvJitLinkDestroy failed");
 
   // cubin is linked, so now load it
   cudaLibrary_t library;
-  RAFT_CUDA_TRY(
+  RTCX_CUDA_TRY(
     cudaLibraryLoadData(&library, cubin.get(), nullptr, nullptr, 0, nullptr, nullptr, 0));
 
   cudaKernel_t kernel;
-  RAFT_CUDA_TRY(cudaLibraryGetKernel(&kernel, library, this->entrypoint.c_str()));
+  RTCX_CUDA_TRY(cudaLibraryGetKernel(&kernel, library, this->entrypoint.c_str()));
 
   return std::make_shared<AlgorithmLauncher>(kernel, library);
 }

CMakeLists.txt

@@ -62,9 +62,6 @@ if(NOT TARGET nvtx3::nvtx3-cpp)
   rapids_cpm_nvtx3()
 endif()
 
-rapids_cpm_find(raft "${RAPIDS_VERSION_MAJOR_MINOR}" GLOBAL_TARGETS raft::raft GIT_REPOSITORY
-                https://github.com/rapidsai/raft.git)
-
 add_library(rtcx STATIC hash.cpp rtcx.cpp AlgorithmLauncher.cpp AlgorithmPlanner.cpp
                         FragmentEntry.cpp nvjitlink_checker.cpp NVRTCLTOFragmentCompiler.cpp)
 add_library(rtcx::rtcx ALIAS rtcx)
@@ -79,7 +76,7 @@ set_target_properties(rtcx
 target_include_directories(rtcx PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
                            PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
 
-target_link_libraries(rtcx PRIVATE zstd ${CMAKE_DL_LIBS} nvtx3::nvtx3-cpp raft::raft)
+target_link_libraries(rtcx PRIVATE zstd ${CMAKE_DL_LIBS} nvtx3::nvtx3-cpp)
 
 if(RTCX_STATIC_LINK_NVRTC)
   target_link_libraries(rtcx PRIVATE CUDA::nvrtc_static CUDA::nvrtc_builtins_static)

NVRTCLTOFragmentCompiler.cpp

@@ -4,11 +4,10 @@
  */
 
 #include "cuda.h"
+#include "macros.hpp"
 
 #include <NVRTCLTOFragmentCompiler.hpp>
 #include <nvrtc.h>
-#include <raft/core/error.hpp>
-#include <raft/util/cuda_rt_essentials.hpp>
 
 #include <mutex>
 
@@ -17,17 +16,17 @@
     nvrtcResult result = _call;                                                \
     std::string error_string =                                                 \
       std::string("nvrtc error: ") + std::string(nvrtcGetErrorString(result)); \
-    RAFT_EXPECTS(result == NVRTC_SUCCESS, "%s", error_string.c_str());         \
+    RTCX_EXPECTS(result == NVRTC_SUCCESS, "%s", error_string.c_str());         \
   }
 
 NVRTCLTOFragmentCompiler::NVRTCLTOFragmentCompiler()
 {
   int device = 0;
   int major  = 0;
   int minor  = 0;
-  RAFT_CUDA_TRY(cudaGetDevice(&device));
-  RAFT_CUDA_TRY(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
-  RAFT_CUDA_TRY(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
+  RTCX_CUDA_TRY(cudaGetDevice(&device));
+  RTCX_CUDA_TRY(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
+  RTCX_CUDA_TRY(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
 
   this->standard_compile_opts = {
     std::string{"-arch=sm_" + std::to_string((major * 10 + minor))},
@@ -80,7 +79,7 @@ std::unique_ptr<UDFFatbinFragment> NVRTCLTOFragmentCompiler::compile(std::string
       NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &log_size));
       std::unique_ptr<char[]> log{new char[log_size]};
       NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log.get()));
-      RAFT_FAIL("nvrtc compile error log: \n%s", log.get());
+      RTCX_FAIL("nvrtc compile error log: \n%s", log.get());
     }
   } catch (...) {
     NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

macros.hpp

@@ -0,0 +1,61 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <cuda_runtime.h>
+
+#include <cstdio>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+#define SET_ERROR_MSG(msg, location_prefix, fmt, ...)                                            \
+  do {                                                                                           \
+    int size1 = std::snprintf(nullptr, 0, "%s", location_prefix);                                \
+    int size2 = std::snprintf(nullptr, 0, "file=%s line=%d: ", __FILE__, __LINE__);              \
+    int size3 = std::snprintf(nullptr, 0, fmt, ##__VA_ARGS__);                                   \
+    if (size1 < 0 || size2 < 0 || size3 < 0)                                                     \
+      throw std::runtime_error("Error in snprintf, cannot format rtcx exception.");              \
+    auto size = size1 + size2 + size3 + 1; /* +1 for final '\0' */                               \
+    std::vector<char> buf(size);                                                                 \
+    std::snprintf(buf.data(), size1 + 1 /* +1 for '\0' */, "%s", location_prefix);               \
+    std::snprintf(                                                                               \
+      buf.data() + size1, size2 + 1 /* +1 for '\0' */, "file=%s line=%d: ", __FILE__, __LINE__); \
+    std::snprintf(buf.data() + size1 + size2, size3 + 1 /* +1 for '\0' */, fmt, ##__VA_ARGS__);  \
+    msg += std::string(buf.data(), buf.data() + size - 1); /* -1 to remove final '\0' */         \
+  } while (0)
+
+#define RTCX_CUDA_TRY(call)                        \
+  do {                                             \
+    cudaError_t const status = call;               \
+    if (status != cudaSuccess) {                   \
+      cudaGetLastError();                          \
+      std::string msg{};                           \
+      SET_ERROR_MSG(msg,                           \
+                    "CUDA error encountered at: ", \
+                    "call='%s', Reason=%s:%s",     \
+                    #call,                         \
+                    cudaGetErrorName(status),      \
+                    cudaGetErrorString(status));   \
+      throw std::runtime_error(msg);               \
+    }                                              \
+  } while (0)
+
+#define RTCX_EXPECTS(cond, fmt, ...)                              \
+  do {                                                            \
+    if (!(cond)) {                                                \
+      std::string msg{};                                          \
+      SET_ERROR_MSG(msg, "RTCX failure at ", fmt, ##__VA_ARGS__); \
+      throw std::logic_error(msg);                                \
+    }                                                             \
+  } while (0)
+
+#define RTCX_FAIL(fmt, ...)                                     \
+  do {                                                          \
+    std::string msg{};                                          \
+    SET_ERROR_MSG(msg, "RTCX failure at ", fmt, ##__VA_ARGS__); \
+    throw std::logic_error(msg);                                \
+  } while (0)

nvjitlink_checker.cpp

@@ -3,9 +3,10 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
+#include "macros.hpp"
+
 #include <nvJitLink.h>
 #include <nvjitlink_checker.hpp>
-#include <raft/core/error.hpp>
 
 #include <memory>
 #include <string>
@@ -21,6 +22,6 @@ void check_nvjitlink_result(nvJitLinkHandle handle, nvJitLinkResult result)
       result = nvJitLinkGetErrorLog(handle, log.get());
       if (result == NVJITLINK_SUCCESS) { error_msg += "\n" + std::string(log.get()); }
     }
-    RAFT_FAIL("AlgorithmPlanner nvJITLink error log: %s", error_msg.c_str());
+    RTCX_FAIL("AlgorithmPlanner nvJITLink error log: %s", error_msg.c_str());
   }
 }