diff --git a/cpp/include/raft/core/detail/macros.hpp b/cpp/include/raft/core/detail/macros.hpp index fb0d597850..390acea697 100644 --- a/cpp/include/raft/core/detail/macros.hpp +++ b/cpp/include/raft/core/detail/macros.hpp @@ -60,6 +60,20 @@ #define RAFT_INLINE_CONDITIONAL inline #endif // RAFT_COMPILED +// The RAFT_WEAK_FUNCTION specificies that: +// +// 1. A function may be defined in multiple translation units (like inline) +// +// 2. Must still emit an external symbol (unlike inline). This enables declaring +// a function signature in an `-ext` header and defining it in a source file. +// +// From +// https://gcc.gnu.org/onlinedocs/gcc/Common-Function-Attributes.html#Common-Function-Attributes: +// +// "The weak attribute causes a declaration of an external symbol to be emitted +// as a weak symbol rather than a global." +#define RAFT_WEAK_FUNCTION __attribute__((weak)) + /** * Some macro magic to remove optional parentheses of a macro argument. * See https://stackoverflow.com/a/62984543 diff --git a/cpp/include/raft/neighbors/detail/ivf_pq_compute_similarity-ext.cuh b/cpp/include/raft/neighbors/detail/ivf_pq_compute_similarity-ext.cuh index 0d5ca90297..41e9fda701 100644 --- a/cpp/include/raft/neighbors/detail/ivf_pq_compute_similarity-ext.cuh +++ b/cpp/include/raft/neighbors/detail/ivf_pq_compute_similarity-ext.cuh @@ -17,6 +17,7 @@ #pragma once #include // __half +#include // RAFT_WEAK_FUNCTION #include // raft::distance::DistanceType #include // raft::neighbors::ivf_pq::detail::fp_8bit #include // raft::neighbors::ivf_pq::codebook_gen @@ -30,7 +31,8 @@ namespace raft::neighbors::ivf_pq::detail { // is_local_topk_feasible is not inline here, because we would have to define it // here as well. That would run the risk of the definitions here and in the // -inl.cuh header diverging. -auto is_local_topk_feasible(uint32_t k, uint32_t n_probes, uint32_t n_queries) -> bool; +auto RAFT_WEAK_FUNCTION is_local_topk_feasible(uint32_t k, uint32_t n_probes, uint32_t n_queries) + -> bool; template = 32) && !(kMaxCapacity & (kMaxCapacity - 1)), "kMaxCapacity must be a power of two, not smaller than the WarpSize."); // using weak attribute here, because it may be compiled multiple times. -auto __attribute__((weak)) is_local_topk_feasible(uint32_t k, uint32_t n_probes, uint32_t n_queries) +auto RAFT_WEAK_FUNCTION is_local_topk_feasible(uint32_t k, uint32_t n_probes, uint32_t n_queries) -> bool { if (k > kMaxCapacity) { return false; } // warp_sort not possible