From 75f90a8503766ce37d62f4418449315874bf2d87 Mon Sep 17 00:00:00 2001 From: Arha Gatram Date: Fri, 10 Jul 2026 15:08:08 -0700 Subject: [PATCH 1/2] Remove raft dependency --- AlgorithmLauncher.cpp | 7 +++-- AlgorithmPlanner.cpp | 17 +++++----- CMakeLists.txt | 5 +-- NVRTCLTOFragmentCompiler.cpp | 13 ++++---- macros.hpp | 61 ++++++++++++++++++++++++++++++++++++ nvjitlink_checker.cpp | 5 +-- 6 files changed, 83 insertions(+), 25 deletions(-) create mode 100644 macros.hpp diff --git a/AlgorithmLauncher.cpp b/AlgorithmLauncher.cpp index 981575f..960ddce 100644 --- a/AlgorithmLauncher.cpp +++ b/AlgorithmLauncher.cpp @@ -3,8 +3,9 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "macros.hpp" + #include -#include 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)); } diff --git a/AlgorithmPlanner.cpp b/AlgorithmPlanner.cpp index 9d9da43..da419c1 100644 --- a/AlgorithmPlanner.cpp +++ b/AlgorithmPlanner.cpp @@ -4,12 +4,11 @@ */ #include "cuda_runtime.h" +#include "macros.hpp" #include "nvJitLink.h" #include #include -#include -#include #include #include @@ -53,7 +52,7 @@ std::shared_ptr AlgorithmPlanner::get_launcher() log_message += std::string{fragment->get_key()} + ","; } log_message.pop_back(); - RAFT_LOG_DEBUG("%s", log_message.c_str()); + // RTCX_LOG_DEBUG("%s", log_message.c_str()); auto launcher = this->build(); launchers[launch_key] = launcher; return launcher; @@ -64,9 +63,9 @@ std::shared_ptr 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 +100,15 @@ std::shared_ptr 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(kernel, library); } diff --git a/CMakeLists.txt b/CMakeLists.txt index ccdfe93..712e308 100644 --- a/CMakeLists.txt +++ b/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 $ 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) diff --git a/NVRTCLTOFragmentCompiler.cpp b/NVRTCLTOFragmentCompiler.cpp index e64f64b..06b187c 100644 --- a/NVRTCLTOFragmentCompiler.cpp +++ b/NVRTCLTOFragmentCompiler.cpp @@ -4,11 +4,10 @@ */ #include "cuda.h" +#include "macros.hpp" #include #include -#include -#include #include @@ -17,7 +16,7 @@ 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() @@ -25,9 +24,9 @@ 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 NVRTCLTOFragmentCompiler::compile(std::string NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &log_size)); std::unique_ptr 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)); diff --git a/macros.hpp b/macros.hpp new file mode 100644 index 0000000..752b1e1 --- /dev/null +++ b/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 + +#include +#include +#include +#include + +#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 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) diff --git a/nvjitlink_checker.cpp b/nvjitlink_checker.cpp index da53108..c91194d 100644 --- a/nvjitlink_checker.cpp +++ b/nvjitlink_checker.cpp @@ -3,9 +3,10 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "macros.hpp" + #include #include -#include #include #include @@ -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()); } } From 7804f6417cebad7ae0549b1618d42f175429f5eb Mon Sep 17 00:00:00 2001 From: Arha Gatram Date: Tue, 14 Jul 2026 09:18:10 -0700 Subject: [PATCH 2/2] Remove debug log --- AlgorithmPlanner.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/AlgorithmPlanner.cpp b/AlgorithmPlanner.cpp index da419c1..c6858c1 100644 --- a/AlgorithmPlanner.cpp +++ b/AlgorithmPlanner.cpp @@ -46,13 +46,6 @@ std::shared_ptr AlgorithmPlanner::get_launcher() std::unique_lock 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(); - // RTCX_LOG_DEBUG("%s", log_message.c_str()); auto launcher = this->build(); launchers[launch_key] = launcher; return launcher;