diff --git a/.github/workflows/lib_qec.yaml b/.github/workflows/lib_qec.yaml index 40c29284b..81cfe3189 100644 --- a/.github/workflows/lib_qec.yaml +++ b/.github/workflows/lib_qec.yaml @@ -313,7 +313,8 @@ jobs: - name: Run multi-GPU C++ tests run: | cd $GITHUB_WORKSPACE/build_qec - # The matching tests will be added separately; until then, this - # command should not fail solely because the filter finds no tests. + # CudaDeviceId also matches the TRT 2-GPU placement tests + # (TRTDecoderTest.CudaDeviceId_*); TwoTrtDecodersConcurrently is the + # TRT concurrent-decode test that skips on single-GPU lanes. ctest --output-on-failure \ - -R "DecoderPool|HardwarePinningGpu|PoolRunsTwoTrt|CudaDeviceId" + -R "HardwarePinningGpu|CudaDeviceId|TwoTrtDecodersConcurrently" diff --git a/libs/qec/include/cudaq/qec/decoder.h b/libs/qec/include/cudaq/qec/decoder.h index 0a2f4c0c0..0662d3a24 100644 --- a/libs/qec/include/cudaq/qec/decoder.h +++ b/libs/qec/include/cudaq/qec/decoder.h @@ -13,13 +13,16 @@ #include "cuda-qx/core/tensor.h" #include "sparse_binary_matrix.h" #include "cudaq/qec/detector_error_model.h" +#include "cudaq/qec/device_affinity.h" #include +#include #include #include #include #include #include #include +#include #include #include #include @@ -173,6 +176,13 @@ class decoder /// length of the syndrome vector should be equal to `syndrome_size`. /// @returns Vector of length `block_size` with soft probabilities of errors /// in each index. + /// Note: this entry point has no automatic CUDA/NUMA guard. It is a pure + /// virtual that plugin implementations override directly, so the base + /// class cannot wrap it without an ABI-breaking vtable change. + /// decode_batch(), decode_async(), decode(tensor), and enqueue_syndrome() + /// all apply the guard around their own call into this function; a caller + /// invoking this overload directly on an unbound decoder is responsible + /// for its own placement (e.g. call bind_current_thread() first). virtual decoder_result decode(const std::vector &syndrome) = 0; /// @brief Decode a single syndrome @@ -187,6 +197,8 @@ class decoder /// value is the probability that the syndrome measurement is a |1>. /// @returns std::future of a vector of length `block_size` with soft /// probabilities of errors in each index. + /// @note The caller must ensure the decoder outlives the returned future. + /// Destroying the decoder before calling .get() is undefined behaviour. virtual std::future decode_async(const std::vector &syndrome); @@ -245,6 +257,51 @@ class decoder std::size_t get_block_size() { return block_size; } std::size_t get_syndrome_size() { return syndrome_size; } + /// @brief Store hardware affinity parameters read from the constructor params + /// map. Called by decoder::get() after the plugin constructor returns. + /// Not intended for direct use by plugin authors. + void set_hardware_params(const cudaqx::heterogeneous_map ¶ms); + + /// @brief Target NUMA node for this decoder (-1 = no binding). + int numa_node_id() const { return numa_node_id_; } + + /// @brief Target CUDA device for this decoder (-1 = inherit caller's device). + int cuda_device_id() const { return cuda_device_id_; } + + /// @brief NUMA memory-policy mode for this decoder (preferred = soft, + /// bind = strict). + cudaq::qec::mempolicy_mode mempolicy() const { return mempolicy_; } + + /// @brief Explicit CPU list for this decoder (empty = derive from NUMA node). + const std::vector &cpu_affinity() const { return cpu_affinity_; } + + /// @brief Persistently bind the CALLING thread to this decoder's NUMA node + /// (and CUDA device). Call once from the thread that will own decode()/ + /// enqueue for this decoder (e.g. a realtime worker). No restore. + /// @return the node bound to, or -1 if nothing was applied. Warns if a node + /// was requested (>= 0) but could not be honored. + int bind_current_thread(); + + /// @brief Run decode(syndrome) on a fresh worker thread bound (via + /// bind_current_thread) to this decoder's CUDA device / NUMA node, and return + /// the result. Convenience for a one-off pinned decode; for sustained + /// throughput drive decode on a long-lived bound thread instead. + decoder_result decode_on_pinned_thread(const std::vector &syndrome); + + /// @brief Single-syndrome decode with the same automatic CUDA/NUMA guard as + /// decode_batch(). Non-virtual; dispatches to the virtual decode() inside + /// the guard. Intended for host-language bindings and callers that cannot + /// use bind_current_thread(); bound threads skip the guard as usual. + decoder_result decode_guarded(const std::vector &syndrome); + + /// @brief Forget any bind_current_thread() registration on this decoder so + /// guarded entry points stop skipping their per-call guard. Call before the + /// bound thread exits (e.g. session teardown): thread ids are recycled by + /// the runtime, and a stale registration would let an unrelated new thread + /// silently skip the guard. Does not undo the OS-level placement of the + /// (exiting) bound thread. + void unbind_thread(); + // -- Begin realtime decoding API -- // Note: all of the current realtime decoding API is designed to be used with @@ -357,6 +414,27 @@ class decoder /// @brief The decoder's D matrix in sparse format std::vector> D_sparse; + /// Target CUDA device for this decoder. -1 = inherit caller's device (no-op). + int cuda_device_id_ = -1; + /// Target NUMA node for this decoder. -1 = no binding. + int numa_node_id_ = -1; + /// Thread that bind_current_thread() pinned, if any. Guarded call sites + /// (decode_batch, decode(tensor), enqueue_syndrome) skip their per-call + /// CUDA/NUMA guard only when the calling thread matches. A + /// default-constructed id means no thread is bound. + std::atomic bound_thread_{}; + /// NUMA memory-policy mode for this decoder (default soft PREFERRED). + cudaq::qec::mempolicy_mode mempolicy_ = cudaq::qec::mempolicy_mode::preferred; + /// Explicit CPU cores for this decoder's owning thread (empty = use node + /// cpuset). + std::vector cpu_affinity_; + + /// True if bind_current_thread() was called, and the CURRENT thread is the + /// one that called it -- the only thread allowed to skip the per-call + /// CUDA/NUMA guard. Protected so plugin overrides of decode_batch() can + /// apply the same skip as the base-class entry points. + bool is_bound_here() const; + private: decode_result_type result_type_ = decode_result_type::decode_to_errs; }; diff --git a/libs/qec/include/cudaq/qec/device_affinity.h b/libs/qec/include/cudaq/qec/device_affinity.h new file mode 100644 index 000000000..3cf24f841 --- /dev/null +++ b/libs/qec/include/cudaq/qec/device_affinity.h @@ -0,0 +1,88 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ + +#pragma once + +#include "cuda-qx/core/heterogeneous_map.h" +#include +#include +#include + +namespace cudaq::qec { + +namespace detail { +/// Returns -1 if key is absent; throws if stored value is negative. +/// Handles both int (YAML path) and std::size_t (Python kwargs path) storage. +inline int read_pin_key(const cudaqx::heterogeneous_map ¶ms, + const std::string &key) { + if (!params.contains(key)) + return -1; + int value = params.get(key); + if (value < 0) + throw std::runtime_error(key + " must be >= 0 (got " + + std::to_string(value) + ")"); + return value; +} +} // namespace detail + +/// @brief GPU device id for a decoder. -1 = inherit current device (no-op). +inline int read_cuda_device_id(const cudaqx::heterogeneous_map ¶ms) { + return detail::read_pin_key(params, "cuda_device_id"); +} + +/// @brief NUMA node id for a decoder. -1 = no binding. +inline int read_numa_node_id(const cudaqx::heterogeneous_map ¶ms) { + return detail::read_pin_key(params, "numa_node_id"); +} + +/// @brief Public NUMA memory-policy selector. preferred = soft; bind = strict. +enum class mempolicy_mode { preferred, bind }; + +/// @brief Read "mempolicy": "bind"->bind, "preferred"/absent->preferred, else +/// throw. +inline mempolicy_mode read_mempolicy(const cudaqx::heterogeneous_map ¶ms) { + if (!params.contains("mempolicy")) + return mempolicy_mode::preferred; + const std::string v = params.get("mempolicy"); + if (v == "bind") + return mempolicy_mode::bind; + if (v == "preferred") + return mempolicy_mode::preferred; + throw std::runtime_error( + "mempolicy must be \"preferred\" or \"bind\" (got \"" + v + "\")"); +} + +/// @brief Read "cpu_affinity": a list of CPU core ids. Absent -> empty (no +/// override). Accepts vector (C++/YAML) and vector with integral +/// values (Python kwargs path — lists arrive as doubles). +inline std::vector +read_cpu_affinity(const cudaqx::heterogeneous_map ¶ms) { + if (!params.contains("cpu_affinity")) + return {}; + try { + return params.get>("cpu_affinity"); + } catch (...) { + } + const auto vals = params.get>("cpu_affinity"); + std::vector cores; + cores.reserve(vals.size()); + for (double v : vals) { + if (v != static_cast(static_cast(v))) + throw std::runtime_error("cpu_affinity core ids must be integers (got " + + std::to_string(v) + ")"); + cores.push_back(static_cast(v)); + } + return cores; +} + +/// @brief NUMA node local to a CUDA device (via PCIe locality), or -1 if the +/// device id is negative or the topology can't be resolved. Defined in +/// decoder.cpp. +int numa_node_for_cuda_device(int cuda_device_id); + +} // namespace cudaq::qec diff --git a/libs/qec/include/cudaq/qec/realtime/decoding_config.h b/libs/qec/include/cudaq/qec/realtime/decoding_config.h index 5d1568b24..d7f02a45d 100644 --- a/libs/qec/include/cudaq/qec/realtime/decoding_config.h +++ b/libs/qec/include/cudaq/qec/realtime/decoding_config.h @@ -179,6 +179,10 @@ struct decoder_config { std::vector H_sparse; std::vector O_sparse; std::vector D_sparse; + std::optional cuda_device_id; + std::optional numa_node_id; + std::optional mempolicy; // "preferred" | "bind" + std::optional> cpu_affinity; // explicit core ids std::variant diff --git a/libs/qec/lib/decoder.cpp b/libs/qec/lib/decoder.cpp index 0419bd89f..85d5bdb62 100644 --- a/libs/qec/lib/decoder.cpp +++ b/libs/qec/lib/decoder.cpp @@ -8,14 +8,33 @@ #include "cudaq/qec/decoder.h" #include "cuda-qx/core/library_utils.h" +#include "cudaq/qec/device_affinity.h" #include "cudaq/qec/logger.h" #include "cudaq/qec/plugin_loader.h" #include "cudaq/qec/version.h" #include +#include +#include +#include #include #include #include +#include #include +#if defined(__linux__) +#include +#include +#include +#include +#include +#include +#include +#endif +#include "hardware_affinity.h" +#include "hardware_guards.h" + +using cudaq::qec::detail_affinity::CudaDeviceGuard; +using cudaq::qec::detail_affinity::NumaGuard; INSTANTIATE_REGISTRY(cudaq::qec::decoder, const cudaq::qec::decoder_init &, const cudaqx::heterogeneous_map &) @@ -87,9 +106,178 @@ decoder::decoder(cudaq::qec::sparse_binary_matrix H) pimpl->should_log = ch[0] == '1' || ch[0] == 'y' || ch[0] == 'Y'; } +int numa_node_for_cuda_device(int cuda_device_id) { + if (cuda_device_id < 0) + return -1; +#if defined(__linux__) + char busid[32] = {0}; + if (cudaDeviceGetPCIBusId(busid, sizeof(busid), cuda_device_id) != + cudaSuccess) { + cudaq::qec::detail_affinity::affinity_info( + "numa_node_id auto-derive for cuda_device_id " + + std::to_string(cuda_device_id) + + " could not read the PCI bus id; NUMA binding skipped"); + return -1; + } + for (char *c = busid; *c; ++c) + *c = static_cast(std::tolower(static_cast(*c))); + std::ifstream f(std::string("/sys/bus/pci/devices/") + busid + "/numa_node"); + int node = -1; + if (!f.is_open() || !(f >> node) || node < 0) { + cudaq::qec::detail_affinity::affinity_info( + "numa_node_id auto-derive for cuda_device_id " + + std::to_string(cuda_device_id) + + " resolved to no node (single-node host or unresolved topology); " + "NUMA binding skipped"); + return -1; + } + return node; +#else + return -1; +#endif +} + +void decoder::set_hardware_params(const cudaqx::heterogeneous_map ¶ms) { + cuda_device_id_ = cudaq::qec::read_cuda_device_id(params); + numa_node_id_ = cudaq::qec::read_numa_node_id(params); + mempolicy_ = cudaq::qec::read_mempolicy(params); + cpu_affinity_ = cudaq::qec::read_cpu_affinity(params); + // Soft-derive: user pinned a GPU but not a node -> use the GPU's node (or -1 + // if unresolved). + if (numa_node_id_ < 0 && cuda_device_id_ >= 0) + numa_node_id_ = numa_node_for_cuda_device(cuda_device_id_); +} + +int decoder::bind_current_thread() { + int prev_dev = -1; + bool dev_switched = false; + if (cuda_device_id_ >= 0) { + // Persistent set on this thread (CudaDeviceGuard restores on scope exit, so + // set directly here and let it stick). + int count = 0; + if (cudaGetDeviceCount(&count) != cudaSuccess || cuda_device_id_ >= count) + throw std::runtime_error( + "cuda_device_id " + std::to_string(cuda_device_id_) + + " out of range or CUDA unavailable in bind_current_thread"); + if (cudaGetDevice(&prev_dev) != cudaSuccess) + prev_dev = -1; + cudaError_t e = cudaSetDevice(cuda_device_id_); + if (e != cudaSuccess) + throw std::runtime_error("bind_current_thread: cudaSetDevice(" + + std::to_string(cuda_device_id_) + + ") failed: " + cudaGetErrorString(e)); + dev_switched = (prev_dev >= 0 && prev_dev != cuda_device_id_); + } + // Capture thread placement before the NUMA bind so every side effect can be + // rolled back if any later step throws and bound_thread_ is never stored. + auto prev_mempol_bind = + cudaq::qec::detail_affinity::capture_thread_mempolicy(); +#if defined(__linux__) + cpu_set_t prev_affinity; + CPU_ZERO(&prev_affinity); + const bool has_prev_affinity = + (sched_getaffinity(0, sizeof(prev_affinity), &prev_affinity) == 0); +#endif + try { + cudaq::qec::detail_affinity::bind_this_thread_to_numa_node(numa_node_id_, + mempolicy_); + if (numa_node_id_ >= 0) { + // Best-effort confirmation: warn if the calling thread is not actually + // running on the requested node's CPUs (e.g. locked cpuset in a + // container). +#if defined(__linux__) + cpu_set_t want, have; + CPU_ZERO(&want); + CPU_ZERO(&have); + if (cudaq::qec::detail_affinity::build_node_cpuset(numa_node_id_, want) && + sched_getaffinity(0, sizeof(have), &have) == 0) { + bool on_node = false; + for (int c = 0; c < CPU_SETSIZE; ++c) + if (CPU_ISSET(c, &have) && CPU_ISSET(c, &want)) { + on_node = true; + break; + } + if (!on_node) + cudaq::qec::detail_affinity::affinity_warn( + "bind_current_thread: numa_node_id " + + std::to_string(numa_node_id_) + + " requested but the calling thread could not be pinned to it"); + } +#endif + } + if (!cpu_affinity_.empty()) + cudaq::qec::detail_affinity::set_thread_cpu_affinity(cpu_affinity_); + } catch (...) { + // Roll back every side effect applied above so a failed bind leaves the + // thread exactly as it was: mempolicy, CPU affinity, and CUDA device. + cudaq::qec::detail_affinity::restore_thread_mempolicy(prev_mempol_bind); +#if defined(__linux__) + if (has_prev_affinity) + sched_setaffinity(0, sizeof(prev_affinity), &prev_affinity); +#endif + if (dev_switched && cudaSetDevice(prev_dev) != cudaSuccess) + cudaq::qec::detail_affinity::affinity_warn( + "bind_current_thread: failed to restore prior CUDA device " + + std::to_string(prev_dev) + " after a failed bind"); + throw; + } + // Only mark this thread as bound once every step that can throw has + // succeeded, so a failed bind never leaves the decoder falsely marked as + // bound to this thread. + bound_thread_.store(std::this_thread::get_id(), std::memory_order_release); + return numa_node_id_; +} + +bool decoder::is_bound_here() const { + return bound_thread_.load(std::memory_order_acquire) == + std::this_thread::get_id(); +} + +void decoder::unbind_thread() { + bound_thread_.store(std::thread::id{}, std::memory_order_release); +} + +decoder_result +decoder::decode_on_pinned_thread(const std::vector &syndrome) { + decoder_result result; + std::exception_ptr err; + // Snapshot the caller's binding before the worker overwrites it, so we can + // restore it after join. + const auto caller_id = bound_thread_.load(std::memory_order_acquire); + std::thread worker([&] { + try { + bind_current_thread(); + result = decode(syndrome); + } catch (...) { + err = std::current_exception(); + } + // Clear the one-shot worker's own binding on both success and exception + // paths so a recycled thread id never spuriously skips the guards. + auto me = std::this_thread::get_id(); + bound_thread_.compare_exchange_strong(me, std::thread::id{}, + std::memory_order_acq_rel); + }); + worker.join(); + // Restore caller_id only if the slot is still empty (worker cleared it). + // If a concurrent bind_current_thread() wrote a different id between the + // worker's CAS-clear and here, leave that binding intact. + std::thread::id empty{}; + bound_thread_.compare_exchange_strong(empty, caller_id, + std::memory_order_acq_rel); + if (err) + std::rethrow_exception(err); + return result; +} + // Provide a trivial implementation of for tensor decode call. Child // classes should override this if they never want to pass through floats. decoder_result decoder::decode(const cudaqx::tensor &syndrome) { + // Guards are constructed before ANY input processing so the soft-syndrome + // temporaries below are allocated on the decoder's device/NUMA node. They + // are exception-safe RAII, so placement is restored even on the rank throw. + const bool skip_guard = is_bound_here(); + CudaDeviceGuard dev(skip_guard ? -1 : cuda_device_id_); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); // Check tensor is of order-1 // If order >1, we could check that other modes are of dim = 1 such that // n x 1, or 1 x n tensors are still valid. @@ -103,10 +291,25 @@ decoder_result decoder::decode(const cudaqx::tensor &syndrome) { return decode(soft_syndrome); } +decoder_result decoder::decode_guarded(const std::vector &syndrome) { + const bool skip_guard = is_bound_here(); + CudaDeviceGuard dev(skip_guard ? -1 : cuda_device_id_); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); + return decode(syndrome); +} + // Provide a trivial implementation of the multi-syndrome decoder. Child classes // should override this if they can do it more efficiently than this. std::vector decoder::decode_batch(const std::vector> &syndrome) { + // Apply affinity once for the whole batch (not per-syndrome) to avoid + // repeated syscall overhead on the hot path. + // Skip the guard only on the exact thread that called bind_current_thread() + // -- a different thread must still be guarded even if this decoder was + // bound elsewhere. + const bool skip_guard = is_bound_here(); + CudaDeviceGuard dev(skip_guard ? -1 : cuda_device_id_); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); std::vector result; result.reserve(syndrome.size()); for (auto &s : syndrome) @@ -123,8 +326,21 @@ std::string decoder::get_version() const { std::future decoder::decode_async(const std::vector &syndrome) { + // The three affinity scalars are captured by value so the lambda doesn't + // dereference decoder members on the hot path. + // The syndrome copy is captured on the CALLING thread; its pages follow the + // caller's placement. The worker's guards cover everything decode-side. + // LIFETIME PRECONDITION: the decoder must outlive the returned future. + // Destroying the decoder before calling .get() is undefined behaviour. + const int cuda_id = cuda_device_id_; + const int numa_id = numa_node_id_; + const cudaq::qec::mempolicy_mode mempolicy = mempolicy_; return std::async(std::launch::async, - [this, syndrome] { return this->decode(syndrome); }); + [this, syndrome, cuda_id, numa_id, mempolicy] { + CudaDeviceGuard dev(cuda_id); + NumaGuard numa(numa_id, mempolicy); + return this->decode(syndrome); + }); } std::unique_ptr @@ -138,7 +354,37 @@ decoder::get(const std::string &name, const decoder_init &init, "invalid decoder requested: " + name + ". Run with CUDAQ_LOG_LEVEL=info (environment variable) to see " "additional plugin diagnostics at startup."); - return iter->second(init, param_map); + // Guards during construction so allocations land on the right hardware. + // Restored before this function returns; decode-time affinity is + // re-applied per call at every guarded entry point: decode_batch(), + // decode_async(), decode(const cudaqx::tensor &), and + // enqueue_syndrome(). + const int dev = cudaq::qec::read_cuda_device_id(param_map); + int node = cudaq::qec::read_numa_node_id(param_map); + if (node < 0 && dev >= 0) + node = numa_node_for_cuda_device(dev); + const auto ctor_mempolicy = cudaq::qec::read_mempolicy(param_map); + CudaDeviceGuard ctor_dev(dev); + NumaGuard ctor_numa(node, ctor_mempolicy); + // The affinity knobs above are consumed by the base class. Strip them from + // the map handed to the plugin constructor so a decoder that strictly + // validates its own parameter keys does not reject them. + auto is_affinity_key = [](const std::string &k) { + return k == "cuda_device_id" || k == "numa_node_id" || k == "mempolicy" || + k == "cpu_affinity"; + }; + cudaqx::heterogeneous_map plugin_params; + for (const auto &kv : param_map) + if (!is_affinity_key(kv.first)) + plugin_params.insert(kv.first, kv.second); + auto d = iter->second(init, plugin_params); + // Inject the pre-computed NUMA node so set_hardware_params() doesn't + // re-derive it via a second numa_node_for_cuda_device() sysfs read. + cudaqx::heterogeneous_map hw_params = param_map; + if (cudaq::qec::read_numa_node_id(param_map) < 0 && node >= 0) + hw_params.insert("numa_node_id", node); + d->set_hardware_params(hw_params); + return d; } namespace details { @@ -198,6 +444,10 @@ set_sparse_from_vec(const std::vector &vec_in, } void decoder::set_O_sparse(const std::vector> &O_sparse) { + // The corrections buffer allocated here lives for the whole session, so it + // must follow the decoder's placement (host allocation: no device guard). + const bool skip_guard = is_bound_here(); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); this->O_sparse = O_sparse; validate_sparse_column_indices(this->O_sparse, block_size, "O_sparse"); this->pimpl->corrections.clear(); @@ -205,6 +455,10 @@ void decoder::set_O_sparse(const std::vector> &O_sparse) { } void decoder::set_O_sparse(const std::vector &O_sparse_vec_in) { + // The corrections buffer allocated here lives for the whole session, so it + // must follow the decoder's placement (host allocation: no device guard). + const bool skip_guard = is_bound_here(); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); set_sparse_from_vec(O_sparse_vec_in, this->O_sparse); validate_sparse_column_indices(this->O_sparse, block_size, "O_sparse"); this->pimpl->corrections.clear(); @@ -254,11 +508,21 @@ void set_D_sparse_common(decoder *decoder, } void decoder::set_D_sparse(const std::vector> &D_sparse) { + // The msyn/detector buffers allocated here live for the whole session, so + // they must follow the decoder's placement (host allocation: no device + // guard). + const bool skip_guard = is_bound_here(); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); this->D_sparse = D_sparse; set_D_sparse_common(this, D_sparse, pimpl.get()); } void decoder::set_D_sparse(const std::vector &D_sparse_vec_in) { + // The msyn/detector buffers allocated here live for the whole session, so + // they must follow the decoder's placement (host allocation: no device + // guard). + const bool skip_guard = is_bound_here(); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); set_sparse_from_vec(D_sparse_vec_in, this->D_sparse); set_D_sparse_common(this, this->D_sparse, pimpl.get()); } @@ -356,7 +620,13 @@ bool decoder::enqueue_syndrome(const uint8_t *syndrome, // Send the data to the decoder. convert_vec_hard_to_soft(pimpl->persistent_detector_buffer, pimpl->persistent_soft_detector_buffer); - auto decoded_result = decode(pimpl->persistent_soft_detector_buffer); + decoder_result decoded_result; + { + const bool skip_guard = is_bound_here(); + CudaDeviceGuard dev(skip_guard ? -1 : cuda_device_id_); + NumaGuard numa(skip_guard ? -1 : numa_node_id_, mempolicy_); + decoded_result = decode(pimpl->persistent_soft_detector_buffer); + } // If we didn't get a decoded result, just return if (pimpl->is_sliding_window) { diff --git a/libs/qec/lib/decoders/plugins/chromobius/chromobius.cpp b/libs/qec/lib/decoders/plugins/chromobius/chromobius.cpp index 54712c57d..71222d8ca 100644 --- a/libs/qec/lib/decoders/plugins/chromobius/chromobius.cpp +++ b/libs/qec/lib/decoders/plugins/chromobius/chromobius.cpp @@ -140,6 +140,9 @@ class chromobius : public decoder { // observable-reduction logic treat each predicted bit as its own observable // correction. this->set_O_sparse(identity_sparse(num_observables)); + // Chromobius returns observable flips directly; declare this so the base + // class knows not to project through O_sparse a second time. + set_result_type(decode_to_obs); } decoder_result decode(const std::vector &syndrome) override { diff --git a/libs/qec/lib/decoders/plugins/trt_decoder/CMakeLists.txt b/libs/qec/lib/decoders/plugins/trt_decoder/CMakeLists.txt index 84a5a7b76..b34307f75 100644 --- a/libs/qec/lib/decoders/plugins/trt_decoder/CMakeLists.txt +++ b/libs/qec/lib/decoders/plugins/trt_decoder/CMakeLists.txt @@ -124,6 +124,7 @@ if(CUDAQ_QEC_TRT_DECODER_ENABLED) ${CMAKE_SOURCE_DIR}/libs/qec/include ${CMAKE_SOURCE_DIR}/libs/core/include PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/../../.. ${TENSORRT_INCLUDE_DIR} ${CUDAToolkit_INCLUDE_DIRS} ) diff --git a/libs/qec/lib/decoders/plugins/trt_decoder/trt_decoder.cpp b/libs/qec/lib/decoders/plugins/trt_decoder/trt_decoder.cpp index 208a185c5..c1c5c5414 100644 --- a/libs/qec/lib/decoders/plugins/trt_decoder/trt_decoder.cpp +++ b/libs/qec/lib/decoders/plugins/trt_decoder/trt_decoder.cpp @@ -19,6 +19,11 @@ #include #include #include +#if defined(__linux__) +#include +#endif +#include "hardware_affinity.h" +#include "hardware_guards.h" // TensorRT headers #include "NvInfer.h" @@ -850,17 +855,26 @@ decoder_result trt_decoder::decode(const std::vector &syndrome) { "Model has batch_size={}, zero-padding single syndrome to fill batch", model_batch_size_); - // Create a batch with the real syndrome plus zero-padded syndromes + // Create a batch with the real syndrome plus zero-padded syndromes. + // Declared before the guard scope (the empty vector allocates nothing); + // every allocation happens inside the scope below. std::vector> padded_batch; - padded_batch.reserve(model_batch_size_); - - // First syndrome is the real one - padded_batch.push_back(syndrome); - - // Fill remaining batch slots with zero syndromes - std::vector zero_syndrome(syndrome_size_per_sample_, 0.0f); - for (size_t i = 1; i < model_batch_size_; ++i) { - padded_batch.push_back(zero_syndrome); + { + // The padding buffers must land on the decoder's NUMA node like every + // other decode-side allocation. The guard scope closes before the + // delegation because decode_batch() re-applies its own guards. + cudaq::qec::detail_affinity::NumaGuard place( + is_bound_here() ? -1 : numa_node_id_, mempolicy()); + padded_batch.reserve(model_batch_size_); + + // First syndrome is the real one + padded_batch.push_back(syndrome); + + // Fill remaining batch slots with zero syndromes + std::vector zero_syndrome(syndrome_size_per_sample_, 0.0f); + for (size_t i = 1; i < model_batch_size_; ++i) { + padded_batch.push_back(zero_syndrome); + } } auto results = decode_batch(padded_batch); @@ -888,6 +902,19 @@ decoder_result trt_decoder::decode(const std::vector &syndrome) { std::vector trt_decoder::decode_batch(const std::vector> &syndromes) { + // Skip device/NUMA guards when the calling thread already called + // bind_current_thread() — mirrors the base class decode_batch() behaviour. + const bool skip_guard = is_bound_here(); + + // This override bypasses decoder::decode_batch()'s guards, so it applies + // the same shared pair as the base-class entry points (lib-private header, + // reachable across the plugin boundary): device switch + NUMA bind, both + // restored on scope exit. Out-of-range ids throw; an unreadable current + // device warns and skips the switch. + namespace da = cudaq::qec::detail_affinity; + da::CudaDeviceGuard dev_guard(skip_guard ? -1 : cuda_device_id_); + da::NumaGuard numa_guard(skip_guard ? -1 : numa_node_id_, mempolicy()); + // Validate that we have syndromes to decode if (syndromes.empty()) { return {}; diff --git a/libs/qec/lib/hardware_affinity.h b/libs/qec/lib/hardware_affinity.h new file mode 100644 index 000000000..ad5fa1842 --- /dev/null +++ b/libs/qec/lib/hardware_affinity.h @@ -0,0 +1,259 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ + +// Lib-private (NOT installed): hardware affinity primitives shared by the base +// decoder and the realtime session. Never include from a public header. +#pragma once + +#include "cudaq/qec/device_affinity.h" +#include +#include +#include +#include +#include +#include +#include + +#if defined(__linux__) +#include +#include +#include +#include +#include +#include +#include +#include +#endif + +namespace cudaq::qec::detail_affinity { + +// Decoupled diagnostics: plain C, no cudaq dependency, so this header can be +// reused independently. +inline void affinity_warn(const std::string &msg) { + std::fprintf(stderr, "[cudaq-qec affinity] WARNING: %s\n", msg.c_str()); +} +// INFO/debug: only shown when CUDAQ_QEC_AFFINITY_DEBUG is set (a legitimate +// no-op like a single-node auto-derive must not spam; it is debug-loggable, not +// a failure). +inline void affinity_info(const std::string &msg) { + if (std::getenv("CUDAQ_QEC_AFFINITY_DEBUG") != nullptr) + std::fprintf(stderr, "[cudaq-qec affinity] INFO: %s\n", msg.c_str()); +} + +#if defined(__linux__) + +inline int mempolicy_syscall_mode(cudaq::qec::mempolicy_mode m) { + return (m == cudaq::qec::mempolicy_mode::bind) ? MPOL_BIND : MPOL_PREFERRED; +} + +// Parse /sys/devices/system/node/node/cpulist e.g. "0-7,16-23". +// Returns false if the file is missing or empty. +inline bool build_node_cpuset(int node, cpu_set_t &out) { + std::ifstream f("/sys/devices/system/node/node" + std::to_string(node) + + "/cpulist"); + if (!f.is_open()) + return false; + std::string list; + std::getline(f, list); + if (list.empty()) + return false; + std::stringstream ss(list); + std::string range; + bool any = false; + while (std::getline(ss, range, ',')) { + auto dash = range.find('-'); + int lo = std::stoi(range.substr(0, dash)); + int hi = + (dash == std::string::npos) ? lo : std::stoi(range.substr(dash + 1)); + for (int c = lo; c <= hi; ++c) { + if (c < CPU_SETSIZE) { // guard against OOB on >1024-CPU machines + CPU_SET(c, &out); + any = true; + } + } + } + return any; +} + +// Persistently bind the CALLING thread's CPU affinity + memory policy to a NUMA +// node. No restore. node < 0 = no-op. Only pins CPU affinity when prior +// affinity is readable (avoids permanent pinning in a locked-cpuset container). +// apply_mempolicy=false skips the set_mempolicy half; temporary guards use it +// when the prior policy could not be captured, so they never apply a policy +// they cannot restore (the same only-if-restorable rule as the affinity half). +// Throws if node cannot be encoded in the mempolicy nodemask (node >= 64). +// Warns (does not throw) if the OS declines an otherwise well-formed request. +inline void bind_this_thread_to_numa_node( + int node, + cudaq::qec::mempolicy_mode mode = cudaq::qec::mempolicy_mode::preferred, + bool apply_mempolicy = true) { + if (node < 0) + return; + if (node >= static_cast(sizeof(unsigned long) * 8)) + throw std::runtime_error("numa_node_id " + std::to_string(node) + + " exceeds the maximum encodable NUMA node (" + + std::to_string(sizeof(unsigned long) * 8 - 1) + + "); cannot bind memory policy"); + cpu_set_t prev; + CPU_ZERO(&prev); + const bool can_restore = (sched_getaffinity(0, sizeof(prev), &prev) == 0); + cpu_set_t node_set; + CPU_ZERO(&node_set); + if (can_restore && build_node_cpuset(node, node_set)) { + if (sched_setaffinity(0, sizeof(node_set), &node_set) != 0) + affinity_warn( + "numa_node_id " + std::to_string(node) + + " requested but sched_setaffinity could not pin the thread: " + + std::strerror(errno) + "; running unpinned (locked cpuset?)"); + } else { + affinity_warn("numa_node_id " + std::to_string(node) + + " requested but its CPU list could not be read (or prior " + "affinity unreadable); CPU affinity left unpinned"); + } + if (!apply_mempolicy) + return; + unsigned long nodemask = 1UL << node; + if (syscall(SYS_set_mempolicy, mempolicy_syscall_mode(mode), &nodemask, + static_cast(sizeof(nodemask) * 8)) != 0) + affinity_warn("numa_node_id " + std::to_string(node) + + " requested but set_mempolicy failed: " + + std::strerror(errno) + "; memory not bound to node"); +} + +// Migrate an already-allocated region onto a NUMA node. The ring buffers are +// calloc'd (already faulted) on the setup thread, so MPOL_MF_MOVE is required +// to relocate the pages. node < 0 / null / 0 bytes = no-op. +// Throws if node cannot be encoded in the mempolicy nodemask (node >= 64). +// Warns (does not throw) if the OS declines an otherwise well-formed request. +inline void bind_region_to_numa_node( + void *p, std::size_t bytes, int node, + cudaq::qec::mempolicy_mode mode = cudaq::qec::mempolicy_mode::preferred) { + if (node < 0 || p == nullptr || bytes == 0) + return; + if (node >= static_cast(sizeof(unsigned long) * 8)) + throw std::runtime_error( + "numa_node_id " + std::to_string(node) + + " exceeds the maximum encodable NUMA node; cannot migrate region"); + unsigned long nodemask = 1UL << node; + if (syscall(SYS_mbind, p, static_cast(bytes), + mempolicy_syscall_mode(mode), &nodemask, + static_cast(sizeof(nodemask) * 8), + MPOL_MF_MOVE) != 0) + affinity_warn("mbind of region to numa_node_id " + std::to_string(node) + + " failed: " + std::strerror(errno) + + "; pages not migrated (missing CAP_SYS_NICE?)"); +} + +// Query the calling thread's current memory-policy mode (MPOL_* constant). +// Returns -1 if the query fails. +inline int current_thread_mempolicy_mode() { + int mode = -1; + if (syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) != 0) + return -1; + return mode; +} + +// Full state needed to restore a thread's memory policy exactly as it was. +// mode == -1 means the capture failed (nothing to restore). +struct mempolicy_state { + int mode = -1; + unsigned long nodemask[16] = {0}; // up to 1024 nodes +}; + +// Capture the calling thread's current memory policy (mode + nodemask) so it +// can be restored exactly later. +inline mempolicy_state capture_thread_mempolicy() { + mempolicy_state s; + if (syscall(SYS_get_mempolicy, &s.mode, s.nodemask, + static_cast(sizeof(s.nodemask) * 8), nullptr, + 0UL) != 0) + s.mode = -1; + return s; +} + +// Restore a previously-captured memory policy exactly. No-op if the capture +// failed. Warns (does not throw) if the OS declines. +inline void restore_thread_mempolicy(const mempolicy_state &s) { + if (s.mode < 0) + return; + if (syscall(SYS_set_mempolicy, s.mode, s.nodemask, + static_cast(sizeof(s.nodemask) * 8)) != 0) + affinity_warn("failed to restore prior thread mempolicy: " + + std::string(std::strerror(errno)) + + "; thread may remain bound to its temporary policy"); +} + +// Pin the CALLING thread's CPU affinity to an explicit list of core ids. +// No restore. cores.empty() = no-op. Throws std::invalid_argument if any core +// id is out of range. Warns (does not throw) if the OS declines an otherwise +// well-formed request. +inline void set_thread_cpu_affinity(const std::vector &cores) { + if (cores.empty()) + return; + for (int c : cores) + if (c < 0 || c >= CPU_SETSIZE) + throw std::invalid_argument("cpu_affinity core id " + std::to_string(c) + + " is out of range [0, " + + std::to_string(CPU_SETSIZE) + ")"); + cpu_set_t set; + CPU_ZERO(&set); + for (int c : cores) + CPU_SET(c, &set); + if (sched_setaffinity(0, sizeof(set), &set) != 0) + affinity_warn("cpu_affinity requested but sched_setaffinity failed: " + + std::string(std::strerror(errno)) + + "; thread left unpinned (locked cpuset?)"); +} + +// Query the calling thread's current CPU affinity set as a sorted list of core +// ids. Returns an empty vector if the query fails. +inline std::vector current_thread_cpuset() { + std::vector out; + cpu_set_t set; + CPU_ZERO(&set); + if (sched_getaffinity(0, sizeof(set), &set) != 0) + return out; + for (int c = 0; c < CPU_SETSIZE; ++c) + if (CPU_ISSET(c, &set)) + out.push_back(c); + return out; +} + +#else // non-Linux: no-ops + +inline void warn_numa_unsupported_once() { + static bool warned = false; + if (!warned) { + affinity_warn( + "numa_node_id ignored: NUMA binding is only supported on Linux"); + warned = true; + } +} + +inline void bind_this_thread_to_numa_node( + int node, + cudaq::qec::mempolicy_mode = cudaq::qec::mempolicy_mode::preferred, + bool /*apply_mempolicy*/ = true) { + if (node < 0) + return; + warn_numa_unsupported_once(); +} +inline void bind_region_to_numa_node( + void *, std::size_t, int, + cudaq::qec::mempolicy_mode = cudaq::qec::mempolicy_mode::preferred) {} +inline int current_thread_mempolicy_mode() { return -1; } +inline void set_thread_cpu_affinity(const std::vector &) {} +inline std::vector current_thread_cpuset() { return {}; } + +struct mempolicy_state {}; +inline mempolicy_state capture_thread_mempolicy() { return {}; } +inline void restore_thread_mempolicy(const mempolicy_state &) {} + +#endif + +} // namespace cudaq::qec::detail_affinity diff --git a/libs/qec/lib/hardware_guards.h b/libs/qec/lib/hardware_guards.h new file mode 100644 index 000000000..9fb0406b9 --- /dev/null +++ b/libs/qec/lib/hardware_guards.h @@ -0,0 +1,128 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ + +// Lib-private (NOT installed): RAII guards that place the calling thread on a +// decoder's CUDA device / NUMA node and restore on scope exit. Shared by the +// base decoder entry points and the realtime session. Never include from a +// public header. +#pragma once + +#include "hardware_affinity.h" +#include + +namespace cudaq::qec::detail_affinity { + +// RAII: sets the calling thread's CUDA current device, restores on destruction. +// target < 0 = no-op. Throws std::runtime_error if target is out of range of +// the visible device count or cudaSetDevice() fails while switching to target. +// If cudaGetDevice() cannot read the current device, warns and skips the +// switch (the decode proceeds on the caller's current device). +struct CudaDeviceGuard { + int prev_ = -1; + bool active_ = false; + + explicit CudaDeviceGuard(int target) { + if (target < 0) + return; + int count = 0; + if (cudaGetDeviceCount(&count) != cudaSuccess || target >= count) + throw std::runtime_error( + "cuda_device_id " + std::to_string(target) + + " out of range (device_count=" + std::to_string(count) + ")"); + if (cudaGetDevice(&prev_) != cudaSuccess) { + // Can't determine the current device; warn and skip the switch so the + // decode proceeds on whatever device the caller is already on. + cudaq::qec::detail_affinity::affinity_warn( + "CudaDeviceGuard: cudaGetDevice failed for cuda_device_id " + + std::to_string(target) + "; skipping device switch"); + return; + } + if (prev_ != target) { + cudaError_t e = cudaSetDevice(target); + if (e != cudaSuccess) + throw std::runtime_error("CudaDeviceGuard: cudaSetDevice(" + + std::to_string(target) + + ") failed: " + cudaGetErrorString(e)); + active_ = true; + } + } + ~CudaDeviceGuard() { + if (!active_) + return; + if (cudaSetDevice(prev_) != cudaSuccess) + cudaq::qec::detail_affinity::affinity_warn( + "CudaDeviceGuard: failed to restore prior CUDA device " + + std::to_string(prev_) + "; thread may remain on wrong device"); + } + CudaDeviceGuard(const CudaDeviceGuard &) = delete; + CudaDeviceGuard &operator=(const CudaDeviceGuard &) = delete; +}; + +#if defined(__linux__) +// RAII: binds the calling thread to a NUMA node and restores on destruction. +// node < 0 = no-op. Uses the shared persistent-bind primitive for the set half +// and remembers prior affinity for the restore. +struct NumaGuard { + bool affinity_set_ = false; + bool has_prev_affinity_ = false; + bool mempol_set_ = false; + cpu_set_t prev_set_{}; + cudaq::qec::detail_affinity::mempolicy_state prev_mempolicy_; + + explicit NumaGuard(int node, cudaq::qec::mempolicy_mode mode = + cudaq::qec::mempolicy_mode::preferred) { + if (node < 0) + return; + CPU_ZERO(&prev_set_); + has_prev_affinity_ = + (sched_getaffinity(0, sizeof(prev_set_), &prev_set_) == 0); + if (node < static_cast(sizeof(unsigned long) * 8)) { + prev_mempolicy_ = cudaq::qec::detail_affinity::capture_thread_mempolicy(); + mempol_set_ = (prev_mempolicy_.mode >= 0); + if (!mempol_set_) + cudaq::qec::detail_affinity::affinity_warn( + "NumaGuard: prior thread mempolicy unreadable (get_mempolicy " + "failed); temporary NUMA memory policy skipped for this decode"); + } + // A temporary guard must never apply a policy it cannot restore: skip the + // mempolicy half when the capture failed (affinity is still applied, its + // restore path is independent). + cudaq::qec::detail_affinity::bind_this_thread_to_numa_node( + node, mode, /*apply_mempolicy=*/mempol_set_); + // Arm the affinity restore only after a successful bind; if bind throws the + // affinity was not changed and there is nothing to undo. + affinity_set_ = has_prev_affinity_; + } + + ~NumaGuard() { + if (mempol_set_) + cudaq::qec::detail_affinity::restore_thread_mempolicy(prev_mempolicy_); + // has_prev_affinity_ is implied: it is set iff affinity_set_ is set. + if (affinity_set_) + if (sched_setaffinity(0, sizeof(prev_set_), &prev_set_) != 0) + cudaq::qec::detail_affinity::affinity_warn( + "NumaGuard restore: failed to reset thread affinity: " + + std::string(std::strerror(errno)) + "; thread may remain pinned"); + } + + NumaGuard(const NumaGuard &) = delete; + NumaGuard &operator=(const NumaGuard &) = delete; +}; +#else +struct NumaGuard { + explicit NumaGuard(int node, cudaq::qec::mempolicy_mode mode = + cudaq::qec::mempolicy_mode::preferred) { + (void)mode; + if (node < 0) + return; + cudaq::qec::detail_affinity::warn_numa_unsupported_once(); + } +}; +#endif + +} // namespace cudaq::qec::detail_affinity diff --git a/libs/qec/lib/realtime/config.cpp b/libs/qec/lib/realtime/config.cpp index 3042b5d5f..bd5018204 100644 --- a/libs/qec/lib/realtime/config.cpp +++ b/libs/qec/lib/realtime/config.cpp @@ -709,6 +709,10 @@ struct MappingTraits { io.mapRequired("H_sparse", config.H_sparse); io.mapRequired("O_sparse", config.O_sparse); io.mapRequired("D_sparse", config.D_sparse); + io.mapOptional("cuda_device_id", config.cuda_device_id); + io.mapOptional("numa_node_id", config.numa_node_id); + io.mapOptional("mempolicy", config.mempolicy); + io.mapOptional("cpu_affinity", config.cpu_affinity); // Validate that the number of rows in the H_sparse vector is equal to // syndrome_size. diff --git a/libs/qec/lib/realtime/qec_realtime_session.cpp b/libs/qec/lib/realtime/qec_realtime_session.cpp index 80c9d62b0..f9ca55a65 100644 --- a/libs/qec/lib/realtime/qec_realtime_session.cpp +++ b/libs/qec/lib/realtime/qec_realtime_session.cpp @@ -10,6 +10,10 @@ #include "qec_realtime_session.h" +// Lib-private headers one level up (libs/qec/lib); this file lives in +// realtime/. +#include "../hardware_affinity.h" +#include "../hardware_guards.h" #include "cudaq/qec/logger.h" #include "cudaq/qec/realtime/decoder_rpc_ids.h" #include "cudaq/qec/realtime/graph_resources.h" @@ -20,8 +24,10 @@ #include #include #include +#include #include #include +#include namespace cudaq::qec::realtime { @@ -326,10 +332,107 @@ void qec_realtime_session::initialize() { // finalize() is null-safe at every step, so we can roll a half-built session // back from any throw. try { + // Resolve the session's placement consensus (NUMA node, CUDA device, CPU + // list, mempolicy) from the decoders before the CUDA resource setup + // below: graph capture must run on the agreed device. A single shared + // dispatch thread can honor only one of each; warn and disable pinning + // on disagreement. + { + int chosen = -1; + bool conflict = false; + int chosen_dev = -1; + bool dev_conflict = false; + std::vector chosen_cpus; + bool affinity_conflict = false; + bool have_mempolicy = false; + bool mempolicy_conflict = false; + cudaq::qec::mempolicy_mode chosen_mempolicy = + cudaq::qec::mempolicy_mode::preferred; + for (const auto &d : decoders_) { + if (!d) + continue; + const int n = d->numa_node_id(); + if (n >= 0) { + if (chosen < 0) + chosen = n; + else if (chosen != n) + conflict = true; + } + const int dev = d->cuda_device_id(); + if (dev >= 0) { + if (chosen_dev < 0) + chosen_dev = dev; + else if (chosen_dev != dev) + dev_conflict = true; + } + // bind_current_thread() also applies each decoder's explicit + // cpu_affinity list persistently; disagreeing non-empty lists would + // silently resolve to last-bind-wins. Empty lists never disagree. + const std::vector &cpus = d->cpu_affinity(); + if (!cpus.empty()) { + if (chosen_cpus.empty()) + chosen_cpus = cpus; + else if (chosen_cpus != cpus) + affinity_conflict = true; + } + // mempolicy participates only when the decoder set a placement knob; + // the default (preferred) never conflicts with itself. + if (n >= 0 || dev >= 0 || !cpus.empty()) { + if (!have_mempolicy) { + chosen_mempolicy = d->mempolicy(); + have_mempolicy = true; + } else if (chosen_mempolicy != d->mempolicy()) { + mempolicy_conflict = true; + } + } + } + if (conflict) { + CUDA_QEC_WARN( + "realtime decoders request different numa_node_ids; the shared " + "host dispatch thread can honor only one. NUMA pinning " + "disabled for this session."); + session_numa_node_ = -1; + } else { + session_numa_node_ = chosen; + } + // HOST mode only: DEVICE mode throws on dev_conflict below, and that + // message supersedes this one (per-call guards never get to run there). + if (dev_conflict && !device_mode_) + CUDA_QEC_WARN( + "realtime decoders request different cuda_device_ids; the shared " + "dispatch thread cannot hold one device for all of them. " + "Guard-skip binding disabled; per-call guards stay active."); + if (affinity_conflict) + CUDA_QEC_WARN( + "realtime decoders request different cpu_affinity lists; the " + "shared dispatch thread can honor only one. Guard-skip binding " + "disabled; per-call guards stay active."); + if (mempolicy_conflict) + CUDA_QEC_WARN( + "realtime decoders request different mempolicy modes; the shared " + "dispatch thread can honor only one. Guard-skip binding disabled; " + "per-call guards stay active."); + bind_decoders_to_host_loop_ = + !dev_conflict && !conflict && !affinity_conflict && + !mempolicy_conflict && + (chosen >= 0 || chosen_dev >= 0 || !chosen_cpus.empty()); + session_cuda_device_ = dev_conflict ? -1 : chosen_dev; + if (device_mode_ && dev_conflict) + throw std::runtime_error( + "qec_realtime_session::initialize: decoders request different " + "cuda_device_ids; DEVICE mode runs one dispatch kernel on one " + "device and cannot honor them. Use one device per session."); + } if (device_mode_) capture_decoder_graphs(); - allocate_ring_buffer(); - populate_function_table(); + { + // Ring buffers, function tables, and control words are polled every + // dispatch iteration; first-touch them on the session node. Scoped: + // restores the config thread's placement before the loops start. + cudaq::qec::detail_affinity::NumaGuard place(session_numa_node_); + allocate_ring_buffer(); + populate_function_table(); + } if (device_mode_) start_device_loop(); start_host_loop(); @@ -451,6 +554,11 @@ void qec_realtime_session::finalize() { //============================================================================== void qec_realtime_session::capture_decoder_graphs() { + // Graph capture allocates streams/exec graphs on the CURRENT device; place + // it on the session's agreed device so a pinned decoder's graph does not + // land on the ambient device of whichever thread called initialize(). + cudaq::qec::detail_affinity::CudaDeviceGuard place(session_cuda_device_); + captured_graphs_.assign(decoders_.size(), nullptr); num_decoders_with_graph_ = 0; @@ -589,29 +697,41 @@ void qec_realtime_session::allocate_ring_buffer() { shutdown_flag_dev_ = static_cast(d); } } else { - // HOST mode: plain host memory; the device-visible pointers alias the host - // backings (no GPU required at runtime). The host loop reads only the - // *_host views; the producer's address-as-flag publish uses rx_data_dev() - // (== rx_data_host_ here), which the host loop dereferences as host memory. - auto alloc_u64 = [&](volatile std::uint64_t *&host, - volatile std::uint64_t *&dev, const char *what) { - void *p = std::calloc(num_slots_, sizeof(std::uint64_t)); + // HOST mode: plain page-aligned host memory; the device-visible pointers + // alias the host backings (no GPU required at runtime). Page alignment is + // required by the mbind(2) migration below — a calloc pointer is not + // page-aligned and made the migration fail with EINVAL. + const std::size_t page = static_cast(sysconf(_SC_PAGESIZE)); + // calloc (the previous allocator here) detected multiplication overflow; + // aligned_alloc does not, so guard the num_slots_ * size products. + auto checked_mul = [](std::size_t a, std::size_t b) -> std::size_t { + if (b != 0 && a > std::numeric_limits::max() / b) + throw std::runtime_error( + "qec_realtime_session::initialize: ring size overflow"); + return a * b; + }; + auto alloc_zeroed_pages = [&](std::size_t bytes, + const char *what) -> void * { + const std::size_t rounded = (bytes + page - 1) & ~(page - 1); + void *p = std::aligned_alloc(page, rounded); if (!p) throw std::runtime_error( std::string( "qec_realtime_session::initialize: failed to allocate ") + what); + std::memset(p, 0, rounded); + return p; + }; + auto alloc_u64 = [&](volatile std::uint64_t *&host, + volatile std::uint64_t *&dev, const char *what) { + void *p = alloc_zeroed_pages( + checked_mul(num_slots_, sizeof(std::uint64_t)), what); host = static_cast(p); dev = host; }; auto alloc_u8 = [&](std::uint8_t *&host, std::uint8_t *&dev, const char *what) { - void *p = std::calloc(num_slots_, slot_size_); - if (!p) - throw std::runtime_error( - std::string( - "qec_realtime_session::initialize: failed to allocate ") + - what); + void *p = alloc_zeroed_pages(checked_mul(num_slots_, slot_size_), what); host = static_cast(p); dev = host; }; @@ -619,6 +739,23 @@ void qec_realtime_session::allocate_ring_buffer() { alloc_u64(tx_flags_host_, tx_flags_dev_, "tx_flags"); alloc_u8(rx_data_host_, rx_data_dev_, "RX ring data"); alloc_u8(tx_data_host_, tx_data_dev_, "TX ring data"); + // Migrate the syndrome ring buffers onto the session node so the (pinned) + // dispatch thread reads them on-node. The pages are zero-filled (faulted) + // at allocation, so MPOL_MF_MOVE relocates them; unfaulted pages follow + // the VMA policy set by mbind. + if (session_numa_node_ >= 0) { + using cudaq::qec::detail_affinity::bind_region_to_numa_node; + bind_region_to_numa_node((void *)rx_flags_host_, + num_slots_ * sizeof(std::uint64_t), + session_numa_node_); + bind_region_to_numa_node((void *)tx_flags_host_, + num_slots_ * sizeof(std::uint64_t), + session_numa_node_); + bind_region_to_numa_node((void *)rx_data_host_, num_slots_ * slot_size_, + session_numa_node_); + bind_region_to_numa_node((void *)tx_data_host_, num_slots_ * slot_size_, + session_numa_node_); + } } std::memset(&ringbuffer_, 0, sizeof(ringbuffer_)); @@ -752,13 +889,18 @@ void qec_realtime_session::populate_function_table() { //============================================================================== void qec_realtime_session::start_device_loop() { + // Dispatcher setup below (including the device_stats_dev_ cudaMalloc) runs + // on the CURRENT device; place it on the session's agreed device so the + // persistent dispatch kernel's resources live with the captured graphs. + cudaq::qec::detail_affinity::CudaDeviceGuard place(session_cuda_device_); + if (cudaq_dispatch_manager_create(&device_manager_) != CUDAQ_OK) throw std::runtime_error( "qec_realtime_session::initialize: cudaq_dispatch_manager_create " "failed"); cudaq_dispatcher_config_t dev_config{}; - dev_config.device_id = 0; + dev_config.device_id = session_cuda_device_ >= 0 ? session_cuda_device_ : 0; dev_config.num_blocks = 1; dev_config.threads_per_block = 64; dev_config.num_slots = static_cast(num_slots_); @@ -843,8 +985,40 @@ void qec_realtime_session::start_host_loop() { host_ctx_.skip_stream_sweep = true; shutdown_flag_ = 0; - host_loop_thread_ = - std::thread([this]() { cudaq_host_dispatcher_loop(&host_ctx_); }); + const int node = session_numa_node_; + const bool bind_decoders = bind_decoders_to_host_loop_; + host_loop_thread_ = std::thread([this, node, bind_decoders]() { + // Persistent bind via decoder::bind_current_thread() so that + // is_bound_here() returns true and per-call NUMA/CUDA guards are + // suppressed in the decode hot path. Only when initialize() proved all + // decoders agree on all placement knobs (device, node, cpu list, + // mempolicy); agreement on any knob engages the bind. Otherwise keep + // per-call guards active and give the thread plain NUMA locality at + // most. + if (bind_decoders) { + for (auto &d : decoders_) { + if (!d) + continue; + try { + d->bind_current_thread(); + } catch (const std::exception &e) { + cudaq::qec::detail_affinity::affinity_warn( + "host loop thread bind failed for decoder " + + std::to_string(d->get_decoder_id()) + ": " + + std::string(e.what()) + " — continuing without guard-skip"); + } + } + } else if (node >= 0) { + try { + cudaq::qec::detail_affinity::bind_this_thread_to_numa_node(node); + } catch (const std::exception &e) { + cudaq::qec::detail_affinity::affinity_warn( + "host loop thread NUMA bind failed: " + std::string(e.what()) + + " — continuing without affinity"); + } + } + cudaq_host_dispatcher_loop(&host_ctx_); + }); return; } @@ -869,60 +1043,72 @@ void qec_realtime_session::start_host_loop() { void **mailbox_bank = nullptr; - std::size_t slot = 0; - for (std::size_t i = 0; i < decoders_.size(); ++i) { - if (!captured_graphs_[i]) - continue; - auto *gres = static_cast( - captured_graphs_[i]); - - cudaStream_t stream = nullptr; - if (cudaStreamCreate(&stream) != cudaSuccess) - throw std::runtime_error( - "qec_realtime_session::initialize: cudaStreamCreate for HOST_LOOP " - "worker " + - std::to_string(slot) + " (decoder_id=" + std::to_string(i) + - ") failed"); - host_worker_streams_[i] = stream; - - auto &w = host_workers_[slot]; - w.graph_exec = gres->graph_exec; - w.stream = stream; - w.function_id = cudaq::qec::decoding::rpc::kEnqueueSyndromesFunctionId; - w.routing_key = static_cast(i); - w.pre_launch_fn = nullptr; - w.pre_launch_data = nullptr; - w.post_launch_fn = nullptr; - w.post_launch_data = nullptr; - - if (slot == 0) - mailbox_bank = gres->h_mailbox; - ++slot; - } - - host_idle_mask_storage_ = new std::uint64_t( - num_decoders_with_graph_ < 64 - ? ((std::uint64_t{1} << num_decoders_with_graph_) - 1) - : ~std::uint64_t{0}); - host_live_dispatched_storage_ = new std::uint64_t(0); - host_inflight_slot_tags_ = new int[num_decoders_with_graph_]; - for (std::size_t i = 0; i < num_decoders_with_graph_; ++i) - host_inflight_slot_tags_[i] = -1; - - // Per-worker GraphIOContext array (pinned-mapped so both CPU monitor and GPU - // graph see the same backing). { - void *h = nullptr; - void *d = nullptr; - const std::size_t bytes = - num_decoders_with_graph_ * sizeof(cudaq::realtime::GraphIOContext); - if (!allocate_pinned_mapped(bytes, &h, &d)) - throw std::runtime_error("qec_realtime_session::start_host_loop: failed " - "to allocate per-worker " - "GraphIOContext array"); - std::memset(h, 0, bytes); - io_ctxs_host_ = static_cast(h); - io_ctxs_dev_ = static_cast(d); + // Worker streams and the pinned-mapped GraphIOContext array are created + // on the CURRENT device; place them on the session's agreed device so + // they live where the captured graphs run (no-op when no decoder pinned + // a device). The host-visible halves (GraphIOContext array, control + // words) are first-touched here, so also place them on the session node. + // The host_loop_thread_ spawned below stays outside this guard: it binds + // itself. + cudaq::qec::detail_affinity::CudaDeviceGuard place(session_cuda_device_); + cudaq::qec::detail_affinity::NumaGuard numa_place(session_numa_node_); + + std::size_t slot = 0; + for (std::size_t i = 0; i < decoders_.size(); ++i) { + if (!captured_graphs_[i]) + continue; + auto *gres = static_cast( + captured_graphs_[i]); + + cudaStream_t stream = nullptr; + if (cudaStreamCreate(&stream) != cudaSuccess) + throw std::runtime_error( + "qec_realtime_session::initialize: cudaStreamCreate for HOST_LOOP " + "worker " + + std::to_string(slot) + " (decoder_id=" + std::to_string(i) + + ") failed"); + host_worker_streams_[i] = stream; + + auto &w = host_workers_[slot]; + w.graph_exec = gres->graph_exec; + w.stream = stream; + w.function_id = cudaq::qec::decoding::rpc::kEnqueueSyndromesFunctionId; + w.routing_key = static_cast(i); + w.pre_launch_fn = nullptr; + w.pre_launch_data = nullptr; + w.post_launch_fn = nullptr; + w.post_launch_data = nullptr; + + if (slot == 0) + mailbox_bank = gres->h_mailbox; + ++slot; + } + + host_idle_mask_storage_ = new std::uint64_t( + num_decoders_with_graph_ < 64 + ? ((std::uint64_t{1} << num_decoders_with_graph_) - 1) + : ~std::uint64_t{0}); + host_live_dispatched_storage_ = new std::uint64_t(0); + host_inflight_slot_tags_ = new int[num_decoders_with_graph_]; + for (std::size_t i = 0; i < num_decoders_with_graph_; ++i) + host_inflight_slot_tags_[i] = -1; + + // Per-worker GraphIOContext array (pinned-mapped so both CPU monitor and + // GPU graph see the same backing). + { + void *h = nullptr; + void *d = nullptr; + const std::size_t bytes = + num_decoders_with_graph_ * sizeof(cudaq::realtime::GraphIOContext); + if (!allocate_pinned_mapped(bytes, &h, &d)) + throw std::runtime_error( + "qec_realtime_session::start_host_loop: failed to allocate " + "per-worker GraphIOContext array"); + std::memset(h, 0, bytes); + io_ctxs_host_ = static_cast(h); + io_ctxs_dev_ = static_cast(d); + } } std::memset(&host_ctx_, 0, sizeof(host_ctx_)); @@ -951,8 +1137,41 @@ void qec_realtime_session::start_host_loop() { host_ctx_.io_ctxs_dev = io_ctxs_dev_; host_ctx_.skip_stream_sweep = false; - host_loop_thread_ = - std::thread([this]() { cudaq_host_dispatcher_loop(&host_ctx_); }); + { + const int node = session_numa_node_; + const bool bind_decoders = bind_decoders_to_host_loop_; + host_loop_thread_ = std::thread([this, node, bind_decoders]() { + // Mirror HOST-mode: bind via decoder::bind_current_thread() so that + // guard-skip is activated for the decode hot path. Only when + // initialize() proved all decoders agree on all placement knobs + // (device, node, cpu list, mempolicy); agreement on any knob engages + // the bind. Otherwise keep per-call guards active and give the thread + // plain NUMA locality at most. + if (bind_decoders) { + for (auto &d : decoders_) { + if (!d) + continue; + try { + d->bind_current_thread(); + } catch (const std::exception &e) { + cudaq::qec::detail_affinity::affinity_warn( + "host loop thread bind failed for decoder " + + std::to_string(d->get_decoder_id()) + ": " + + std::string(e.what()) + " — continuing without guard-skip"); + } + } + } else if (node >= 0) { + try { + cudaq::qec::detail_affinity::bind_this_thread_to_numa_node(node); + } catch (const std::exception &e) { + cudaq::qec::detail_affinity::affinity_warn( + "host loop thread NUMA bind failed: " + std::string(e.what()) + + " — continuing without affinity"); + } + } + cudaq_host_dispatcher_loop(&host_ctx_); + }); + } } //============================================================================== @@ -975,6 +1194,13 @@ void qec_realtime_session::stop_loops() { if (host_loop_thread_.joinable()) host_loop_thread_.join(); + // The dispatch thread may have registered itself on the decoders via + // bind_current_thread(); it is gone now — clear the registrations before + // its thread id can be recycled by an unrelated new thread. + for (auto &d : decoders_) + if (d) + d->unbind_thread(); + if (device_dispatcher_) { cudaq_dispatcher_stop(device_dispatcher_); cudaq_dispatcher_destroy(device_dispatcher_); @@ -985,6 +1211,16 @@ void qec_realtime_session::stop_loops() { device_manager_ = nullptr; } + // Drain every worker stream before anything frees graph IO buffers: + // stream handles carry their device, so this synchronizes correctly even + // when the tearing-down thread's current device differs. In-flight graph + // launches from the non-blocking submit path must complete before + // release_decode_graph()/finalize() free the memory they read. + for (auto s : host_worker_streams_) { + if (s) + cudaStreamSynchronize(s); + } + for (auto s : host_worker_streams_) { if (s) cudaStreamDestroy(s); diff --git a/libs/qec/lib/realtime/qec_realtime_session.h b/libs/qec/lib/realtime/qec_realtime_session.h index 833aa8805..9f1204feb 100644 --- a/libs/qec/lib/realtime/qec_realtime_session.h +++ b/libs/qec/lib/realtime/qec_realtime_session.h @@ -208,6 +208,22 @@ class __attribute__((visibility("default"))) qec_realtime_session { // ---- HOST_LOOP wiring (both modes) ---- cudaq_host_dispatch_loop_ctx_t host_ctx_{}; std::thread host_loop_thread_; + /// Single NUMA node this session pins to (dispatch thread + ring buffers). + /// -1 = no pinning (unset, or decoders disagree — see initialize()). + int session_numa_node_ = -1; + /// Single CUDA device this session's decoders agreed on (-1 = none set, + /// or decoders disagree in HOST mode). DEVICE mode captures graphs, creates + /// worker streams, and launches the dispatch kernel on this device; + /// conflicting devices are rejected there. + int session_cuda_device_ = -1; + /// True when every decoder that sets a placement knob agrees on ALL of + /// cuda_device_id, numa_node_id, cpu_affinity, and mempolicy, and at least + /// one of numa_node_id / cuda_device_id / cpu_affinity is actually set — + /// only then may the shared dispatch thread register itself via + /// bind_current_thread() (a bound decoder skips its per-call guard, so + /// binding under any placement conflict would run some decoder on the + /// wrong GPU, node, CPU set, or memory policy). + bool bind_decoders_to_host_loop_ = false; std::uint64_t host_stats_counter_ = 0; // Plain (non-pinned) shutdown flag for HOST mode (no device kernel shares // it). diff --git a/libs/qec/lib/realtime/realtime_decoding.cpp b/libs/qec/lib/realtime/realtime_decoding.cpp index 8ff58ba3b..5471dc164 100644 --- a/libs/qec/lib/realtime/realtime_decoding.cpp +++ b/libs/qec/lib/realtime/realtime_decoding.cpp @@ -183,6 +183,14 @@ namespace cudaq::qec::decoding::host { cudaqx::heterogeneous_map prepare_decoder_params( const cudaq::qec::decoding::config::decoder_config &decoder_config) { auto params = decoder_config.decoder_custom_args_to_heterogeneous_map(); + if (decoder_config.cuda_device_id.has_value()) + params.insert("cuda_device_id", decoder_config.cuda_device_id.value()); + if (decoder_config.numa_node_id.has_value()) + params.insert("numa_node_id", decoder_config.numa_node_id.value()); + if (decoder_config.mempolicy.has_value()) + params.insert("mempolicy", decoder_config.mempolicy.value()); + if (decoder_config.cpu_affinity.has_value()) + params.insert("cpu_affinity", decoder_config.cpu_affinity.value()); if (decoder_config.type != "trt_decoder") return params; diff --git a/libs/qec/python/bindings/py_decoder.cpp b/libs/qec/python/bindings/py_decoder.cpp index 69af10ad9..87ecb955a 100644 --- a/libs/qec/python/bindings/py_decoder.cpp +++ b/libs/qec/python/bindings/py_decoder.cpp @@ -222,13 +222,20 @@ class PyDecoderRegistry { static std::unordered_map> registry; + // Whether the registered class defines decode_batch/decode_async in Python + // (recorded at registration time from the class namespace). Such overrides + // are dispatched by Python attribute lookup, bypassing the guarded C++ + // entry points, so hardware placement kwargs cannot take effect on them. + static std::unordered_map shadow_flags; public: static void register_decoder(const std::string &name, - std::function factory) { + std::function factory, + bool shadows_guarded_entries) { cudaq::qec::info("Registering Pythonic Decoder with name {}", name); registry[name] = factory; + shadow_flags[name] = shadows_guarded_entries; } static nb::object get_decoder(const std::string &name, nb::object H, @@ -244,13 +251,81 @@ class PyDecoderRegistry { static bool contains(const std::string &name) { return registry.find(name) != registry.end(); } + + static bool shadows_guarded_entries(const std::string &name) { + auto it = shadow_flags.find(name); + return it != shadow_flags.end() && it->second; + } }; std::unordered_map> PyDecoderRegistry::registry; +std::unordered_map PyDecoderRegistry::shadow_flags; namespace { +bool is_affinity_kwarg(const std::string &k) { + return k == "cuda_device_id" || k == "numa_node_id" || k == "mempolicy" || + k == "cpu_affinity"; +} + +// Mirror decoder::get()'s affinity contract for Python-registered decoders: +// strip the affinity keys before __init__ (a strict decoder must not reject +// them) and store them on the C++ base object so the guarded entry points +// (decode_batch / decode(tensor) / decode_async / enqueue_syndrome) pin +// correctly. Construction-time placement is NOT guarded on this path — the +// RAII guards are private to decoder.cpp and the bindings stay CUDA-free. +// Only the affinity keys are converted through hetMapFromKwargs: converting +// every kwarg would throw on Python types it cannot map (callables, None, +// tuples, ...) that a BYOD __init__ may legitimately accept. +nb::object get_py_registered_decoder(const std::string &name, nb::object H, + nb::kwargs options) { + nb::dict stripped; + nb::dict affinity_kwargs; + for (auto item : options) { + if (is_affinity_kwarg(nb::cast(item.first))) + affinity_kwargs[item.first] = item.second; + else + stripped[item.first] = item.second; + } + // A Python class that overrides decode_batch/decode_async in Python is + // called directly by Python attribute lookup — the guarded C++ entry + // points never run, so hardware kwargs would be accepted and silently + // dead. Reject loudly instead (before construction; the flag was recorded + // from the class namespace at decorator-registration time). + if (affinity_kwargs.size() > 0 && + PyDecoderRegistry::shadows_guarded_entries(name)) + throw std::runtime_error( + "Decoder '" + name + + "' overrides decode_batch/decode_async in " + "Python, so the hardware placement kwargs (cuda_device_id/" + "numa_node_id/mempolicy/cpu_affinity) cannot be applied to those " + "calls. Remove the kwargs, or implement only decode() and use the " + "guarded base-class batch entry points."); + nb::object obj = PyDecoderRegistry::get_decoder( + name, H, nb::borrow(stripped.ptr())); + // No affinity kwargs: preserve exact pre-existing behavior (including BYOD + // classes that never call qec.Decoder.__init__). + if (affinity_kwargs.size() == 0) + return obj; + // Narrow try: only the base-class cast maps to the "uninitialized base" + // message. hetMapFromKwargs can itself throw nb::cast_error (e.g. a negative + // Python int for cuda_device_id) and must not be misreported as a missing + // qec.Decoder.__init__ call. + decoder *base = nullptr; + try { + base = &nb::cast(obj); + } catch (const nb::cast_error &) { + throw std::runtime_error( + "Python decoder '" + name + + "' did not initialize its qec.Decoder base; call " + "qec.Decoder.__init__(self, H) in __init__ before using hardware " + "affinity kwargs (cuda_device_id/numa_node_id/mempolicy/cpu_affinity)"); + } + base->set_hardware_params( + hetMapFromKwargs(nb::borrow(affinity_kwargs.ptr()))); + return obj; +} struct batch_decoder_result { // Python-facing constructor for decoder plugin authors. nanobind enforces @@ -683,10 +758,16 @@ void bindDecoder(nb::module_ &mod) { .def( "decode", [](decoder &decoder, const std::vector &syndrome) { - return decoder.decode(syndrome); + return decoder.decode_guarded(syndrome); }, - "Decode the given syndrome to determine the error correction", - nb::arg("syndrome")) + "Decode the given syndrome to determine the error correction. " + "Applies the decoder's hardware placement automatically.", + nb::arg("syndrome"), + // Runs without the GIL so other Python threads can decode + // concurrently. Safe: the lambda is pure C++, and a Python-side + // decode override re-acquires the GIL in nanobind's trampoline + // (PyGILState_Ensure in trampoline_enter). + nb::call_guard()) .def( "decode_async", [](decoder &dec, @@ -700,11 +781,45 @@ void bindDecoder(nb::module_ &mod) { "decode_batch", [](decoder &decoder, const std::vector> &syndrome) { - auto results = decoder.decode_batch(syndrome); + std::vector results; + { + // Release the GIL only around the C++ decode: a call_guard + // would also cover makeBatchDecoderResult, which builds Python + // objects and must hold the GIL. A Python-side decode override + // reached from the batch loop re-acquires the GIL in nanobind's + // trampoline (PyGILState_Ensure in trampoline_enter). + nb::gil_scoped_release release; + results = decoder.decode_batch(syndrome); + } return makeBatchDecoderResult(results); }, "Decode multiple syndromes and return the results", nb::arg("syndrome")) + .def("cuda_device_id", &decoder::cuda_device_id, + "Target CUDA device for this decoder (-1 = inherit caller's " + "device).") + .def("numa_node_id", &decoder::numa_node_id, + "Target NUMA node for this decoder (-1 = no binding).") + .def( + "mempolicy", + [](decoder &d) { + return d.mempolicy() == cudaq::qec::mempolicy_mode::bind + ? "bind" + : "preferred"; + }, + "NUMA memory-policy mode for this decoder.") + .def("cpu_affinity", &decoder::cpu_affinity, + "Explicit CPU core list for this decoder's owning thread " + "(empty = node cpuset).") + .def("bind_current_thread", &decoder::bind_current_thread, + "Persistently pin the calling thread to this decoder's device/" + "node; guarded entry points then skip their per-call guard on " + "this thread. Call unbind_thread() before the thread exits.", + // Pure C++ (affinity/mempolicy/CUDA-device syscalls); runs without + // the GIL so pinning cannot stall other Python threads. + nb::call_guard()) + .def("unbind_thread", &decoder::unbind_thread, + "Forget a bind_current_thread() registration.") .def("get_block_size", &decoder::get_block_size, "Get the size of the code block") .def("get_syndrome_size", &decoder::get_syndrome_size, @@ -840,6 +955,14 @@ void bindDecoder(nb::module_ &mod) { if (!nb::hasattr(decoder_class, "decode")) throw std::runtime_error("Decoder class must implement decode method"); + // Record whether the class defines decode_batch/decode_async in + // PYTHON. The check must run here, on the class namespace: hasattr on + // an instance (or on new_class) also sees the guarded C++ bindings + // inherited from qec.Decoder and would always be true. + const bool shadows_guarded_entries = + namespace_dict.contains("decode_batch") || + namespace_dict.contains("decode_async"); + // Use Python's type() so the correct metaclass (nanobind's) is resolved nb::object type_fn = nb::module_::import_("builtins").attr("type"); nb::object new_class = @@ -847,10 +970,12 @@ void bindDecoder(nb::module_ &mod) { // Register the new class in the decoder registry PyDecoderRegistry::register_decoder( - name, [new_class](nb::object H, nb::kwargs options) { + name, + [new_class](nb::object H, nb::kwargs options) { nb::object instance = new_class(H, **options); return instance; - }); + }, + shadows_guarded_entries); return new_class; }); }); @@ -870,7 +995,7 @@ void bindDecoder(nb::module_ &mod) { options["error_rate_vec"] = copyToPyArray(*defaults.error_rate_vec); nb::object H_obj = copyToPyArray(dem.detector_error_matrix); - return PyDecoderRegistry::get_decoder(name, H_obj, options); + return get_py_registered_decoder(name, H_obj, options); } return get_decoder(name, decoder_init{dem_text}, hetMapFromKwargs(options)); @@ -887,7 +1012,7 @@ void bindDecoder(nb::module_ &mod) { } if (PyDecoderRegistry::contains(name)) { - return PyDecoderRegistry::get_decoder(name, H, options); + return get_py_registered_decoder(name, H, options); } cudaq::qec::sparse_binary_matrix H_sparse; diff --git a/libs/qec/python/bindings/py_decoding_config.cpp b/libs/qec/python/bindings/py_decoding_config.cpp index 283c90af1..e64a01982 100644 --- a/libs/qec/python/bindings/py_decoding_config.cpp +++ b/libs/qec/python/bindings/py_decoding_config.cpp @@ -284,6 +284,10 @@ void bindDecodingConfig(nb::module_ &mod) { .def_rw("H_sparse", &decoder_config::H_sparse) .def_rw("O_sparse", &decoder_config::O_sparse) .def_rw("D_sparse", &decoder_config::D_sparse) + .def_rw("cuda_device_id", &decoder_config::cuda_device_id) + .def_rw("numa_node_id", &decoder_config::numa_node_id) + .def_rw("mempolicy", &decoder_config::mempolicy) + .def_rw("cpu_affinity", &decoder_config::cpu_affinity) .def_rw("decoder_custom_args", &decoder_config::decoder_custom_args) .def( "set_decoder_custom_args", diff --git a/libs/qec/python/tests/test_decoder.py b/libs/qec/python/tests/test_decoder.py index 3201de1c6..658d181e0 100644 --- a/libs/qec/python/tests/test_decoder.py +++ b/libs/qec/python/tests/test_decoder.py @@ -1019,5 +1019,188 @@ def test_get_decoder_stim_dem_without_observables_returns_errors(): assert list(result.result) == [1.0] +def test_python_decoder_affinity_kwargs_stripped_and_stored(): + + @qec.decoder("strict_affinity_byod") + class StrictAffinityDecoder: + # Deliberately NO **kwargs: affinity keys must be stripped before + # __init__, mirroring the C++ decoder::get() contract. + def __init__(self, H): + qec.Decoder.__init__(self, H) + + def decode(self, syndrome): + r = qec.DecoderResult() + r.converged = True + r.result = [0.0, 0.0, 0.0] + return r + + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + # numa_node_id must not reach __init__ (no TypeError) and must be stored + # on the C++ base so guarded entry points pin correctly. + d = qec.get_decoder("strict_affinity_byod", H, numa_node_id=0) + assert d.numa_node_id() == 0 + assert d.cuda_device_id() == -1 + results = d.decode_batch([[0.0, 0.0], [0.1, 0.1]]) + assert len(results) == 2 + + +def test_python_decoder_affinity_kwargs_coexist_with_arbitrary_kwargs(): + received = {} + + @qec.decoder("affinity_plus_callback_byod") + class AffinityPlusCallbackDecoder: + + def __init__(self, H, **kwargs): + qec.Decoder.__init__(self, H) + received.update(kwargs) + + def decode(self, syndrome): + r = qec.DecoderResult() + r.converged = True + r.result = [0.0, 0.0, 0.0] + return r + + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + cb = lambda x: x + # Non-affinity kwargs of arbitrary Python types (e.g. callables) must pass + # through to __init__ untouched — only the affinity keys are converted to + # the C++ heterogeneous map, so this must construct without raising. + d = qec.get_decoder("affinity_plus_callback_byod", + H, + numa_node_id=0, + callback=cb) + assert d.numa_node_id() == 0 + assert received["callback"] is cb + assert "numa_node_id" not in received + + +def test_native_decoder_accepts_affinity_kwargs_from_python(): + # Native (C++-registered) decoders route through decoder::get(), which + # strips the affinity keys before the plugin constructor and stores them + # on the base via set_hardware_params(). + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + d = qec.get_decoder("multi_error_lut", H, numa_node_id=0) + assert d.numa_node_id() == 0 + assert d.cuda_device_id() == -1 + r = d.decode([0.0, 0.0]) + assert r is not None + + +def test_dem_text_overload_strips_affinity_kwargs_for_python_decoder(): + + @qec.decoder("dem_strict_affinity_byod_neg") + class DemStrictAffinityDecoder: + # Strict signature, deliberately NO **kwargs: the DEM-text path + # injects O and error_rate_vec kwargs, so exactly those (plus H) + # must arrive here — the affinity keys must already be stripped. + def __init__(self, H, O, error_rate_vec): + qec.Decoder.__init__(self, H) + self.num_obs = O.shape[0] + + def decode(self, syndrome): + r = qec.DecoderResult() + r.converged = True + r.result = [0.0] * self.num_obs + return r + + dem_text = ("error(0.1) D0 L0\n" + "error(0.1) D1 L0\n" + "error(0.05) D0 D1\n") + d = qec.get_decoder("dem_strict_affinity_byod_neg", + dem_text, + numa_node_id=0) + assert d.numa_node_id() == 0 + r = d.decode([0.0, 0.0]) + assert r.converged is True + + +def test_uninitialized_base_with_affinity_kwargs_raises_actionable_error(): + + @qec.decoder("no_base_byod_neg") + class NoBaseDecoder: + + def __init__(self, H): + pass # deliberately never calls qec.Decoder.__init__ + + def decode(self, syndrome): + return None + + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + with pytest.raises(RuntimeError, match="qec.Decoder.__init__"): + qec.get_decoder("no_base_byod_neg", H, numa_node_id=0) + + +def test_explicit_negative_cuda_device_id_raises(): + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + # The failure comes from the kwargs converter: hetMapFromKwargs stores + # Python ints as std::size_t, and nanobind's cast rejects negative ints + # for unsigned storage. That nb::cast_error surfaces as TypeError or + # RuntimeError depending on the nanobind translation, with no stable + # message text to match. + with pytest.raises((RuntimeError, TypeError)): + qec.get_decoder("multi_error_lut", H, cuda_device_id=-1) + + +def test_non_integral_cpu_affinity_raises_naming_integers(): + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + # Python lists arrive in the heterogeneous map as vector; + # read_cpu_affinity() rejects non-integral core ids at construction. + with pytest.raises(RuntimeError, match="must be integers"): + qec.get_decoder("multi_error_lut", + H, + numa_node_id=0, + cpu_affinity=[0.5]) + + +def test_invalid_mempolicy_string_raises_naming_the_knob(): + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + # read_mempolicy() accepts only "preferred" or "bind"; anything else is + # malformed explicit input and must throw, naming the knob. + with pytest.raises(RuntimeError, match="mempolicy"): + qec.get_decoder("multi_error_lut", H, numa_node_id=0, mempolicy="bnid") + + +def test_registered_decoder_with_python_decode_batch_rejects_affinity_kwargs(): + + @qec.decoder("shadowing_byod_neg") + class ShadowingDecoder: + + def __init__(self, H, **kwargs): + qec.Decoder.__init__(self, H) + + def decode(self, syndrome): + r = qec.DecoderResult() + r.converged = True + r.result = [0.0, 0.0, 0.0] + return r + + # A Python-level decode_batch is dispatched by attribute lookup and + # bypasses the guarded C++ entry, so placement kwargs cannot apply. + def decode_batch(self, syndromes): + return [self.decode(s) for s in syndromes] + + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + with pytest.raises(RuntimeError, match="decode_batch"): + qec.get_decoder("shadowing_byod_neg", H, numa_node_id=0) + # Without affinity kwargs the same class constructs fine: + d = qec.get_decoder("shadowing_byod_neg", H) + assert d is not None + + +def test_python_bind_unbind_and_readback_surface(): + H = np.array([[1, 1, 0], [0, 1, 1]], dtype=np.uint8) + d = qec.get_decoder("multi_error_lut", + H, + numa_node_id=0, + mempolicy="bind", + cpu_affinity=[0]) + assert d.mempolicy() == "bind" + assert d.cpu_affinity() == [0] + d.bind_current_thread() + r = d.decode([0.0, 0.0]) # routes through the guarded single-shot entry + assert r is not None + d.unbind_thread() + + if __name__ == "__main__": pytest.main() diff --git a/libs/qec/unittests/CMakeLists.txt b/libs/qec/unittests/CMakeLists.txt index bddc8dd6e..5a269c706 100644 --- a/libs/qec/unittests/CMakeLists.txt +++ b/libs/qec/unittests/CMakeLists.txt @@ -35,20 +35,45 @@ find_package(CUDAToolkit REQUIRED) add_compile_options(-Wno-attributes) add_executable(test_decoders test_decoders.cpp decoders/sample_decoder.cpp) -target_link_libraries(test_decoders PRIVATE GTest::gtest_main cudaq-qec cudaq-qec-realtime-decoding cudaq::cudaq libstim) +target_link_libraries(test_decoders PRIVATE GTest::gtest_main cudaq-qec cudaq-qec-realtime-decoding cudaq::cudaq libstim CUDA::cudart) +target_include_directories(test_decoders PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) add_dependencies(CUDAQXQECUnitTests test_decoders) gtest_discover_tests(test_decoders) + add_executable(test_decoders_yaml test_decoders_yaml.cpp decoders/sample_decoder.cpp) target_link_libraries(test_decoders_yaml PRIVATE GTest::gtest_main cudaq-qec cudaq-qec-realtime-decoding cudaq-qec-realtime-decoding-simulation - cudaq::cudaq) + cudaq::cudaq + CUDA::cudart) +target_include_directories(test_decoders_yaml PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) add_dependencies(CUDAQXQECUnitTests test_decoders_yaml) gtest_discover_tests(test_decoders_yaml) +# LD_PRELOAD interposer + benchmark/syscall-gate test for hardware pinning. +add_library(affinity_syscall_shim SHARED support/affinity_syscall_shim.cpp) +target_link_libraries(affinity_syscall_shim PRIVATE ${CMAKE_DL_LIBS}) + +add_executable(test_pinning_benchmark test_pinning_benchmark.cpp) +target_link_libraries(test_pinning_benchmark PRIVATE + GTest::gtest_main cudaq-qec cudaq::cudaq CUDA::cudart) +target_include_directories(test_pinning_benchmark PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) +add_dependencies(test_pinning_benchmark affinity_syscall_shim) +add_dependencies(CUDAQXQECUnitTests test_pinning_benchmark) +gtest_discover_tests(test_pinning_benchmark + PROPERTIES ENVIRONMENT "LD_PRELOAD=$") + +add_executable(test_device_affinity test_device_affinity.cpp) +target_link_libraries(test_device_affinity PRIVATE GTest::gtest_main cudaq-qec cudaq::cudaq CUDA::cudart) +target_include_directories(test_device_affinity PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/../lib + ${CUDAToolkit_INCLUDE_DIRS}) +add_dependencies(CUDAQXQECUnitTests test_device_affinity) +gtest_discover_tests(test_device_affinity) + add_executable(test_qec test_qec.cpp) target_link_libraries(test_qec PRIVATE GTest::gtest_main cudaq-qec cudaq::cudaq-stim-target) add_dependencies(CUDAQXQECUnitTests test_qec) @@ -72,8 +97,9 @@ target_link_libraries(test_logger PRIVATE GTest::gtest_main cudaq-qec cudaq::cud add_dependencies(CUDAQXQECUnitTests test_logger) gtest_discover_tests(test_logger) -# TensorRT decoder test is only built for x86 architectures -if(CUDAQ_QEC_TRT_DECODER_ENABLED AND CMAKE_SYSTEM_PROCESSOR MATCHES "(x86_64)|(AMD64|amd64)|(^i.86$)") +# Built wherever the TRT decoder plugin is enabled (x86_64 and arm64); the +# flag is false on platforms without TensorRT (e.g. arm64 + CUDA 12). +if(CUDAQ_QEC_TRT_DECODER_ENABLED) add_executable(test_trt_decoder ./decoders/trt_decoder/test_trt_decoder.cpp) # Find TensorRT for the test @@ -151,7 +177,11 @@ if(CUDAQ_QEC_TRT_DECODER_ENABLED AND CMAKE_SYSTEM_PROCESSOR MATCHES "(x86_64)|(A target_compile_definitions(test_trt_decoder PRIVATE TRT_TEST_ONNX_PATH="${CMAKE_CURRENT_SOURCE_DIR}/../../../assets/tests/surface_code_decoder.onnx") add_dependencies(CUDAQXQECUnitTests test_trt_decoder) - gtest_discover_tests(test_trt_decoder) + # Preload the affinity syscall shim so the guard tests can assert syscall + # counts (they degrade to placement-only asserts without it). + add_dependencies(test_trt_decoder affinity_syscall_shim) + gtest_discover_tests(test_trt_decoder + PROPERTIES ENVIRONMENT "LD_PRELOAD=$") endif() # ============================================================================== diff --git a/libs/qec/unittests/decoders/pymatching/test_pymatching_realtime.cpp b/libs/qec/unittests/decoders/pymatching/test_pymatching_realtime.cpp index c345fc95a..0123261ce 100644 --- a/libs/qec/unittests/decoders/pymatching/test_pymatching_realtime.cpp +++ b/libs/qec/unittests/decoders/pymatching/test_pymatching_realtime.cpp @@ -14,9 +14,14 @@ #include +#include +#include + #include +#include #include #include +#include #include #include @@ -162,6 +167,24 @@ TEST(PyMatchingRealtime, RejectsOversizedSyndromeRequest) { session.finalize(); } +TEST(PyMatchingRealtime, InitializeToleratesSparseDecoderIds) { + auto decoders = make_pymatching_decoders(/*h_vec=*/{1, 0, 1, 1, 0, 1}, + /*syndrome_size=*/3, + /*block_size=*/2); + auto other = make_pymatching_decoders(/*h_vec=*/{1, 0, 1, 1, 0, 1}, + /*syndrome_size=*/3, + /*block_size=*/2); + other[0]->set_decoder_id(2); + // Gap at index 1: decoder ids {0, 2}, mirroring a production config where + // decoder_config ids are not contiguous. + decoders.push_back(nullptr); + decoders.push_back(std::move(other[0])); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session.initialize(); + session.finalize(); +} + TEST(PyMatchingRealtime, ConfiguresViaRealtimeDecoderConfig) { namespace config = cudaq::qec::decoding::config; @@ -208,3 +231,301 @@ TEST(PyMatchingRealtime, ConfiguresViaRealtimeDecoderConfig) { config::finalize_decoders(); } + +//============================================================================== +// Session-level hardware-pinning conflict tests (consensus gate in +// qec_realtime_session::initialize()). Contract under test: decoders that +// disagree on an affinity knob must degrade LOUDLY (a CUDA_QEC_WARN naming the +// knob) but never break the session -- initialize() succeeds, every decoder +// still decodes correctly through the per-call guards, finalize() is clean. +//============================================================================== + +namespace { + +// H (3x2) shared by the conflict tests; syndrome {1,1,0} decodes to {1,0} +// (the same known-good pymatching mapping CheckRegularEdges relies on). +const std::vector kConflictH = {1, 0, 1, 1, 0, 1}; +constexpr std::size_t kConflictSyndromeSize = 3; +constexpr std::size_t kConflictBlockSize = 2; + +// Two identical pymatching decoders with ids {0, 1}. Affinity knobs are +// applied by each test via the public decoder::set_hardware_params, so the +// knobless (B4) path never touches the knob API at all. +DecoderVec make_session_pair() { + auto a = make_pymatching_decoders(kConflictH, kConflictSyndromeSize, + kConflictBlockSize); + auto b = make_pymatching_decoders(kConflictH, kConflictSyndromeSize, + kConflictBlockSize); + b[0]->set_decoder_id(1); + DecoderVec decoders; + decoders.push_back(std::move(a[0])); + decoders.push_back(std::move(b[0])); + return decoders; +} + +// HOST-mode sessions are process-global-exclusive (g_active_decoders), so the +// slot must be released on EVERY exit path, assertion failures included. +// finalize() is documented idempotent + destructor-safe, so an extra explicit +// finalize() in the happy path is harmless. +struct session_finalizer { + cudaq::qec::realtime::qec_realtime_session &session; + ~session_finalizer() { session.finalize(); } +}; + +// initialize() while capturing stderr. Never lets an exception escape while +// the gtest capture is active (a dangling capture would poison later tests); +// the caller asserts on `threw` after the capture is closed. +std::string +initialize_capturing_stderr(cudaq::qec::realtime::qec_realtime_session &session, + bool &threw, std::string &what) { + threw = false; + what.clear(); + testing::internal::CaptureStderr(); + try { + session.initialize(); + } catch (const std::exception &e) { + threw = true; + what = e.what(); + } catch (...) { + threw = true; + what = "non-std::exception thrown"; + } + return testing::internal::GetCapturedStderr(); +} + +// Round-trip one known syndrome through `decoder_id` and require the exact +// pymatching correction for kConflictH. +// CUDA device count for the B2 gate. This test target does not link cudart +// directly (a direct cudaGetDeviceCount reference fails to link with "DSO +// missing from command line") and CMake edits are out of scope, but +// libcudart.so is already in the process image as a dependency of +// libcudaq-qec -- so resolve the symbol dynamically for the probe only. +// Returns -1 when the symbol or the CUDA runtime is unavailable. +int probe_cuda_device_count() { + void *sym = dlsym(RTLD_DEFAULT, "cudaGetDeviceCount"); + if (!sym) + return -1; + // ABI: cudaError_t is an int-sized enum; cudaSuccess == 0. + auto get_count = reinterpret_cast(sym); + int count = 0; + if (get_count(&count) != 0) + return -1; + return count; +} + +void expect_decoder_still_decodes( + cudaq::qec::realtime::qec_realtime_session &session, std::size_t decoder_id, + std::uint64_t tag) { + const std::vector syndrome{1, 1, 0}; + cudaq::qec::decoding::rpc_producer::enqueue_syndromes( + session, decoder_id, syndrome.data(), syndrome.size(), tag); + std::vector corrections(kConflictBlockSize, 0xCC); + cudaq::qec::decoding::rpc_producer::get_corrections( + session, decoder_id, corrections.data(), corrections.size(), + /*reset=*/1); + EXPECT_EQ(corrections, (std::vector{1, 0})) + << "decoder " << decoder_id + << " must still decode correctly after a session-level pinning conflict"; +} + +} // namespace + +// Conflict gate: same NUMA node, different non-empty cpu_affinity lists. +TEST(PyMatchingRealtime, SessionCpuAffinityConflictWarnsAndStillDecodes) { +#if defined(__linux__) + cpu_set_t allowed; + CPU_ZERO(&allowed); + ASSERT_EQ(sched_getaffinity(0, sizeof(allowed), &allowed), 0); + if (!CPU_ISSET(0, &allowed) || !CPU_ISSET(1, &allowed)) + GTEST_SKIP() << "CPUs 0 and 1 are not both in the allowed cpuset"; + + auto decoders = make_session_pair(); + cudaqx::heterogeneous_map knobs_a; + knobs_a.insert("numa_node_id", 0); // node 0 always exists + knobs_a.insert("cpu_affinity", std::vector{0}); + decoders[0]->set_hardware_params(knobs_a); + cudaqx::heterogeneous_map knobs_b; + knobs_b.insert("numa_node_id", 0); + knobs_b.insert("cpu_affinity", std::vector{1}); + decoders[1]->set_hardware_params(knobs_b); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + bool threw = false; + std::string what; + const std::string err = initialize_capturing_stderr(session, threw, what); + ASSERT_FALSE(threw) << "conflicting cpu_affinity lists must degrade to " + "per-call guards, not throw: " + << what; + EXPECT_NE(err.find("different cpu_affinity lists"), std::string::npos) + << "conflict must be loud; captured stderr: [" << err << "]"; + + expect_decoder_still_decodes(session, /*decoder_id=*/0, /*tag=*/1); + expect_decoder_still_decodes(session, /*decoder_id=*/1, /*tag=*/2); + EXPECT_NO_THROW(session.finalize()); +#else + GTEST_SKIP() << "Linux-only (sched_getaffinity probe)"; +#endif +} + +// Conflict gate: different cuda_device_ids (no numa knob -- exercises the +// dev-conflict leg; numa_node_id may be soft-derived from the device but on a +// conflict-free node set that leg stays quiet). +TEST(PyMatchingRealtime, SessionCudaDeviceConflictWarnsAndStillDecodes) { + const int device_count = probe_cuda_device_count(); + if (device_count < 2) + GTEST_SKIP() << "needs >= 2 CUDA devices, have " << device_count; + + auto decoders = make_session_pair(); + cudaqx::heterogeneous_map knobs_a; + knobs_a.insert("cuda_device_id", 0); + decoders[0]->set_hardware_params(knobs_a); + cudaqx::heterogeneous_map knobs_b; + knobs_b.insert("cuda_device_id", 1); + decoders[1]->set_hardware_params(knobs_b); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + bool threw = false; + std::string what; + const std::string err = initialize_capturing_stderr(session, threw, what); + ASSERT_FALSE(threw) << "conflicting cuda_device_ids must degrade to " + "per-call guards, not throw: " + << what; + EXPECT_NE(err.find("different cuda_device_ids"), std::string::npos) + << "conflict must be loud; captured stderr: [" << err << "]"; + + expect_decoder_still_decodes(session, /*decoder_id=*/0, /*tag=*/1); + expect_decoder_still_decodes(session, /*decoder_id=*/1, /*tag=*/2); + EXPECT_NO_THROW(session.finalize()); +} + +// Conflict gate: different numa_node_ids. Requires a multi-node host (CI +// lane); must SKIP cleanly on single-node boxes. +TEST(PyMatchingRealtime, SessionNumaNodeConflictWarnsAndStillDecodes) { + std::ifstream node1("/sys/devices/system/node/node1/cpulist"); + if (!node1.is_open()) + GTEST_SKIP() << "single NUMA node host; numa conflict not testable"; + + auto decoders = make_session_pair(); + cudaqx::heterogeneous_map knobs_a; + knobs_a.insert("numa_node_id", 0); + decoders[0]->set_hardware_params(knobs_a); + cudaqx::heterogeneous_map knobs_b; + knobs_b.insert("numa_node_id", 1); + decoders[1]->set_hardware_params(knobs_b); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + bool threw = false; + std::string what; + const std::string err = initialize_capturing_stderr(session, threw, what); + ASSERT_FALSE(threw) << "conflicting numa_node_ids must disable session NUMA " + "pinning, not throw: " + << what; + EXPECT_NE(err.find("different numa_node_ids"), std::string::npos) + << "conflict must be loud; captured stderr: [" << err << "]"; + + expect_decoder_still_decodes(session, /*decoder_id=*/0, /*tag=*/1); + expect_decoder_still_decodes(session, /*decoder_id=*/1, /*tag=*/2); + EXPECT_NO_THROW(session.finalize()); +} + +// Conflict gate: same NUMA node, different mempolicy modes. One dispatch +// thread can hold only one memory policy, so guard-skip binding must stay +// off, loudly, while per-call guards keep every decoder decoding correctly. +TEST(PyMatchingRealtime, SessionMempolicyConflictWarnsAndStillDecodes) { + auto decoders = make_session_pair(); + cudaqx::heterogeneous_map knobs_a; + knobs_a.insert("numa_node_id", 0); // node 0 always exists + knobs_a.insert("mempolicy", std::string("preferred")); + decoders[0]->set_hardware_params(knobs_a); + cudaqx::heterogeneous_map knobs_b; + knobs_b.insert("numa_node_id", 0); + knobs_b.insert("mempolicy", std::string("bind")); + decoders[1]->set_hardware_params(knobs_b); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + bool threw = false; + std::string what; + const std::string err = initialize_capturing_stderr(session, threw, what); + ASSERT_FALSE(threw) << "conflicting mempolicy modes must degrade to " + "per-call guards, not throw: " + << what; + EXPECT_NE(err.find("different mempolicy modes"), std::string::npos) + << "conflict must be loud; captured stderr: [" << err << "]"; + + expect_decoder_still_decodes(session, /*decoder_id=*/0, /*tag=*/1); + expect_decoder_still_decodes(session, /*decoder_id=*/1, /*tag=*/2); + EXPECT_NO_THROW(session.finalize()); +} + +// Agreement on ONLY cpu_affinity (no numa_node_id, no cuda_device_id) must +// still engage the dispatch-thread bind rather than silently dropping the +// knob: initialize is conflict-warning-free, both decoders decode correctly +// through the bound dispatch thread, and finalize round-trips cleanly. (This +// binary runs without the affinity syscall interposer, so the contract is +// asserted through the warning-silence + decode + finalize surface.) +TEST(PyMatchingRealtime, SessionCpuAffinityOnlyStillBinds) { +#if defined(__linux__) + cpu_set_t allowed; + CPU_ZERO(&allowed); + ASSERT_EQ(sched_getaffinity(0, sizeof(allowed), &allowed), 0); + if (!CPU_ISSET(0, &allowed)) + GTEST_SKIP() << "CPU 0 is not in the allowed cpuset"; + + auto decoders = make_session_pair(); + cudaqx::heterogeneous_map knobs; + knobs.insert("cpu_affinity", std::vector{0}); + decoders[0]->set_hardware_params(knobs); + decoders[1]->set_hardware_params(knobs); + + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + bool threw = false; + std::string what; + const std::string err = initialize_capturing_stderr(session, threw, what); + ASSERT_FALSE(threw) << "agreeing cpu_affinity lists must initialize " + "cleanly: " + << what; + EXPECT_EQ(err.find("request different"), std::string::npos) + << "agreeing knobs must not raise a conflict warning; captured " + "stderr: [" + << err << "]"; + EXPECT_EQ(err.find("[cudaq-qec affinity]"), std::string::npos) + << "binding to an allowed CPU must be silent; captured stderr: [" << err + << "]"; + + expect_decoder_still_decodes(session, /*decoder_id=*/0, /*tag=*/1); + expect_decoder_still_decodes(session, /*decoder_id=*/1, /*tag=*/2); + EXPECT_NO_THROW(session.finalize()); +#else + GTEST_SKIP() << "Linux-only (sched_getaffinity probe)"; +#endif +} + +// Silence contract: a session over decoders that never +// touched a pinning knob must emit ZERO affinity noise ("[cudaq-qec +// affinity]" is the fprintf marker of hardware_affinity.h's affinity_warn) +// across construct + initialize + finalize. +TEST(PyMatchingRealtime, SessionKnoblessDecodersEmitNoAffinityNoise) { + auto decoders = make_session_pair(); // set_hardware_params never called + bool threw = false; + std::string what; + testing::internal::CaptureStderr(); + { + cudaq::qec::realtime::qec_realtime_session session(decoders); + session_finalizer fin{session}; + try { + session.initialize(); + } catch (const std::exception &e) { + threw = true; + what = e.what(); + } + } // finalize() runs here, inside the capture + const std::string err = testing::internal::GetCapturedStderr(); + ASSERT_FALSE(threw) << "knobless session must initialize cleanly: " << what; + EXPECT_EQ(err.find("[cudaq-qec affinity]"), std::string::npos) + << "knobless session must be affinity-silent, got: [" << err << "]"; +} diff --git a/libs/qec/unittests/decoders/sample_decoder.cpp b/libs/qec/unittests/decoders/sample_decoder.cpp index 0d357b5d7..f09d71cde 100644 --- a/libs/qec/unittests/decoders/sample_decoder.cpp +++ b/libs/qec/unittests/decoders/sample_decoder.cpp @@ -7,8 +7,14 @@ ******************************************************************************/ #include "cudaq/qec/decoder.h" +#include #include +#if defined(__linux__) +#include +#include +#endif + using namespace cudaqx; namespace cudaq::qec { @@ -50,4 +56,68 @@ class sample_decoder : public decoder { CUDAQ_EXT_PT_REGISTER_TYPE(sample_decoder) +/// @brief Test-only decoder that records the calling thread's placement -- +/// the CUDA device a real allocation lands on and the raw MPOL_* mempolicy +/// mode -- both DURING construction (inside decoder::get()'s guarded window) +/// and inside each decode() call, and echoes them through result.result: +/// [0] = CUDA device at construction [1] = CUDA device inside decode() +/// [2] = mempolicy mode at construction [3] = mempolicy mode inside decode() +/// Unavailable probes (no CUDA / non-Linux / blocked syscall) report -1. +class placement_probe_decoder : public decoder { +private: + int ctor_device_ = -1; + int ctor_mempolicy_ = -1; + + static int current_cuda_device() { + int dev = -1; + void *p = nullptr; + if (cudaMalloc(&p, 16) == cudaSuccess && p) { + cudaPointerAttributes attr{}; + if (cudaPointerGetAttributes(&attr, p) == cudaSuccess) + dev = attr.device; // device the allocation actually landed on + cudaFree(p); + } else { + cudaGetDevice(&dev); + } + return dev; + } + + static int current_mempolicy_mode() { + int mode = -1; +#if defined(__linux__) + syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL); +#endif + return mode; + } + +public: + placement_probe_decoder(const cudaq::qec::sparse_binary_matrix &H, + const cudaqx::heterogeneous_map ¶ms) + : decoder(H), ctor_device_(current_cuda_device()), + ctor_mempolicy_(current_mempolicy_mode()) {} + + virtual decoder_result decode(const std::vector &syndrome) override { + decoder_result result; + result.converged = true; + result.result = + std::vector{static_cast(ctor_device_), + static_cast(current_cuda_device()), + static_cast(ctor_mempolicy_), + static_cast(current_mempolicy_mode())}; + return result; + } + + virtual ~placement_probe_decoder() {} + + CUDAQ_EXTENSION_CUSTOM_CREATOR_FUNCTION( + placement_probe_decoder, static std::unique_ptr create( + const cudaq::qec::decoder_init &init, + const cudaqx::heterogeneous_map ¶ms) { + return cudaq::qec::make_pcm_decoder(init, + params); + }) +}; + +CUDAQ_EXT_PT_REGISTER_TYPE(placement_probe_decoder) + } // namespace cudaq::qec diff --git a/libs/qec/unittests/decoders/trt_decoder/test_trt_decoder.cpp b/libs/qec/unittests/decoders/trt_decoder/test_trt_decoder.cpp index 05b21fd2a..eca2e4cc0 100644 --- a/libs/qec/unittests/decoders/trt_decoder/test_trt_decoder.cpp +++ b/libs/qec/unittests/decoders/trt_decoder/test_trt_decoder.cpp @@ -11,17 +11,31 @@ #include "cudaq/qec/trt_decoder_internal.h" #include #include +#include #include #include #include #include #include +#include #include #include +#if defined(__linux__) +#include +#include +#include +#include +#endif + using namespace cudaq::qec; +// Resolved from the LD_PRELOAD affinity shim when preloaded; weak so the +// binary still links (and counter asserts are skipped) without it. +extern "C" __attribute__((weak)) void cudaqx_affinity_syscall_reset(); +extern "C" __attribute__((weak)) long cudaqx_set_mempolicy_count(); + static bool gpu_available() { int count = 0; return cudaGetDeviceCount(&count) == cudaSuccess && count > 0; @@ -780,3 +794,415 @@ TEST_F(TRTDecoderTest, CompositeGlobalDecoderCombinesLogicalFrame) { // require actual TensorRT/CUDA initialization which is not available in the // test environment. Only parameter validation and utility function tests are // enabled above. + +TEST_F(TRTDecoderTest, CudaDeviceId_OutOfRangeThrows) { + if (!gpu_available()) + GTEST_SKIP() << "No CUDA GPU available"; + std::string onnx_path = get_onnx_asset_path(); + if (!std::filesystem::exists(onnx_path)) + GTEST_SKIP() << "ONNX model not found: " << onnx_path; + int count = 0; + cudaGetDeviceCount(&count); + cudaqx::tensor H_small({1, 1}); + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", onnx_path); + params.insert("cuda_device_id", count); // one past the last valid device + EXPECT_THROW( + { cudaq::qec::decoder::get("trt_decoder", H_small, params); }, + std::runtime_error); +} + +TEST_F(TRTDecoderTest, CudaDeviceId_DecodeAsyncOnGpu1) { + if (!gpu_available()) + GTEST_SKIP() << "No CUDA GPU available"; + int count = 0; + cudaGetDeviceCount(&count); + if (count < 2) + GTEST_SKIP() << "needs >= 2 GPUs to prove non-default pinning"; + std::string onnx_path = get_onnx_asset_path(); + if (!std::filesystem::exists(onnx_path)) + GTEST_SKIP() << "ONNX model not found: " << onnx_path; + + cudaSetDevice(0); // calling thread stays on the default device + + std::size_t num_detectors = NUM_DETECTORS; + cudaqx::tensor H_mat({num_detectors, num_detectors}); + for (std::size_t i = 0; i < num_detectors; ++i) + H_mat.at({i, i}) = 1; + + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", onnx_path); + params.insert("cuda_device_id", 1); // engine/buffers must land on GPU 1 + + std::unique_ptr d; + ASSERT_NO_THROW( + { d = cudaq::qec::decoder::get("trt_decoder", H_mat, params); }); + + // decode_async spawns a fresh thread (defaults to GPU 0). If trt did NOT + // re-assert the device, the decode would run on GPU 0 while the engine lives + // on GPU 1 -> error/garbage. A correct result proves the per-decode guard. + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + auto fut = d->decode_async(syndrome); + cudaq::qec::decoder_result res; + ASSERT_NO_THROW({ res = fut.get(); }); + ASSERT_FALSE(res.result.empty()); + float trt_output = res.result[0]; + float expected_output = TEST_OUTPUTS[0][0]; + float error = std::abs(trt_output - expected_output); + EXPECT_LT(error, 1e-4f) << "GPU-1 decode differs from expected: got " + << trt_output << ", expected " << expected_output; +} + +// Two real trt decoders, each pinned to a separate GPU, run concurrently from +// two independent threads and both converge to the correct output. +TEST_F(TRTDecoderTest, TwoTrtDecodersConcurrently) { + if (!gpu_available()) + GTEST_SKIP() << "No CUDA GPU available"; + int count = 0; + cudaGetDeviceCount(&count); + if (count < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + std::string onnx_path = get_onnx_asset_path(); + if (!std::filesystem::exists(onnx_path)) + GTEST_SKIP() << "ONNX model not found: " << onnx_path; + + std::size_t num_detectors = NUM_DETECTORS; + cudaqx::tensor H_mat({num_detectors, num_detectors}); + for (std::size_t i = 0; i < num_detectors; ++i) + H_mat.at({i, i}) = 1; + + cudaqx::heterogeneous_map params0; + params0.insert("onnx_load_path", onnx_path); + params0.insert("cuda_device_id", 0); + cudaqx::heterogeneous_map params1; + params1.insert("onnx_load_path", onnx_path); + params1.insert("cuda_device_id", 1); + + std::unique_ptr dec0, dec1; + try { + dec0 = decoder::get("trt_decoder", H_mat, params0); + dec1 = decoder::get("trt_decoder", H_mat, params1); + } catch (const std::exception &e) { + GTEST_SKIP() << "TRT construction failed: " << e.what(); + } + + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + + cudaq::qec::decoder_result res0, res1; + std::exception_ptr ex0, ex1; + std::thread t0([&] { + try { + res0 = dec0->decode(syndrome); + } catch (...) { + ex0 = std::current_exception(); + } + }); + std::thread t1([&] { + try { + res1 = dec1->decode(syndrome); + } catch (...) { + ex1 = std::current_exception(); + } + }); + t0.join(); + t1.join(); + + if (ex0) + std::rethrow_exception(ex0); + if (ex1) + std::rethrow_exception(ex1); + + float expected_output = TEST_OUTPUTS[0][0]; + EXPECT_TRUE(res0.converged); + ASSERT_FALSE(res0.result.empty()); + EXPECT_LT(std::abs(res0.result[0] - expected_output), 1e-4f) + << "GPU-0 decode differs from expected: got " << res0.result[0] + << ", expected " << expected_output; + EXPECT_TRUE(res1.converged); + ASSERT_FALSE(res1.result.empty()); + EXPECT_LT(std::abs(res1.result[0] - expected_output), 1e-4f) + << "GPU-1 decode differs from expected: got " << res1.result[0] + << ", expected " << expected_output; +} + +// Container runtimes commonly block the mempolicy syscalls (seccomp without +// CAP_SYS_NICE). trt_decoder::decode_batch() degrades gracefully there (warns, +// keeps going), but a test that wants to exercise the mempolicy path must skip +// instead of fail. Mirrors the guard in test_device_affinity.cpp. +TEST(TrtDecoder, DecodeBatchAppliesNumaPolicy) { +#if defined(__linux__) + int mode = -1; + if (syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; +#endif + if (!gpu_available()) + GTEST_SKIP() << "No CUDA GPU available"; + std::string onnx_path = get_onnx_asset_path(); + if (!std::filesystem::exists(onnx_path)) + GTEST_SKIP() << "ONNX model not found: " << onnx_path; + + std::size_t num_detectors = NUM_DETECTORS; + cudaqx::tensor H({num_detectors, num_detectors}); + for (std::size_t i = 0; i < num_detectors; ++i) + H.at({i, i}) = 1; + + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", onnx_path); + params.insert("numa_node_id", 0); + + std::unique_ptr trt_decoder; + try { + trt_decoder = decoder::get("trt_decoder", H, params); + } catch (const std::exception &e) { + GTEST_SKIP() << "Failed to create TRT decoder: " << e.what(); + } + + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + std::vector results; + // decode_batch() applies the NUMA guard for the duration of the call and + // restores the thread's prior policy before returning, so this only proves + // no crash/throw occurs with a numa_node_id knob set; the mempolicy syscall + // behavior itself is covered at the primitive level by + // HardwareAffinity.MempolicyBindWhenRequested. + EXPECT_NO_THROW({ results = trt_decoder->decode_batch({syndrome}); }); + ASSERT_EQ(results.size(), 1u); + EXPECT_TRUE(results[0].converged); + ASSERT_FALSE(results[0].result.empty()); +} + +#if defined(__linux__) +// Snapshot of the calling thread's placement: CPU affinity mask plus thread +// memory-policy mode and nodemask (raw SYS_get_mempolicy, matching what the +// decode_batch guard saves/restores). +struct thread_placement { + cpu_set_t mask; + int mempolicy_mode = -1; + unsigned long nodemask[16] = {0}; +}; + +static std::optional capture_thread_placement() { + thread_placement p; + CPU_ZERO(&p.mask); + if (sched_getaffinity(0, sizeof(p.mask), &p.mask) != 0) + return std::nullopt; + if (syscall(SYS_get_mempolicy, &p.mempolicy_mode, p.nodemask, + sizeof(p.nodemask) * 8, nullptr, 0UL) != 0) + return std::nullopt; + return p; +} + +static void expect_same_placement(const thread_placement &before, + const thread_placement &after) { + EXPECT_TRUE(CPU_EQUAL(&before.mask, &after.mask)) + << "CPU affinity mask changed across decode_batch()"; + EXPECT_EQ(before.mempolicy_mode, after.mempolicy_mode) + << "thread mempolicy mode changed across decode_batch()"; + EXPECT_EQ( + 0, std::memcmp(before.nodemask, after.nodemask, sizeof(before.nodemask))) + << "thread mempolicy nodemask changed across decode_batch()"; +} + +// Shared skip guard for the placement tests below (GPU + ONNX asset + NUMA +// node 0 + mempolicy syscalls, which container seccomp commonly blocks). +static std::optional placement_test_skip_reason() { + if (!gpu_available()) + return "No CUDA GPU available"; + if (!std::filesystem::exists(get_onnx_asset_path())) + return "ONNX model not found: " + get_onnx_asset_path(); + if (!std::filesystem::exists("/sys/devices/system/node/node0")) + return "NUMA node 0 not present"; + int mode = -1; + if (syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) != 0) + return "mempolicy syscalls unavailable (container seccomp?)"; + return std::nullopt; +} + +static void expect_first_output_correct( + const std::vector &results) { + ASSERT_EQ(results.size(), 1u); + EXPECT_TRUE(results[0].converged); + ASSERT_FALSE(results[0].result.empty()); + EXPECT_LT(std::abs(results[0].result[0] - TEST_OUTPUTS[0][0]), 1e-4f) + << "decode_batch output differs from expected: got " + << results[0].result[0] << ", expected " << TEST_OUTPUTS[0][0]; +} +#endif + +// A thread that called bind_current_thread() owns its placement: the +// decode_batch() guard must not fire on it. Placement is captured after the +// bind and must be bit-identical after decode_batch() — regression test for +// the pinning-defeat bug where the guard re-bound (and widened) an +// already-bound thread. +TEST(TrtDecoder, BoundThreadSkipsGuardInDecodeBatch) { +#if defined(__linux__) + if (auto reason = placement_test_skip_reason()) + GTEST_SKIP() << *reason; + + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", get_onnx_asset_path()); + params.insert("cuda_device_id", 0); + params.insert("numa_node_id", 0); + std::unique_ptr trt_decoder; + try { + trt_decoder = + decoder::get("trt_decoder", make_identity_h(NUM_DETECTORS), params); + } catch (const std::exception &e) { + GTEST_SKIP() << "Failed to create TRT decoder: " << e.what(); + } + + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + // bind_current_thread() pins persistently (no restore), so run on a worker + // thread to keep the gtest main thread's placement intact for later tests. + std::optional before, after; + std::vector results; + std::exception_ptr err; + // A guard that fires and restores perfectly is indistinguishable from a + // skipped guard by before/after placement alone; the set_mempolicy counter + // is the observable difference (TRT itself never calls set_mempolicy). + const bool have_shim = cudaqx_affinity_syscall_reset != nullptr && + cudaqx_set_mempolicy_count != nullptr; + long setmem_during_decode = 0; + std::thread worker([&] { + try { + trt_decoder->bind_current_thread(); + before = capture_thread_placement(); + if (have_shim) + cudaqx_affinity_syscall_reset(); + results = trt_decoder->decode_batch({syndrome}); + if (have_shim) + setmem_during_decode = cudaqx_set_mempolicy_count(); + after = capture_thread_placement(); + } catch (...) { + err = std::current_exception(); + } + }); + worker.join(); + if (err) + std::rethrow_exception(err); + + expect_first_output_correct(results); + ASSERT_TRUE(before.has_value()); + ASSERT_TRUE(after.has_value()); + expect_same_placement(*before, *after); + if (have_shim) + EXPECT_EQ(setmem_during_decode, 0) + << "bound decode_batch() issued set_mempolicy: the guard fired on a " + "thread that already called bind_current_thread()"; +#else + GTEST_SKIP() << "Linux-only placement test"; +#endif +} + +// Unbound caller: decode_batch() applies the NUMA guard for the duration of +// the call and must restore the caller's placement exactly before returning. +TEST(TrtDecoder, UnboundDecodeBatchRestoresPlacement) { +#if defined(__linux__) + if (auto reason = placement_test_skip_reason()) + GTEST_SKIP() << *reason; + + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", get_onnx_asset_path()); + params.insert("cuda_device_id", 0); + params.insert("numa_node_id", 0); + std::unique_ptr trt_decoder; + try { + trt_decoder = + decoder::get("trt_decoder", make_identity_h(NUM_DETECTORS), params); + } catch (const std::exception &e) { + GTEST_SKIP() << "Failed to create TRT decoder: " << e.what(); + } + + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + auto before = capture_thread_placement(); + ASSERT_TRUE(before.has_value()); + auto results = trt_decoder->decode_batch({syndrome}); + auto after = capture_thread_placement(); + ASSERT_TRUE(after.has_value()); + + expect_first_output_correct(results); + expect_same_placement(*before, *after); +#else + GTEST_SKIP() << "Linux-only placement test"; +#endif +} + +// An explicit cpu_affinity={0} bind is strictly narrower than node 0's cpuset: +// if decode_batch()'s _ScopedNuma guard fired on the bound thread it would +// widen the mask to the whole node, so "exactly CPU 0 before AND after" proves +// the explicit pin survives decode_batch(). +TEST(TrtDecoder, ExplicitCpuAffinityHonoredWhenBound) { +#if defined(__linux__) + if (auto reason = placement_test_skip_reason()) + GTEST_SKIP() << *reason; + + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", get_onnx_asset_path()); + params.insert("numa_node_id", 0); + params.insert("cpu_affinity", std::vector{0}); + std::unique_ptr trt_decoder; + try { + trt_decoder = + decoder::get("trt_decoder", make_identity_h(NUM_DETECTORS), params); + } catch (const std::exception &e) { + GTEST_SKIP() << "Failed to create TRT decoder: " << e.what(); + } + + std::vector syndrome(TEST_INPUTS[0].begin(), + TEST_INPUTS[0].end()); + cpu_set_t only_cpu0; + CPU_ZERO(&only_cpu0); + CPU_SET(0, &only_cpu0); + + std::optional before, after; + std::vector results; + std::exception_ptr err; + std::thread worker([&] { + try { + trt_decoder->bind_current_thread(); + before = capture_thread_placement(); + results = trt_decoder->decode_batch({syndrome}); + after = capture_thread_placement(); + } catch (...) { + err = std::current_exception(); + } + }); + worker.join(); + if (err) + std::rethrow_exception(err); + + expect_first_output_correct(results); + ASSERT_TRUE(before.has_value()); + ASSERT_TRUE(after.has_value()); + EXPECT_TRUE(CPU_EQUAL(&only_cpu0, &before->mask)) + << "bind_current_thread() did not honor cpu_affinity={0}"; + EXPECT_TRUE(CPU_EQUAL(&only_cpu0, &after->mask)) + << "decode_batch() widened the explicit cpu_affinity={0} pin"; + expect_same_placement(*before, *after); +#else + GTEST_SKIP() << "Linux-only placement test"; +#endif +} + +// Dependency-free regression guard: decoder::get()'s own construction-time +// CudaDeviceGuard rejects an out-of-range cuda_device_id before the plugin is +// ever instantiated, so decode_batch()'s own device-guard error-checking is +// never reached for this particular failure mode. Kept as a guard on the +// overall contract (out-of-range device ids must never silently proceed). +TEST(TrtDecoder, DecodeBatchThrowsOnOutOfRangeCudaDeviceId) { + int count = 0; + cudaGetDeviceCount(&count); + cudaqx::tensor H({2, 2}); + H.at({0, 0}) = 1; + H.at({1, 1}) = 1; + cudaqx::heterogeneous_map params; + params.insert("onnx_load_path", get_onnx_asset_path()); + params.insert("cuda_device_id", count + 5); // deliberately out of range + EXPECT_THROW( + { cudaq::qec::decoder::get("trt_decoder", H, params); }, + std::runtime_error); +} diff --git a/libs/qec/unittests/support/affinity_syscall_shim.cpp b/libs/qec/unittests/support/affinity_syscall_shim.cpp new file mode 100644 index 000000000..3c6cdeb0b --- /dev/null +++ b/libs/qec/unittests/support/affinity_syscall_shim.cpp @@ -0,0 +1,298 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ +// Test-only LD_PRELOAD interposer: counts the placement-related syscalls so a +// test can assert a bound decode loop issues none (and an unbound one does). +// Counts five syscalls: sched_setaffinity / sched_getaffinity (glibc symbols) +// and SYS_set_mempolicy / SYS_get_mempolicy / SYS_mbind (issued through +// glibc's syscall(2) wrapper, so that wrapper is interposed too). Also counts +// opens of /sys/devices/system/node/* (the per-guard cpulist read) via the +// openat/openat64 glibc symbols. Not linked into the library. +// +// Fault injection (getenv per call, so a test can toggle around one decode): +// CUDAQX_SHIM_FAIL_SETAFFINITY sched_setaffinity -> EPERM, -1 +// CUDAQX_SHIM_FAIL_GETAFFINITY sched_getaffinity -> EPERM, -1 +// CUDAQX_SHIM_FAIL_SET_MEMPOLICY SYS_set_mempolicy -> EPERM, -1 +// CUDAQX_SHIM_FAIL_GET_MEMPOLICY SYS_get_mempolicy -> EPERM, -1 +// CUDAQX_SHIM_FAIL_MBIND SYS_mbind -> EPERM, -1 +// CUDAQX_SHIM_FAIL_ALL_PLACEMENT all of the above -> EPERM, -1 +// Every counter counts ATTEMPTS: the count is incremented before injection. +// The sysfs open counters have NO fault injection. +#ifndef _GNU_SOURCE +#define _GNU_SOURCE +#endif +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { +std::atomic g_setaff{0}; +std::atomic g_getaff{0}; +std::atomic g_setmem{0}; +std::atomic g_getmem{0}; +std::atomic g_mbind{0}; +std::atomic g_node_sysfs_open{0}; + +// True when the call must be failed: either its specific injection variable or +// the blanket CUDAQX_SHIM_FAIL_ALL_PLACEMENT is set. getenv per call so a test +// can set/unset around a single decode. +bool fail_injected(const char *specific) { + return std::getenv(specific) != nullptr || + std::getenv("CUDAQX_SHIM_FAIL_ALL_PLACEMENT") != nullptr; +} + +// dlsym may itself invoke interposed functions; the thread-local flag breaks +// that recursion (the inner call sees "resolving" and reports ENOSYS instead +// of re-entering dlsym forever). +using syscall_fn = long (*)(long, long, long, long, long, long, long); +syscall_fn real_syscall_lazy() { + static std::atomic cached{nullptr}; + syscall_fn fn = cached.load(std::memory_order_acquire); + if (fn) + return fn; + static thread_local bool resolving = false; + if (resolving) + return nullptr; + resolving = true; + fn = reinterpret_cast(dlsym(RTLD_NEXT, "syscall")); + resolving = false; + if (fn) + cached.store(fn, std::memory_order_release); + return fn; +} + +// openat's optional 4th argument exists only for these flags. +bool open_needs_mode(int flags) { + if (flags & O_CREAT) + return true; +#ifdef O_TMPFILE + if ((flags & O_TMPFILE) == O_TMPFILE) + return true; +#endif + return false; +} + +void count_node_sysfs_path(const char *pathname) { + if (pathname && std::strstr(pathname, "/sys/devices/system/node/")) + g_node_sysfs_open.fetch_add(1, std::memory_order_relaxed); +} +} // namespace + +extern "C" { + +// Interpose the glibc symbols; forward to the real ones via RTLD_NEXT. +// Counts first (attempts), then applies fault injection. +int sched_setaffinity(pid_t pid, size_t cpusetsize, const cpu_set_t *mask) { + static int (*real)(pid_t, size_t, const cpu_set_t *) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "sched_setaffinity")); + g_setaff.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_SETAFFINITY")) { + errno = EPERM; + return -1; + } + return real(pid, cpusetsize, mask); +} + +int sched_getaffinity(pid_t pid, size_t cpusetsize, cpu_set_t *mask) { + static int (*real)(pid_t, size_t, cpu_set_t *) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "sched_getaffinity")); + g_getaff.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_GETAFFINITY")) { + errno = EPERM; + return -1; + } + return real(pid, cpusetsize, mask); +} + +// The mempolicy calls have no glibc wrappers; the library reaches them through +// syscall(2), so interpose that. Counts ATTEMPTS (even ones seccomp rejects). +long syscall(long number, ...) { + va_list ap; + va_start(ap, number); + long a0 = va_arg(ap, long), a1 = va_arg(ap, long), a2 = va_arg(ap, long); + long a3 = va_arg(ap, long), a4 = va_arg(ap, long), a5 = va_arg(ap, long); + va_end(ap); + switch (number) { + case SYS_set_mempolicy: + g_setmem.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_SET_MEMPOLICY")) { + errno = EPERM; + return -1; + } + break; + case SYS_get_mempolicy: + g_getmem.fetch_add(1, std::memory_order_relaxed); + // Fault injection: simulate a seccomp profile that blocks get_mempolicy + // while allowing set_mempolicy (the un-restorable-policy scenario). + if (fail_injected("CUDAQX_SHIM_FAIL_GET_MEMPOLICY")) { + errno = EPERM; + return -1; + } + break; + case SYS_mbind: + g_mbind.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_MBIND")) { + errno = EPERM; + return -1; + } + break; + case SYS_sched_setaffinity: + g_setaff.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_SETAFFINITY")) { + errno = EPERM; + return -1; + } + break; + case SYS_sched_getaffinity: + g_getaff.fetch_add(1, std::memory_order_relaxed); + if (fail_injected("CUDAQX_SHIM_FAIL_GETAFFINITY")) { + errno = EPERM; + return -1; + } + break; + } + syscall_fn real = real_syscall_lazy(); + if (!real) { // recursion during resolve; refuse rather than loop + errno = ENOSYS; + return -1; + } + return real(number, a0, a1, a2, a3, a4, a5); +} + +// Count opens of the node-topology sysfs tree (build_node_cpuset's per-guard +// /sys/devices/system/node/node/cpulist read). On this platform libstdc++'s +// ifstream is stdio-based (__basic_file wraps FILE*), so the file is opened +// through glibc's fopen -- whose INTERNAL open bypasses every interposable +// open/openat symbol (measured: openat/openat64 alone counted 0, adding +// open/open64 still counted 0). fopen/fopen64 are therefore interposed too; +// the open* spellings are kept so the counter stays robust across libstdc++ +// builds that call them directly. Pass-through only: NO fault injection. +FILE *fopen(const char *pathname, const char *mode) { + static FILE *(*real)(const char *, const char *) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "fopen")); + count_node_sysfs_path(pathname); + return real(pathname, mode); +} + +FILE *fopen64(const char *pathname, const char *mode) { + static FILE *(*real)(const char *, const char *) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "fopen64")); + count_node_sysfs_path(pathname); + return real(pathname, mode); +} + +int open(const char *pathname, int flags, ...) { + static int (*real)(const char *, int, ...) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "open")); + count_node_sysfs_path(pathname); + if (open_needs_mode(flags)) { + va_list ap; + va_start(ap, flags); + mode_t mode = va_arg(ap, mode_t); + va_end(ap); + return real(pathname, flags, mode); + } + return real(pathname, flags); +} + +int open64(const char *pathname, int flags, ...) { + static int (*real)(const char *, int, ...) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "open64")); + count_node_sysfs_path(pathname); + if (open_needs_mode(flags)) { + va_list ap; + va_start(ap, flags); + mode_t mode = va_arg(ap, mode_t); + va_end(ap); + return real(pathname, flags, mode); + } + return real(pathname, flags); +} + +int openat(int dirfd, const char *pathname, int flags, ...) { + static int (*real)(int, const char *, int, ...) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "openat")); + count_node_sysfs_path(pathname); + if (open_needs_mode(flags)) { + va_list ap; + va_start(ap, flags); + mode_t mode = va_arg(ap, mode_t); + va_end(ap); + return real(dirfd, pathname, flags, mode); + } + return real(dirfd, pathname, flags); +} + +int openat64(int dirfd, const char *pathname, int flags, ...) { + static int (*real)(int, const char *, int, ...) = nullptr; + if (!real) + real = reinterpret_cast( + dlsym(RTLD_NEXT, "openat64")); + count_node_sysfs_path(pathname); + if (open_needs_mode(flags)) { + va_list ap; + va_start(ap, flags); + mode_t mode = va_arg(ap, mode_t); + va_end(ap); + return real(dirfd, pathname, flags, mode); + } + return real(dirfd, pathname, flags); +} + +// Read/reset hooks the test resolves (weakly) from the preloaded shim. +// Legacy pair kept for existing tests: count == sched_setaffinity count. +long cudaqx_affinity_syscall_count() { + return g_setaff.load(std::memory_order_relaxed); +} +void cudaqx_affinity_syscall_reset() { + g_setaff.store(0, std::memory_order_relaxed); + g_getaff.store(0, std::memory_order_relaxed); + g_setmem.store(0, std::memory_order_relaxed); + g_getmem.store(0, std::memory_order_relaxed); + g_mbind.store(0, std::memory_order_relaxed); + g_node_sysfs_open.store(0, std::memory_order_relaxed); +} +// Per-syscall counters. +long cudaqx_sched_setaffinity_count() { + return g_setaff.load(std::memory_order_relaxed); +} +long cudaqx_sched_getaffinity_count() { + return g_getaff.load(std::memory_order_relaxed); +} +long cudaqx_set_mempolicy_count() { + return g_setmem.load(std::memory_order_relaxed); +} +long cudaqx_get_mempolicy_count() { + return g_getmem.load(std::memory_order_relaxed); +} +long cudaqx_mbind_count() { return g_mbind.load(std::memory_order_relaxed); } +long cudaqx_node_sysfs_open_count() { + return g_node_sysfs_open.load(std::memory_order_relaxed); +} +} diff --git a/libs/qec/unittests/support/thread_placement.h b/libs/qec/unittests/support/thread_placement.h new file mode 100644 index 000000000..0af736c05 --- /dev/null +++ b/libs/qec/unittests/support/thread_placement.h @@ -0,0 +1,79 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ +// Test-only, header-only snapshot of the calling thread's hardware placement +// (CPU affinity + NUMA mempolicy mode + optional CUDA device) so tests can +// assert a guard restored everything exactly. Included via a relative path; +// not installed. +#ifndef CUDAQ_QEC_UNITTESTS_SUPPORT_THREAD_PLACEMENT_H +#define CUDAQ_QEC_UNITTESTS_SUPPORT_THREAD_PLACEMENT_H + +#include +#include +#include +#include +#if defined(__linux__) +#include +#include +#endif + +namespace cudaq::qec::test { + +struct thread_placement { + cpu_set_t affinity{}; + int mempolicy_mode = -1; // -1: query failed / unavailable (e.g. seccomp) + int cuda_device = -1; // -1: not captured (with_cuda=false) or no CUDA + + // NOTE: issues sched_getaffinity + SYS_get_mempolicy itself; call it outside + // any shim reset/read window. cudaGetDevice only runs when with_cuda is set + // so CPU-only tests never touch the CUDA runtime. + static thread_placement capture(bool with_cuda = false) { + thread_placement p; + CPU_ZERO(&p.affinity); + (void)sched_getaffinity(0, sizeof(p.affinity), &p.affinity); +#if defined(__linux__) + if (syscall(SYS_get_mempolicy, &p.mempolicy_mode, nullptr, 0UL, nullptr, + 0UL) != 0) + p.mempolicy_mode = -1; +#endif + if (with_cuda && cudaGetDevice(&p.cuda_device) != cudaSuccess) + p.cuda_device = -1; + return p; + } + + bool operator==(const thread_placement &o) const { + return CPU_EQUAL(&affinity, &o.affinity) && + mempolicy_mode == o.mempolicy_mode && cuda_device == o.cuda_device; + } + bool operator!=(const thread_placement &o) const { return !(*this == o); } + + // Names which field differs (for assertion messages). + std::string describe_difference(const thread_placement &o) const { + std::string out; + if (!CPU_EQUAL(&affinity, &o.affinity)) + out += "cpu affinity differs (" + std::to_string(CPU_COUNT(&affinity)) + + " vs " + std::to_string(CPU_COUNT(&o.affinity)) + " cpus set); "; + if (mempolicy_mode != o.mempolicy_mode) + out += "mempolicy_mode differs (" + std::to_string(mempolicy_mode) + + " vs " + std::to_string(o.mempolicy_mode) + "); "; + if (cuda_device != o.cuda_device) + out += "cuda_device differs (" + std::to_string(cuda_device) + " vs " + + std::to_string(o.cuda_device) + "); "; + return out.empty() ? "identical" : out; + } +}; + +// gtest hook so failed EXPECT_EQ prints something readable. +inline void PrintTo(const thread_placement &p, std::ostream *os) { + *os << "{cpus_set=" << CPU_COUNT(&p.affinity) + << ", mempolicy_mode=" << p.mempolicy_mode + << ", cuda_device=" << p.cuda_device << "}"; +} + +} // namespace cudaq::qec::test + +#endif // CUDAQ_QEC_UNITTESTS_SUPPORT_THREAD_PLACEMENT_H diff --git a/libs/qec/unittests/test_decoders.cpp b/libs/qec/unittests/test_decoders.cpp index 0f2f4c8d5..442ad2108 100644 --- a/libs/qec/unittests/test_decoders.cpp +++ b/libs/qec/unittests/test_decoders.cpp @@ -12,10 +12,18 @@ #include "cudaq/qec/pcm_utils.h" #include #include +#include +#include #include #include #include #include +#include +#if defined(__linux__) +#include +#include +#include +#endif namespace { class ScopedEnv { @@ -226,6 +234,16 @@ TEST(SampleDecoder, RealtimeApiAndDefaultGraphHooks) { EXPECT_FALSE(decoder->enqueue_syndrome(msyn.data(), msyn.size() + 1)); } +TEST(SampleDecoder, AcceptsNumaNodeId) { + std::size_t block_size = 10, syndrome_size = 5; + cudaqx::tensor H({syndrome_size, block_size}); + auto H_sparse = cudaq::qec::sparse_binary_matrix(H); + cudaqx::heterogeneous_map params; + params.insert("numa_node_id", 0); // node 0 always exists + auto d = cudaq::qec::decoder::get("sample_decoder", H_sparse, params); + ASSERT_NE(d, nullptr); +} + TEST(DecoderPlugins, SingleErrorLutExample_DecodesSingletonColumnSyndromes) { using cudaq::qec::float_t; @@ -956,6 +974,134 @@ TEST(StimDemGetDecoder, ThrowsOnProbabilityOutOfRange) { std::runtime_error); } +TEST(HardwarePinning, AccessorsReturnStoredValues) { + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("numa_node_id", 0); + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + EXPECT_EQ(d->numa_node_id(), 0); + EXPECT_EQ(d->cuda_device_id(), -1); // unset +} + +TEST(HardwarePinning, BindCurrentThreadNoopWhenUnset) { + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; // no numa_node_id + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + EXPECT_EQ(d->bind_current_thread(), -1); // nothing to bind +} + +TEST(HardwarePinning, DecodeStillWorksAfterBind) { + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + d->bind_current_thread(); // no-op at -1, but must not break decode + std::vector> batch = {{0.1f, 0.1f}}; + auto results = d->decode_batch(batch); + EXPECT_EQ(results.size(), 1u); +} + +TEST(CudaDeviceId, OutOfRangeThrowsForAnyDecoder) { + // A device id beyond the available devices is rejected during construction + // for any decoder, including CPU ones. + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device_id", 999999); + EXPECT_THROW(cudaq::qec::decoder::get("multi_error_lut", H, params), + std::runtime_error); +} + +TEST(CudaDeviceId, DecodeAsyncRestoresCallerDevice) { + // Verify decoder::get() does not corrupt the calling thread's current device. + int count = 0; + cudaGetDeviceCount(&count); + if (count < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + + cudaSetDevice(0); // caller stays on device 0 + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device_id", 1); // decoder assigned to device 1 + + // multi_error_lut is a CPU decoder -- it does nothing on device 1 itself, + // but the base class guard must still set device 1 in the async thread + // and then restore device 0 on the calling thread after get(). + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + + std::vector syn = {0.1f, 0.1f}; + auto fut = d->decode_async(syn); + fut.get(); // wait for completion + + int restored = -1; + cudaGetDevice(&restored); + EXPECT_EQ(restored, 0) + << "calling thread device not restored after decode_async"; +} + +// On a multi-GPU box: after binding the caller to device 1, decode_batch must +// NOT restore the caller to device 0 (it should leave the persistent binding +// in place). Skipped where < 2 GPUs. +TEST(HardwarePinning, DecodeBatchLeavesPersistentDeviceBinding) { + int count = 0; + cudaGetDeviceCount(&count); + if (count < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + + // Persistent bind must not leak into later tests: run on a worker thread. + std::thread t([] { + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + + d->bind_current_thread(); // caller now persistently on device 1 + std::vector> batch = {{0.1f, 0.1f}}; + (void)d->decode_batch(batch); + + int dev = -1; + cudaGetDevice(&dev); + EXPECT_EQ(dev, 1) << "decode_batch restored device despite persistent bind"; + }); + t.join(); +} + +TEST(HardwarePinning, BindCurrentThreadAppliesCpuAffinity) { + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cpu_affinity", std::vector{0}); + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + std::thread t([&] { + d->bind_current_thread(); +#if defined(__linux__) + cpu_set_t have; + CPU_ZERO(&have); + sched_getaffinity(0, sizeof(have), &have); + EXPECT_TRUE(CPU_ISSET(0, &have)); + EXPECT_EQ(CPU_COUNT(&have), 1); +#endif + }); + t.join(); +} + +// NOTE: construction (decoder::get) already runs CudaDeviceGuard ctor_dev(dev) +// before the decoder object is created/returned, so an out-of-range +// cuda_device_id throws AT CONSTRUCTION, not inside bind_current_thread(). +// bind_current_thread() is hardened the same way (see decoder.cpp) for the +// case where a decoder is constructed with cuda_device_id unset/valid and a +// caller later attempts to bind to an out-of-range device directly, but this +// test's OOB device is supplied via decoder::get()'s param_map, so the throw +// fires during get() itself. Assert the throw at the layer where it actually +// fires: construction. +TEST(HardwarePinning, BindCurrentThreadThrowsOnOutOfRangeDevice) { + int n = 0; + cudaGetDeviceCount(&n); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device_id", n + 100); // guaranteed OOB + EXPECT_THROW(cudaq::qec::decoder::get("multi_error_lut", H, params), + std::runtime_error); +} + TEST(StimDemGetDecoder, ThrowsOnMalformedStimDem) { EXPECT_THROW(cudaq::qec::get_decoder("single_error_lut", "not a valid DEM"), std::runtime_error); @@ -1190,3 +1336,396 @@ TEST(SlidingWindowDecoder, BaseStreamingCopiesFirstRoundDetectors) { << "First-round detector copy runs, but the sliding window is not full " "yet so no final correction is committed."; } + +// GPU placement: cuda_device_id must be honored across the decode entry points. +// A decoder that overrides only decode() and allocates lazily still lands on +// the assigned GPU when its owning thread is pinned via bind_current_thread(). + +TEST(HardwarePinningGpu, RawDecodeOnPinnedThreadUsesAssignedGpu) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map p; + p.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, p); + int observed = -1; + std::size_t sz = 0; + std::thread worker([&] { + d->bind_current_thread(); // owns this decoder; pins current device to 1 + auto r = d->decode(std::vector{0.1f, 0.1f}); + sz = r.result.size(); + if (sz > 1) + observed = static_cast(r.result[1]); // device inside decode() + }); + worker.join(); + ASSERT_EQ(sz, 4u); + EXPECT_EQ(observed, 1) + << "decode() on a pinned worker did not use the assigned GPU"; +} + +// Two logical patches on two GPUs: each decoder, pinned to its own GPU on its +// own worker thread, decodes on its assigned GPU with no cross-talk. +TEST(HardwarePinningGpu, MultiPatchOnPinnedThreadsUseDistinctGpus) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map p0; + p0.insert("cuda_device_id", 0); + cudaqx::heterogeneous_map p1; + p1.insert("cuda_device_id", 1); + auto d0 = cudaq::qec::decoder::get("placement_probe_decoder", H, p0); + auto d1 = cudaq::qec::decoder::get("placement_probe_decoder", H, p1); + int o0 = -1, o1 = -1; + std::thread t0([&] { + d0->bind_current_thread(); + auto r = d0->decode(std::vector{0.1f, 0.1f}); + if (r.result.size() > 1) + o0 = static_cast(r.result[1]); // device inside decode() + }); + std::thread t1([&] { + d1->bind_current_thread(); + auto r = d1->decode(std::vector{0.1f, 0.1f}); + if (r.result.size() > 1) + o1 = static_cast(r.result[1]); // device inside decode() + }); + t0.join(); + t1.join(); + EXPECT_EQ(o0, 0) << "patch 0 did not use GPU 0"; + EXPECT_EQ(o1, 1) << "patch 1 did not use GPU 1"; +} + +// decode_async applies the device guard on its worker, so it honors +// cuda_device_id. +TEST(HardwarePinningGpu, DecodeAsyncHonorsCudaDeviceId) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map p; + p.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, p); + auto fut = d->decode_async(std::vector{0.1f, 0.1f}); + auto r = fut.get(); + ASSERT_EQ(r.result.size(), 4u); + EXPECT_EQ(static_cast(r.result[1]), 1) + << "decode_async did not honor cuda_device_id"; +} + +// decode_batch applies the device guard, so it honors cuda_device_id. +TEST(HardwarePinningGpu, DecodeBatchHonorsCudaDeviceId) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map p; + p.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, p); + auto rs = d->decode_batch( + std::vector>{{0.1f, 0.1f}}); + ASSERT_EQ(rs.size(), 1u); + ASSERT_EQ(rs[0].result.size(), 4u); + EXPECT_EQ(static_cast(rs[0].result[1]), 1) + << "decode_batch did not honor cuda_device_id"; +} + +TEST(DecoderAffinity, DecodeTensorAppliesGuardWithoutBindCurrentThread) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map o; + o.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, o); + cudaqx::tensor syndrome({2}); + auto result = d->decode(syndrome); + ASSERT_EQ(result.result.size(), 4u); + EXPECT_EQ(static_cast(result.result[1]), 1) + << "decode(tensor) must land on cuda_device_id even without an " + "explicit bind_current_thread() call"; +} + +// decode_on_pinned_thread spawns a worker, binds it to the decoder's assigned +// GPU/NUMA node, decodes there, and joins before returning the result. +TEST(HardwarePinningGpu, DecodeOnPinnedThreadUsesAssignedGpu) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 2) + GTEST_SKIP() << "needs >= 2 GPUs"; + cudaSetDevice(0); + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map p; + p.insert("cuda_device_id", 1); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, p); + auto r = + d->decode_on_pinned_thread(std::vector{0.1f, 0.1f}); + ASSERT_EQ(r.result.size(), 4u); + EXPECT_EQ(static_cast(r.result[1]), 1) + << "decode_on_pinned_thread did not run on the assigned GPU"; +} + +TEST(HardwarePinning, NumaDerivesFromCudaDeviceOrDegrades) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 1) + GTEST_SKIP() << "needs a GPU"; + int node = cudaq::qec::numa_node_for_cuda_device(0); + EXPECT_GE(node, -1); // real node on multi-node HW, or -1 + EXPECT_EQ(cudaq::qec::numa_node_for_cuda_device(-1), + -1); // negative device -> -1 +} +TEST(HardwarePinning, GetDecoderSoftDerivesNumaWhenOnlyCudaSet) { + int n = 0; + cudaGetDeviceCount(&n); + if (n < 1) + GTEST_SKIP() << "needs a GPU"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device_id", 0); // numa_node_id NOT set + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + EXPECT_EQ(d->numa_node_id(), + cudaq::qec::numa_node_for_cuda_device(0)); // derivation ran +} +TEST(HardwarePinning, NumaAutoDeriveUnknownIsInfoNotThrow) { + EXPECT_NO_THROW( + { EXPECT_EQ(cudaq::qec::numa_node_for_cuda_device(-1), -1); }); +} + +TEST(DecoderAffinity, ConstructionTimeGuardHonorsExplicitBindMempolicy) { +#if defined(__linux__) + int probe = -1; + if (syscall(SYS_get_mempolicy, &probe, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map o; + o.insert("numa_node_id", 0); + o.insert("mempolicy", std::string("bind")); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, o); + auto result = d->decode(std::vector{0.0, 0.0}); + ASSERT_EQ(result.result.size(), 4u); + EXPECT_EQ(static_cast(result.result[2]), MPOL_BIND) + << "decoder::get()'s construction-time guard must honor an explicit " + "mempolicy:\"bind\" knob, not default to preferred"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// decode(tensor) is a guarded entry point, but the guard must be SKIPPED on +// the exact thread that called bind_current_thread(): after the bind, the +// thread's mempolicy is deliberately reset to a sentinel (MPOL_DEFAULT); a +// skipped guard leaves the sentinel visible inside decode() and untouched +// after it, while a wrongly applied guard would overwrite it with MPOL_BIND. +TEST(DecoderAffinity, DecodeTensorSkipsGuardOnBoundThread) { +#if defined(__linux__) + int probe = -1; + if (syscall(SYS_get_mempolicy, &probe, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map o; + o.insert("numa_node_id", 0); + o.insert("mempolicy", std::string("bind")); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, o); + int mode_after_bind = -1, mode_in_decode = -1, mode_after_decode = -1; + std::size_t sz = 0; + std::thread worker([&] { // worker so the bind never leaks into other tests + d->bind_current_thread(); + syscall(SYS_get_mempolicy, &mode_after_bind, nullptr, 0UL, nullptr, 0UL); + // Sentinel distinct from the decoder's "bind" knob. + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); + cudaqx::tensor syndrome({2}); + auto r = d->decode(syndrome); + sz = r.result.size(); + if (sz == 4) + mode_in_decode = static_cast(r.result[3]); + syscall(SYS_get_mempolicy, &mode_after_decode, nullptr, 0UL, nullptr, 0UL); + }); + worker.join(); + EXPECT_EQ(mode_after_bind, MPOL_BIND) + << "bind_current_thread did not establish the requested mempolicy"; + ASSERT_EQ(sz, 4u); + EXPECT_EQ(mode_in_decode, MPOL_DEFAULT) + << "decode(tensor) re-applied the guard on the bound thread"; + EXPECT_EQ(mode_after_decode, MPOL_DEFAULT) + << "decode(tensor) changed the bound thread's placement"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// A failed bind_current_thread (valid numa_node_id, out-of-range cpu_affinity +// core id) must throw std::invalid_argument and roll back every side effect: +// mempolicy mode and CPU affinity mask are restored, and the decoder is NOT +// left bound -- a subsequent guarded decode still applies the guard. +TEST(DecoderAffinity, FailedBindThrowsAndFullyRestoresThread) { +#if defined(__linux__) + int probe = -1; + if (syscall(SYS_get_mempolicy, &probe, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map o; + o.insert("numa_node_id", 0); + o.insert("mempolicy", std::string("bind")); + o.insert("cpu_affinity", std::vector{CPU_SETSIZE + 10}); // out of range + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, o); + int mode_before = -1, mode_after = -2; + bool threw_invalid = false, mask_restored = false; + int mode_in_decode = -1; + std::size_t sz = 0; + std::thread worker([&] { + syscall(SYS_get_mempolicy, &mode_before, nullptr, 0UL, nullptr, 0UL); + cpu_set_t before, after; + CPU_ZERO(&before); + CPU_ZERO(&after); + sched_getaffinity(0, sizeof(before), &before); + try { + d->bind_current_thread(); + } catch (const std::invalid_argument &) { + threw_invalid = true; + } + syscall(SYS_get_mempolicy, &mode_after, nullptr, 0UL, nullptr, 0UL); + sched_getaffinity(0, sizeof(after), &after); + mask_restored = CPU_EQUAL(&before, &after); + // Not bound => the guard must still fire on this thread's guarded decode. + cudaqx::tensor syndrome({2}); + auto r = d->decode(syndrome); + sz = r.result.size(); + if (sz == 4) + mode_in_decode = static_cast(r.result[3]); + }); + worker.join(); + EXPECT_TRUE(threw_invalid) + << "out-of-range cpu_affinity must throw std::invalid_argument"; + EXPECT_EQ(mode_after, mode_before) << "failed bind leaked a mempolicy change"; + EXPECT_TRUE(mask_restored) << "failed bind leaked a CPU-affinity change"; + ASSERT_EQ(sz, 4u); + EXPECT_EQ(mode_in_decode, MPOL_BIND) + << "decoder was left bound after a failed bind_current_thread: the " + "guard did not apply mempolicy on a later decode"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// decode_on_pinned_thread must not steal the caller's binding: after a +// bind_current_thread() + decode_on_pinned_thread() sequence, a decode_batch() +// from the caller still skips the guard (the sentinel MPOL_DEFAULT set after +// the bind stays visible inside decode instead of the guard's MPOL_BIND). +TEST(DecoderAffinity, PinnedThreadDecodePreservesCallersBinding) { +#if defined(__linux__) + int probe = -1; + if (syscall(SYS_get_mempolicy, &probe, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map o; + o.insert("numa_node_id", 0); + o.insert("mempolicy", std::string("bind")); + auto d = cudaq::qec::decoder::get("placement_probe_decoder", H, o); + int pinned_mode = -1, batch_mode = -1; + std::size_t pinned_sz = 0, batch_n = 0, batch_sz = 0; + std::thread caller([&] { + d->bind_current_thread(); + // Sentinel on the caller thread; an applied guard would report MPOL_BIND. + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); + auto rp = d->decode_on_pinned_thread( + std::vector{0.0f, 0.0f}); + pinned_sz = rp.result.size(); + if (pinned_sz == 4) + pinned_mode = static_cast(rp.result[3]); // pinned worker bound + auto rb = d->decode_batch( + std::vector>{{0.0f, 0.0f}}); + batch_n = rb.size(); + if (batch_n == 1) { + batch_sz = rb[0].result.size(); + if (batch_sz == 4) + batch_mode = static_cast(rb[0].result[3]); + } + }); + caller.join(); + ASSERT_EQ(pinned_sz, 4u); + EXPECT_EQ(pinned_mode, MPOL_BIND) + << "decode_on_pinned_thread's worker was not bound to the decoder's " + "mempolicy"; + ASSERT_EQ(batch_n, 1u); + ASSERT_EQ(batch_sz, 4u); + EXPECT_EQ(batch_mode, MPOL_DEFAULT) + << "caller's binding did not survive decode_on_pinned_thread: " + "decode_batch re-applied the guard on the bound caller"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// ---- Knob-contract negative tests -------------------------------------- + +// T1.1: a typo'd knob key must not accidentally pin anything -- the tolerant +// plugin ignores unknown keys, so both affinity knobs stay unset and decoding +// still works. +TEST(HardwarePinningNegative, TypoedKnobKeyIsSilentlyIgnoredByTolerantPlugin) { + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("cuda_device", 1); // typo: should be cuda_device_id + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + EXPECT_EQ(d->cuda_device_id(), -1) << "typo must not accidentally pin"; + EXPECT_EQ(d->numa_node_id(), -1); + std::vector> batch = {{0.1f, 0.1f}}; + EXPECT_NO_THROW((void)d->decode_batch(batch)); +} + +// T1.7: silence contract -- users who never touch the pinning knobs must never +// see an affinity warning. +TEST(HardwarePinningNegative, UnpinnedUsersSeeZeroAffinityWarnings) { + cudaqx::tensor H({2, 3}); + testing::internal::CaptureStderr(); + auto d = cudaq::qec::decoder::get("multi_error_lut", H, + cudaqx::heterogeneous_map{}); + std::vector> batch = {{0.1f, 0.1f}}; + (void)d->decode_batch(batch); + std::string err = testing::internal::GetCapturedStderr(); + EXPECT_EQ(err.find("[cudaq-qec affinity]"), std::string::npos) + << "knobless usage must be warning-silent, got: " << err; +} + +// T2.12: a numa_node_id that is encodable (< 64) but does not exist on this +// host must degrade loudly (warning naming the node) and still decode +// correctly, unpinned. +#if defined(__linux__) +TEST(HardwarePinningNegative, NonexistentNodeWarnsAndRunsUnpinned) { + // Pick a node id in [0,64) that does not exist on this host. + int missing = -1; + for (int n = 8; n < 64; ++n) { + std::ifstream f("/sys/devices/system/node/node" + std::to_string(n) + + "/cpulist"); + if (!f.is_open()) { + missing = n; + break; + } + } + if (missing < 0) + GTEST_SKIP() << "every node id in [8,64) exists on this host"; + cudaqx::tensor H({2, 3}); + cudaqx::heterogeneous_map params; + params.insert("numa_node_id", missing); + testing::internal::CaptureStderr(); + auto d = cudaq::qec::decoder::get("multi_error_lut", H, params); + std::vector> batch = {{0.1f, 0.1f}}; + auto results = d->decode_batch(batch); + std::string err = testing::internal::GetCapturedStderr(); + ASSERT_EQ(results.size(), batch.size()); + EXPECT_TRUE(results[0].converged) << "must degrade, not corrupt"; + EXPECT_NE(err.find("numa_node_id " + std::to_string(missing)), + std::string::npos) + << "degradation must be loud"; +} +#endif diff --git a/libs/qec/unittests/test_decoders_yaml.cpp b/libs/qec/unittests/test_decoders_yaml.cpp index dc8be3e41..2896e634a 100644 --- a/libs/qec/unittests/test_decoders_yaml.cpp +++ b/libs/qec/unittests/test_decoders_yaml.cpp @@ -720,3 +720,73 @@ TEST(DecoderConfigTest, SimulationHostPointerWrappersForwardToHostRuntime) { EXPECT_EQ(corrections, (std::vector{0})); finalize_decoders(); } + +TEST(DecoderYaml, PinFieldsRoundTrip) { + using namespace cudaq::qec::decoding::config; + multi_decoder_config multi_config; + decoder_config cfg; + cfg.id = 0; + cfg.type = "multi_error_lut"; + cfg.block_size = 4; + cfg.syndrome_size = 2; + cfg.H_sparse = {0, -1, 1, -1}; + cfg.O_sparse = {}; + cfg.D_sparse = {}; + cfg.cuda_device_id = 1; + cfg.numa_node_id = 0; + multi_config.decoders.push_back(cfg); + + test_decoder_yaml_roundtrip(multi_config); + + auto reparsed = + multi_decoder_config::from_yaml_str(multi_config.to_yaml_str(200)); + ASSERT_EQ(reparsed.decoders.size(), 1u); + EXPECT_EQ(reparsed.decoders[0].cuda_device_id, 1); + EXPECT_EQ(reparsed.decoders[0].numa_node_id, 0); +} + +TEST(DecoderYaml, PrepareParamsInjectsKeys) { + using namespace cudaq::qec::decoding::config; + decoder_config cfg; + cfg.id = 0; + cfg.type = "multi_error_lut"; + cfg.block_size = 4; + cfg.syndrome_size = 2; + cfg.cuda_device_id = 1; + cfg.numa_node_id = 0; + auto params = cudaq::qec::decoding::host::prepare_decoder_params(cfg); + EXPECT_TRUE(params.contains("cuda_device_id")); + EXPECT_TRUE(params.contains("numa_node_id")); +} + +TEST(DecoderYaml, MempolicyAndCpuAffinityRoundTrip) { + using namespace cudaq::qec::decoding::config; + multi_decoder_config multi_config; + decoder_config cfg; + cfg.id = 0; + cfg.type = "multi_error_lut"; + cfg.block_size = 4; + cfg.syndrome_size = 2; + cfg.H_sparse = {0, -1, 1, -1}; + cfg.O_sparse = {}; + cfg.D_sparse = {}; + cfg.mempolicy = "bind"; + cfg.cpu_affinity = std::vector{0, 2}; + multi_config.decoders.push_back(cfg); + + test_decoder_yaml_roundtrip(multi_config); + + auto reparsed = + multi_decoder_config::from_yaml_str(multi_config.to_yaml_str(200)); + ASSERT_EQ(reparsed.decoders.size(), 1u); + const auto &dc = reparsed.decoders.at(0); + ASSERT_TRUE(dc.mempolicy.has_value()); + EXPECT_EQ(*dc.mempolicy, "bind"); + ASSERT_TRUE(dc.cpu_affinity.has_value()); + EXPECT_EQ(*dc.cpu_affinity, (std::vector{0, 2})); + + auto params = cudaq::qec::decoding::host::prepare_decoder_params(dc); + EXPECT_EQ(cudaq::qec::read_mempolicy(params), + cudaq::qec::mempolicy_mode::bind); + EXPECT_EQ(cudaq::qec::read_cpu_affinity(params), (std::vector{0, 2})); +} diff --git a/libs/qec/unittests/test_device_affinity.cpp b/libs/qec/unittests/test_device_affinity.cpp new file mode 100644 index 000000000..2e2208812 --- /dev/null +++ b/libs/qec/unittests/test_device_affinity.cpp @@ -0,0 +1,211 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ + +#include "cuda-qx/core/heterogeneous_map.h" +#include "cudaq/qec/device_affinity.h" +#include +#include +#include + +#if defined(__linux__) +#include "hardware_affinity.h" +#include +#include +#include +#endif +#include + +using cudaq::qec::read_cuda_device_id; +using cudaq::qec::read_numa_node_id; +using cudaqx::heterogeneous_map; + +TEST(DeviceAffinity, ReadAbsentReturnsMinusOne) { + heterogeneous_map m; + EXPECT_EQ(read_cuda_device_id(m), -1); + EXPECT_EQ(read_numa_node_id(m), -1); +} + +TEST(DeviceAffinity, ReadIntStorage) { // YAML path: std::optional -> int + heterogeneous_map m; + m.insert("cuda_device_id", 3); + m.insert("numa_node_id", 1); + EXPECT_EQ(read_cuda_device_id(m), 3); + EXPECT_EQ(read_numa_node_id(m), 1); +} + +TEST(DeviceAffinity, + ReadSizeTStorage) { // kwargs path: Python int -> std::size_t + heterogeneous_map m; + m.insert("cuda_device_id", std::size_t{2}); + EXPECT_EQ(read_cuda_device_id(m), 2); +} + +TEST(DeviceAffinity, ReadNegativeThrows) { + heterogeneous_map m; + m.insert("cuda_device_id", -1); + EXPECT_THROW(read_cuda_device_id(m), std::runtime_error); +} + +TEST(DeviceAffinity, ReadCpuAffinityAcceptsDoubleStorage) { + // Python kwargs deliver lists as vector; integral values convert. + heterogeneous_map m; + m.insert("cpu_affinity", std::vector{0.0, 2.0}); + EXPECT_EQ(cudaq::qec::read_cpu_affinity(m), (std::vector{0, 2})); + heterogeneous_map bad; + bad.insert("cpu_affinity", std::vector{0.5}); + EXPECT_THROW(cudaq::qec::read_cpu_affinity(bad), std::runtime_error); +} + +#if defined(__linux__) +// Container runtimes commonly block the mempolicy syscalls (seccomp without +// CAP_SYS_NICE). The library degrades gracefully there (warns, keeps going), +// but tests that assert the syscalls' effects must skip instead of fail. +static bool mempolicy_syscalls_usable() { + int mode = -1; + if (syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) != 0) + return false; + return syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL) == 0; +} + +TEST(HardwareAffinity, MempolicyDefaultIsPreferredNotBind) { + namespace da = cudaq::qec::detail_affinity; + if (!mempolicy_syscalls_usable()) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cpu_set_t saved; + CPU_ZERO(&saved); + sched_getaffinity(0, sizeof(saved), &saved); + da::bind_this_thread_to_numa_node(0); + EXPECT_EQ(da::current_thread_mempolicy_mode(), MPOL_PREFERRED); + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); + sched_setaffinity(0, sizeof(saved), &saved); +} +TEST(HardwareAffinity, MempolicyBindWhenRequested) { + namespace da = cudaq::qec::detail_affinity; + if (!mempolicy_syscalls_usable()) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cpu_set_t saved; + CPU_ZERO(&saved); + sched_getaffinity(0, sizeof(saved), &saved); + da::bind_this_thread_to_numa_node(0, cudaq::qec::mempolicy_mode::bind); + EXPECT_EQ(da::current_thread_mempolicy_mode(), MPOL_BIND); + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); + sched_setaffinity(0, sizeof(saved), &saved); +} +TEST(HardwareAffinity, MempolicyRejectsUnknownStringThrows) { + cudaqx::heterogeneous_map p; + p.insert("mempolicy", std::string("bnid")); + EXPECT_THROW(cudaq::qec::read_mempolicy(p), std::runtime_error); +} +TEST(HardwareAffinity, NumaNode64ThrowsOnMempolicyBind) { + namespace da = cudaq::qec::detail_affinity; + EXPECT_THROW(da::bind_this_thread_to_numa_node(64), std::runtime_error); +} +TEST(HardwareAffinity, CpuAffinityPinsToExactCores) { + namespace da = cudaq::qec::detail_affinity; + cpu_set_t saved; + CPU_ZERO(&saved); + ASSERT_EQ(sched_getaffinity(0, sizeof(saved), &saved), 0); + // sched_setaffinity succeeds with the INTERSECTION of the requested and + // allowed masks; assert exact placement only where both cores are allowed. + if (!CPU_ISSET(0, &saved) || !CPU_ISSET(2, &saved)) + GTEST_SKIP() << "CPUs 0 and 2 not both in this process's allowed cpuset"; + da::set_thread_cpu_affinity({0, 2}); + auto cpus = da::current_thread_cpuset(); + EXPECT_EQ(cpus, (std::vector{0, 2})); + sched_setaffinity(0, sizeof(saved), &saved); +} +TEST(HardwareAffinity, CpuAffinityEmptyIsNoop) { + namespace da = cudaq::qec::detail_affinity; + cpu_set_t before; + CPU_ZERO(&before); + sched_getaffinity(0, sizeof(before), &before); + da::set_thread_cpu_affinity({}); + cpu_set_t after; + CPU_ZERO(&after); + sched_getaffinity(0, sizeof(after), &after); + EXPECT_TRUE(CPU_EQUAL(&before, &after)); +} +TEST(HardwareAffinity, CpuAffinityOutOfRangeCoreThrows) { + namespace da = cudaq::qec::detail_affinity; + EXPECT_THROW(da::set_thread_cpu_affinity({CPU_SETSIZE}), + std::invalid_argument); + EXPECT_THROW(da::set_thread_cpu_affinity({-1}), std::invalid_argument); +} +TEST(HardwareAffinity, BindThreadPinsAffinityToNodeCpus) { + namespace da = cudaq::qec::detail_affinity; + cpu_set_t node0; + CPU_ZERO(&node0); + if (!da::build_node_cpuset(0, node0)) + GTEST_SKIP() << "no node0 cpulist"; + cpu_set_t saved; + CPU_ZERO(&saved); + sched_getaffinity(0, sizeof(saved), &saved); + da::bind_this_thread_to_numa_node(0); + cpu_set_t have; + CPU_ZERO(&have); + sched_getaffinity(0, sizeof(have), &have); + for (int c = 0; c < CPU_SETSIZE; ++c) + if (CPU_ISSET(c, &have)) + EXPECT_TRUE(CPU_ISSET(c, &node0)) << "cpu " << c << " not on node 0"; + sched_setaffinity(0, sizeof(saved), &saved); + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); +} +TEST(HardwareAffinity, NegativeNodeIsNoop) { + namespace da = cudaq::qec::detail_affinity; + if (!mempolicy_syscalls_usable()) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + cpu_set_t before; + CPU_ZERO(&before); + sched_getaffinity(0, sizeof(before), &before); + da::bind_this_thread_to_numa_node(-1); + cpu_set_t after; + CPU_ZERO(&after); + sched_getaffinity(0, sizeof(after), &after); + EXPECT_TRUE(CPU_EQUAL(&before, &after)); + EXPECT_EQ(da::current_thread_mempolicy_mode(), MPOL_DEFAULT); +} +TEST(HardwareAffinity, BindRegionSetsPolicyOnBuffer) { + namespace da = cudaq::qec::detail_affinity; + const long page = sysconf(_SC_PAGESIZE); + void *p = std::aligned_alloc(page, static_cast(page)); + ASSERT_NE(p, nullptr); + std::memset(p, 0, static_cast(page)); + da::bind_region_to_numa_node(p, static_cast(page), 0); // preferred + int mode = -1; + long rc = syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, p, MPOL_F_ADDR); + if (rc != 0) { + std::free(p); + GTEST_SKIP() << "get_mempolicy(MPOL_F_ADDR) unavailable (seccomp?)"; + } + EXPECT_EQ(mode, MPOL_PREFERRED) + << "mbind on a page-aligned region must set MPOL_PREFERRED"; + da::bind_region_to_numa_node(p, static_cast(page), -1); // no-op ok + std::free(p); +} +TEST(HardwareAffinity, BindRegionNode64Throws) { + namespace da = cudaq::qec::detail_affinity; + void *p = std::calloc(1, 4096); + ASSERT_NE(p, nullptr); + EXPECT_THROW(da::bind_region_to_numa_node(p, 4096, 64), std::runtime_error); + std::free(p); +} +#endif + +// T1.2: a correctly named knob with the WRONG value type must throw, and the +// error must name the offending key so the user can find it. +TEST(DeviceAffinity, WrongTypedKnobThrowsNamingTheKey) { + heterogeneous_map m; + m.insert("cuda_device_id", std::string("0")); + try { + (void)read_cuda_device_id(m); + FAIL() << "expected a throw for string-typed cuda_device_id"; + } catch (const std::exception &e) { + EXPECT_NE(std::string(e.what()).find("cuda_device_id"), std::string::npos) + << "error must name the key; got: " << e.what(); + } +} diff --git a/libs/qec/unittests/test_pinning_benchmark.cpp b/libs/qec/unittests/test_pinning_benchmark.cpp new file mode 100644 index 000000000..aa0f20123 --- /dev/null +++ b/libs/qec/unittests/test_pinning_benchmark.cpp @@ -0,0 +1,839 @@ +/******************************************************************************* + * Copyright (c) 2025 - 2026 NVIDIA Corporation & Affiliates. * + * All rights reserved. * + * * + * This source code and the accompanying materials are made available under * + * the terms of the Apache License 2.0 which accompanies this distribution. * + ******************************************************************************/ +#include "support/thread_placement.h" +#include "cudaq/qec/decoder.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#if defined(__linux__) +#include +#include +#include +#endif + +// Resolved from the LD_PRELOAD shim when preloaded; weak so the binary still +// links (and the test cleanly skips) when it is not. +extern "C" __attribute__((weak)) long cudaqx_affinity_syscall_count(); +extern "C" __attribute__((weak)) void cudaqx_affinity_syscall_reset(); +// Per-syscall counters (same shim). +extern "C" __attribute__((weak)) long cudaqx_sched_setaffinity_count(); +extern "C" __attribute__((weak)) long cudaqx_sched_getaffinity_count(); +extern "C" __attribute__((weak)) long cudaqx_set_mempolicy_count(); +extern "C" __attribute__((weak)) long cudaqx_get_mempolicy_count(); +extern "C" __attribute__((weak)) long cudaqx_mbind_count(); +extern "C" __attribute__((weak)) long cudaqx_node_sysfs_open_count(); + +namespace { +cudaqx::tensor makeH() { + cudaqx::tensor H({2, 3}); + return H; +} +// One decoder with numa_node_id pinned so the decode-time guard's syscalls fire +// on the unbound path. +std::unique_ptr makePinnedLut() { + cudaqx::heterogeneous_map o; + o.insert("numa_node_id", 0); + return cudaq::qec::decoder::get("multi_error_lut", makeH(), o); +} +const std::vector> kChunk = {{0.1, 0.1}, + {0.2, 0.2}}; + +// ---- invariant-test helpers (shim counters + placement snapshots) ---------- +using cudaq::qec::test::thread_placement; + +bool shimCountersPresent() { + return cudaqx_affinity_syscall_reset && cudaqx_sched_setaffinity_count && + cudaqx_sched_getaffinity_count && cudaqx_set_mempolicy_count && + cudaqx_get_mempolicy_count; +} +struct ShimCounts { + long setaff, getaff, setmem, getmem; +}; +ShimCounts readShimCounts() { + return {cudaqx_sched_setaffinity_count(), cudaqx_sched_getaffinity_count(), + cudaqx_set_mempolicy_count(), cudaqx_get_mempolicy_count()}; +} + +#if defined(__linux__) +// Container seccomp commonly blocks the mempolicy syscalls. When +// SYS_get_mempolicy is blocked, NumaGuard cannot capture the prior policy and +// correctly skips set_mempolicy entirely, so tests asserting setmem > 0 must +// skip, not fail (mirrors test_device_affinity.cpp). Read-only: no side +// effects on the thread's policy. +bool get_mempolicy_usable() { + int mode = -1; + return syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) == 0; +} +#else +bool get_mempolicy_usable() { return false; } +#endif + +// bind_current_thread() is persistent (no restore). Run test bodies that bind +// on a disposable worker so the gtest main thread's placement stays intact +// for later tests. EXPECT_* record correctly from any thread. +template +void run_on_worker(F &&fn) { + std::exception_ptr err; + std::thread t([&] { + try { + fn(); + } catch (...) { + err = std::current_exception(); + } + }); + t.join(); + if (err) + std::rethrow_exception(err); +} + +// Pinned decoder with D/O set so ALL entry points (incl. enqueue_syndrome) +// are callable. +std::unique_ptr makeEntryReadyPinnedLut() { + auto d = makePinnedLut(); + d->set_D_sparse(std::vector>{{0}, {1}}); + d->set_O_sparse(std::vector>{{0, 1}}); + return d; +} + +// Every guarded entry point whose CUDA/NUMA guard runs on the CALLING thread +// (decode_async runs its guard on a worker thread; tested separately). +struct GuardedEntryPoint { + const char *name; + void (*call)(cudaq::qec::decoder &); +}; +const GuardedEntryPoint kCallerGuardedEntryPoints[] = { + {"decode_batch", [](cudaq::qec::decoder &d) { d.decode_batch(kChunk); }}, + {"decode_tensor", + [](cudaq::qec::decoder &d) { + cudaqx::tensor syndrome({2}); // zero-initialized, rank-1 + d.decode(syndrome); + }}, + {"enqueue_syndrome", + [](cudaq::qec::decoder &d) { + std::vector syndrome = {1, 1}; + d.enqueue_syndrome(syndrome); + }}, +}; +void callDecodeAsync(cudaq::qec::decoder &d) { + d.decode_async(kChunk[0]).get(); +} +} // namespace + +// Deterministic gate: after binding, a decode loop must issue ZERO +// sched_setaffinity calls; the unbound path (canary) must issue > 0 (proving +// the counter is wired and the guard really fires when not bound). +TEST(PinningBenchmark, BoundDecodeLoopIssuesNoAffinitySyscalls) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker( + [] { + auto d = makePinnedLut(); + constexpr int N = 200; + + // Canary: unbound -> per-call guard fires. + cudaqx_affinity_syscall_reset(); + for (int i = 0; i < N; ++i) + d->decode_batch(kChunk); + long unbound = cudaqx_affinity_syscall_count(); + EXPECT_GT(unbound, 0) + << "unbound decode loop should issue affinity syscalls " + "(else the gate is vacuous)"; + + // Bound: bind once, then the loop must add none. + d->bind_current_thread(); + cudaqx_affinity_syscall_reset(); + for (int i = 0; i < N; ++i) + d->decode_batch(kChunk); + long bound = cudaqx_affinity_syscall_count(); + EXPECT_EQ(bound, 0) + << "bound decode loop must not re-issue affinity syscalls"; + }); +} + +TEST(PinningBenchmark, BindOnOneThreadStillGuardsAnotherThread) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + d->bind_current_thread(); // bind on the binding worker thread + + long from_other_thread = -1; + std::thread other([&] { + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + from_other_thread = cudaqx_affinity_syscall_count(); + }); + other.join(); + EXPECT_GT(from_other_thread, 0) + << "a thread that never called bind_current_thread() must still be " + "guarded, even though a DIFFERENT thread bound this decoder"; + }); +} + +// Guards the decode_on_pinned_thread() contract: its one-shot worker thread +// binds itself for the duration of a single decode, but that binding must +// not stick to the decoder object once the worker exits. A subsequent +// decode_batch() from the (unbound) main thread must still be guarded. +TEST(PinningBenchmark, OneShotPinnedDecodeDoesNotStickToObject) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + d->decode_on_pinned_thread(kChunk[0]); + + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + EXPECT_GT(cudaqx_affinity_syscall_count(), 0) + << "decode_on_pinned_thread()'s one-shot worker binding must not " + "outlive it; a later decode_batch() from the main thread must " + "still issue affinity syscalls"; +} + +// Verify that decode_on_pinned_thread() does not clobber a concurrent +// bind_current_thread() that runs from a third thread. +TEST(PinningBenchmark, OneShotDoesNotClobberConcurrentBind) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + + // Thread B will call bind_current_thread() while the one-shot worker runs. + // It stays alive until the unbound-thread check below completes: if it + // exited first, the fresh "unbound" thread could recycle its pthread id, + // spuriously pass is_bound_here(), and skip the guard (observed CI flake). + std::atomic b_bound{false}; + std::atomic b_done{false}; + std::thread thread_b([&] { + // Spin until main_worker signals us to bind. + while (!b_bound.load(std::memory_order_acquire)) + std::this_thread::yield(); + d->bind_current_thread(); + while (!b_done.load(std::memory_order_acquire)) + std::this_thread::yield(); + }); + + // Run decode_on_pinned_thread from the main thread; signal B to bind midway. + std::thread main_worker([&] { + // Signal B just before decode so the race window is as wide as possible. + b_bound.store(true, std::memory_order_release); + d->decode_on_pinned_thread(kChunk[0]); + }); + + main_worker.join(); + + // Whoever bound last should have their binding intact. At a minimum, + // thread_b's binding must not permanently suppress guards for other + // threads: a decode from a thread that was never bound must still issue + // syscalls. + cudaqx_affinity_syscall_reset(); + // Run decode from a fresh thread that was never bound (thread_b is still + // alive, so this thread cannot alias its id). + std::thread unbound([&] { d->decode_batch(kChunk); }); + unbound.join(); + // An unbound thread must always fire the guard. + EXPECT_GT(cudaqx_affinity_syscall_count(), 0) + << "guard must still fire for unbound threads after concurrent-bind race"; + + b_done.store(true, std::memory_order_release); + thread_b.join(); +} + +TEST(PinningBenchmark, EnqueueSyndromeAppliesGuardOnUnboundThread) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + d->set_D_sparse(std::vector>{{0}, {1}}); + d->set_O_sparse(std::vector>{{0, 1}}); + cudaqx_affinity_syscall_reset(); + std::vector syndrome = {1, 1}; + d->enqueue_syndrome(syndrome); + EXPECT_GT(cudaqx_affinity_syscall_count(), 0) + << "enqueue_syndrome's internal decode() call must apply the NUMA " + "guard on an unbound thread, same as decode_batch()"; +} + +// Loose A/B (report + gross-regression guard only; timing is noisy). +TEST(PinningBenchmark, BoundThroughputNotWorseThanUnbound) { + run_on_worker([] { + auto d = makePinnedLut(); + constexpr int N = 5000; + auto run = [&] { + auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < N; ++i) + d->decode_batch(kChunk); + return std::chrono::duration( + std::chrono::steady_clock::now() - t0) + .count(); + }; + run(); // warm up + double unbound_us = run(); + d->bind_current_thread(); + run(); // warm up bound + double bound_us = run(); + std::printf( + "[pinning A/B] unbound=%.1fus bound=%.1fus (%d decodes) ratio=%.2f\n", + unbound_us, bound_us, N, bound_us / unbound_us); + // Pinning removes per-call guard work; it must not be materially slower. + EXPECT_LT(bound_us, unbound_us * 1.5); + }); +} + +TEST(PinningBenchmark, DecodeBatchRestoresExactPriorMempolicy) { +#if defined(__linux__) + int mode = -1; + if (syscall(SYS_get_mempolicy, &mode, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + unsigned long nodemask = 1UL; // node 0 + if (syscall(SYS_set_mempolicy, MPOL_BIND, &nodemask, sizeof(nodemask) * 8) != + 0) + GTEST_SKIP() << "set_mempolicy unavailable (container seccomp?)"; + auto d = makePinnedLut(); // numa_node_id=0, unbound (per-call guard runs) + d->decode_batch(kChunk); + int mode_after = -1; + syscall(SYS_get_mempolicy, &mode_after, nullptr, 0UL, nullptr, 0UL); + syscall(SYS_set_mempolicy, MPOL_DEFAULT, nullptr, 0UL); // cleanup regardless + EXPECT_EQ(mode_after, MPOL_BIND) + << "decode_batch's guard must restore the exact prior mempolicy " + "(MPOL_BIND), not force MPOL_DEFAULT"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// The un-restorable-policy scenario: a seccomp profile that blocks +// get_mempolicy while allowing set_mempolicy. A temporary guard that cannot +// capture the prior policy must not change it at all — it could never restore +// it, so the change would silently outlive the decode. The shim's +// CUDAQX_SHIM_FAIL_GET_MEMPOLICY injection simulates the blocked capture. +TEST(PinningInvariants, GuardSkipsMempolicyWhenCaptureFails) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; +#if defined(__linux__) + // Baseline read runs with injection off. + int before = -1; + if (syscall(SYS_get_mempolicy, &before, nullptr, 0UL, nullptr, 0UL) != 0) + GTEST_SKIP() << "mempolicy syscalls unavailable (container seccomp?)"; + + auto d = makePinnedLut(); // numa_node_id=0, unbound -> guard runs per call + setenv("CUDAQX_SHIM_FAIL_GET_MEMPOLICY", "1", 1); + d->decode_batch(kChunk); + unsetenv("CUDAQX_SHIM_FAIL_GET_MEMPOLICY"); + + int after = -1; + ASSERT_EQ(syscall(SYS_get_mempolicy, &after, nullptr, 0UL, nullptr, 0UL), 0); + EXPECT_EQ(after, before) + << "the guard applied a mempolicy it could not capture; the policy " + "leaked past the decode (capture-failure must skip set_mempolicy)"; +#else + GTEST_SKIP() << "Linux-only"; +#endif +} + +// Bound-path invariant, per entry point: after bind_current_thread() each +// caller-thread entry point must issue ZERO placement syscalls (all four +// counters) and leave the calling thread's placement bit-identical. +TEST(PinningInvariants, BoundEntryPointsIssueNoPlacementSyscalls) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + for (const auto &ep : kCallerGuardedEntryPoints) { + SCOPED_TRACE(ep.name); + auto d = makeEntryReadyPinnedLut(); + d->bind_current_thread(); + auto before = thread_placement::capture(); + cudaqx_affinity_syscall_reset(); + ep.call(*d); + auto c = readShimCounts(); + EXPECT_EQ(c.setaff, 0) << "bound path issued sched_setaffinity"; + EXPECT_EQ(c.getaff, 0) << "bound path issued sched_getaffinity"; + EXPECT_EQ(c.setmem, 0) << "bound path issued set_mempolicy"; + EXPECT_EQ(c.getmem, 0) << "bound path issued get_mempolicy"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) << "bound path changed placement: " + << before.describe_difference(after); + } + }); +} + +// decode_async decodes on a fresh std::async worker which can never be the +// bound thread, so its guard runs on the WORKER by design (asserted >0 here: +// deleting that guard fails this test). The caller-side invariant is that the +// calling thread's placement is untouched even though it holds the binding. +TEST(PinningInvariants, BoundDecodeAsyncGuardsWorkerAndLeavesCallerPlacement) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + if (!get_mempolicy_usable()) + GTEST_SKIP() << "get_mempolicy blocked (container seccomp?); the guard " + "correctly skips set_mempolicy then, so setmem > 0 " + "cannot hold"; + run_on_worker([] { + auto d = makeEntryReadyPinnedLut(); + d->bind_current_thread(); + auto before = thread_placement::capture(); + cudaqx_affinity_syscall_reset(); + callDecodeAsync(*d); + EXPECT_GT(readShimCounts().setmem, 0) + << "decode_async's worker-thread guard must still attempt " + "set_mempolicy (the caller's binding must not leak to the worker)"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) << "decode_async changed the CALLING thread's " + "placement: " + << before.describe_difference(after); + }); +} + +// Unbound-path invariant, per entry point (incl. decode_async): the guard must +// actually run (set_mempolicy attempted -- deleting the guard fails this) AND +// fully restore the calling thread's placement afterwards. +TEST(PinningInvariants, UnboundEntryPointsApplyAndFullyRestoreGuard) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + if (!get_mempolicy_usable()) + GTEST_SKIP() << "get_mempolicy blocked (container seccomp?); the guard " + "correctly skips set_mempolicy then, so setmem > 0 " + "cannot hold"; + std::vector entries(std::begin(kCallerGuardedEntryPoints), + std::end(kCallerGuardedEntryPoints)); + entries.push_back({"decode_async", callDecodeAsync}); + for (const auto &ep : entries) { + SCOPED_TRACE(ep.name); + auto d = makeEntryReadyPinnedLut(); // numa_node_id=0, never bound + auto before = thread_placement::capture(); + cudaqx_affinity_syscall_reset(); + ep.call(*d); + EXPECT_GT(readShimCounts().setmem, 0) + << "unbound path must attempt set_mempolicy (guard deleted?)"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) << "guard applied but not fully restored: " + << before.describe_difference(after); + } +} + +// The tensor-decode entry must construct its guards BEFORE building the +// soft-syndrome temporaries, so those allocations land on the decoder's node. +// Observable: with the shim, the FIRST placement syscall must occur before +// any decode work — assert the guard fires even for a rank-check failure, +// which returns before temporaries are built. +TEST(PinningInvariants, TensorDecodeGuardPrecedesTemporaries) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makeEntryReadyPinnedLut(); + cudaqx::tensor bad({2, 2}); // rank-2: must throw + cudaqx_affinity_syscall_reset(); + EXPECT_THROW((void)d->decode(bad), std::runtime_error); + EXPECT_GT(readShimCounts().getaff + readShimCounts().getmem, 0) + << "guard must be constructed before input processing"; + }); +} + +// Post-construction sparse setters allocate session-lifetime buffers +// (corrections, msyn accumulators); those buffers must follow the decoder's +// placement like every other allocating entry point, and the guard must +// restore the calling thread's placement afterwards. Both overload shapes of +// each setter allocate independently, so each is exercised on its own. +TEST(PinningInvariants, SparseSettersApplyNumaGuardWhenUnbound) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + struct SetterCall { + const char *name; + void (*call)(cudaq::qec::decoder &); + }; + const SetterCall kSetters[] = { + {"set_D_sparse(nested)", + [](cudaq::qec::decoder &d) { + d.set_D_sparse(std::vector>{{0}, {1}}); + }}, + {"set_D_sparse(flat)", + [](cudaq::qec::decoder &d) { + // -1 terminates a row: {{0}, {1}}. + d.set_D_sparse(std::vector{0, -1, 1}); + }}, + {"set_O_sparse(nested)", + [](cudaq::qec::decoder &d) { + d.set_O_sparse(std::vector>{{0, 1}}); + }}, + {"set_O_sparse(flat)", + [](cudaq::qec::decoder &d) { + d.set_O_sparse(std::vector{0, 1}); + }}, + }; + run_on_worker([&] { + auto d = makePinnedLut(); // numa_node_id=0, unbound -> guard must run + for (const auto &s : kSetters) { + SCOPED_TRACE(s.name); + auto before = thread_placement::capture(); + cudaqx_affinity_syscall_reset(); + s.call(*d); + auto c = readShimCounts(); + EXPECT_GT(c.setmem + c.setaff, 0) + << "sparse setters must apply the placement guard"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) << "setter guard must restore placement: " + << before.describe_difference(after); + } + }); +} + +// Non-virtual guarded single-syndrome entry (used by the Python binding — +// callers that have not bound a thread must self-place). +TEST(PinningInvariants, DecodeGuardedAppliesAndRestores) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + auto before = thread_placement::capture(); + cudaqx_affinity_syscall_reset(); + (void)d->decode_guarded({0.1f, 0.1f}); + EXPECT_GT(readShimCounts().setmem + readShimCounts().setaff, 0); + EXPECT_EQ(before, thread_placement::capture()); + d->bind_current_thread(); + cudaqx_affinity_syscall_reset(); + (void)d->decode_guarded({0.1f, 0.1f}); + EXPECT_EQ(cudaqx_affinity_syscall_count(), 0) + << "bound thread must skip decode_guarded's guard"; + }); +} + +// unbind_thread() must forget the registration so guards re-engage — the +// session calls it at teardown before the bound thread's id can be recycled. +TEST(PinningBenchmark, UnbindThreadRestoresGuarding) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + d->bind_current_thread(); + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + EXPECT_EQ(cudaqx_affinity_syscall_count(), 0) + << "bound thread must skip the guard"; + d->unbind_thread(); + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + EXPECT_GT(cudaqx_affinity_syscall_count(), 0) + << "after unbind_thread() the guard must fire again"; + }); +} + +// ---- Fault-injection negative tests ------------------------------------ +// Each test sets AND unsets its own CUDAQX_SHIM_FAIL_* variable so no +// injection state leaks across tests. + +// T1.4: a container that blocks EVERY placement syscall (blanket injection). +// The decode must still be correct, and the degradation must be loud. +TEST(PinningNegative, FullyBlockedSyscallsDegradeLoudlyNotWrongly) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + setenv("CUDAQX_SHIM_FAIL_ALL_PLACEMENT", "1", 1); + auto d = makeEntryReadyPinnedLut(); + testing::internal::CaptureStderr(); + auto results = d->decode_batch(kChunk); + std::string err = testing::internal::GetCapturedStderr(); + unsetenv("CUDAQX_SHIM_FAIL_ALL_PLACEMENT"); + ASSERT_EQ(results.size(), kChunk.size()); + for (auto &r : results) + EXPECT_TRUE(r.converged) << "degraded env must not corrupt decode"; + EXPECT_NE(err.find("[cudaq-qec affinity] WARNING"), std::string::npos) + << "degradation must be loud"; + }); +} + +// T2.8: when the prior affinity cannot be read, the guard could never restore +// it, so it must not call sched_setaffinity at all (only-if-restorable rule). +TEST(PinningInvariants, BlockedGetaffinityAppliesNoAffinityAtAll) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + auto before = thread_placement::capture(); + setenv("CUDAQX_SHIM_FAIL_GETAFFINITY", "1", 1); + cudaqx_affinity_syscall_reset(); + auto results = d->decode_batch(kChunk); + unsetenv("CUDAQX_SHIM_FAIL_GETAFFINITY"); + EXPECT_EQ(readShimCounts().setaff, 0) + << "unreadable prior affinity must mean NO sched_setaffinity"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after); + ASSERT_EQ(results.size(), kChunk.size()); + EXPECT_TRUE(results[0].converged); + }); +} + +// T2.9: sched_setaffinity blocked (locked cpuset). The independent mempolicy +// half must still run, the failure must be loud, and nothing may leak. +TEST(PinningInvariants, BlockedSetaffinityStillAppliesMempolicyHalf) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + if (!get_mempolicy_usable()) + GTEST_SKIP() << "get_mempolicy blocked (container seccomp?); the guard " + "correctly skips set_mempolicy then, so setmem > 0 " + "cannot hold"; + run_on_worker([] { + auto d = makePinnedLut(); + auto before = thread_placement::capture(); + setenv("CUDAQX_SHIM_FAIL_SETAFFINITY", "1", 1); + cudaqx_affinity_syscall_reset(); + testing::internal::CaptureStderr(); + auto results = d->decode_batch(kChunk); + std::string err = testing::internal::GetCapturedStderr(); + unsetenv("CUDAQX_SHIM_FAIL_SETAFFINITY"); + EXPECT_GT(readShimCounts().setmem, 0) + << "mempolicy half must still run when sched_setaffinity is blocked"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) + << "blocked sched_setaffinity leaked a placement change: " + << before.describe_difference(after); + ASSERT_EQ(results.size(), kChunk.size()); + EXPECT_TRUE(results[0].converged); + EXPECT_NE(err.find("sched_setaffinity"), std::string::npos) + << "blocked sched_setaffinity must be loud; got: " << err; + }); +} + +// T2.10: set_mempolicy blocked (no CAP_SYS_NICE). The independent affinity +// half must still run, the failure must be loud, and nothing may leak. +TEST(PinningInvariants, BlockedSetMempolicyStillAppliesAffinityHalf) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + if (!get_mempolicy_usable()) + GTEST_SKIP() << "get_mempolicy blocked (container seccomp?); the guard " + "then skips the mempolicy half entirely and never prints " + "\"set_mempolicy failed\""; + run_on_worker([] { + auto d = makePinnedLut(); + auto before = thread_placement::capture(); + setenv("CUDAQX_SHIM_FAIL_SET_MEMPOLICY", "1", 1); + cudaqx_affinity_syscall_reset(); + testing::internal::CaptureStderr(); + auto results = d->decode_batch(kChunk); + std::string err = testing::internal::GetCapturedStderr(); + unsetenv("CUDAQX_SHIM_FAIL_SET_MEMPOLICY"); + EXPECT_GT(readShimCounts().setaff, 0) + << "affinity half must still run when set_mempolicy is blocked"; + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) + << "blocked set_mempolicy leaked a placement change: " + << before.describe_difference(after); + ASSERT_EQ(results.size(), kChunk.size()); + EXPECT_TRUE(results[0].converged); + EXPECT_NE(err.find("set_mempolicy failed"), std::string::npos) + << "blocked set_mempolicy must be loud; got: " << err; + }); +} + +// ---- Binding lifecycle negative tests +// ----------------------------------------------- + +// T3.13: re-binding on a second thread migrates the guard-skip: the new owner +// decodes guard-free, the previous owner is guarded again. promise/future +// handshakes keep both threads alive through every measured decode (no thread +// id can be recycled) and serialize the decodes (the shim counters are global). +TEST(PinningBenchmark, RebindMigratesGuardSkipToNewOwner) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + std::promise a_bound_p, b_done_p; + auto a_bound_f = a_bound_p.get_future(); + auto b_done_f = b_done_p.get_future(); + long a_first = -1, b_count = -1, a_second = -1; + std::thread ta([&] { + d->bind_current_thread(); + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + a_first = cudaqx_affinity_syscall_count(); + a_bound_p.set_value(); + b_done_f.wait(); // B has re-bound and decoded + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + a_second = cudaqx_affinity_syscall_count(); + }); + std::thread tb([&] { + a_bound_f.wait(); // A is the current owner + d->bind_current_thread(); // rebind: ownership migrates to B + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + b_count = cudaqx_affinity_syscall_count(); + b_done_p.set_value(); + }); + ta.join(); + tb.join(); + EXPECT_EQ(a_first, 0) << "owner A must decode guard-free before the rebind"; + EXPECT_EQ(b_count, 0) << "new owner B must decode guard-free after rebinding"; + EXPECT_GT(a_second, 0) + << "previous owner A must be guarded again after B re-bound"; +} + +// T3.13b: binding twice on the same thread is idempotent -- no throw, and the +// thread still decodes guard-free. +TEST(PinningBenchmark, DoubleBindIsIdempotent) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + d->bind_current_thread(); + EXPECT_NO_THROW(d->bind_current_thread()); + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + EXPECT_EQ(cudaqx_affinity_syscall_count(), 0); + }); +} + +// T3.14: unbind_thread() called from a DIFFERENT thread clears the owner; the +// still-alive previously-bound thread must be guarded again. The worker stays +// alive across the unbind via a promise/future handshake (no sleeps). +TEST(PinningBenchmark, UnbindFromAnotherThreadRestoresGuarding) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + std::promise bound_p, unbound_p; + auto bound_f = bound_p.get_future(); + auto unbound_f = unbound_p.get_future(); + long count = -1; + std::thread w([&] { + d->bind_current_thread(); + bound_p.set_value(); + unbound_f.wait(); // main thread has called unbind_thread() + cudaqx_affinity_syscall_reset(); + d->decode_batch(kChunk); + count = cudaqx_affinity_syscall_count(); + }); + bound_f.wait(); + d->unbind_thread(); // from the main thread, not the owner + unbound_p.set_value(); + w.join(); + EXPECT_GT(count, 0) + << "after unbind_thread() from another thread, the previously bound " + "thread must be guarded again"; +} + +// T3.16: two never-bound threads decode concurrently; each thread's placement +// must be fully restored and each decode correct. No shim-counter asserts: +// the counters are global and would race here. +TEST(PinningInvariants, ConcurrentUnboundDecodesRestoreBothThreads) { + auto d = makePinnedLut(); + std::atomic go{false}; + auto body = [&] { + while (!go.load(std::memory_order_acquire)) + std::this_thread::yield(); + auto before = thread_placement::capture(); + auto results = d->decode_batch(kChunk); + auto after = thread_placement::capture(); + EXPECT_EQ(before, after) << "concurrent unbound decode leaked placement: " + << before.describe_difference(after); + ASSERT_EQ(results.size(), kChunk.size()); + EXPECT_TRUE(results[0].converged); + }; + std::thread t1(body), t2(body); + go.store(true, std::memory_order_release); + t1.join(); + t2.join(); +} + +// T3.17: N threads race bind_current_thread() on one decoder. The race itself +// must not throw, and afterwards a FRESH never-bound thread must still be +// guarded and decode correctly. The racers are held alive (promise/shared +// handshake) until the fresh thread finishes so its thread id cannot be a +// recycled racer id that would spuriously skip the guard. No attempt to +// identify the winner. +TEST(PinningBenchmark, BindRaceLeavesFreshThreadsGuarded) { + if (!cudaqx_affinity_syscall_count) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + auto d = makePinnedLut(); + constexpr int kRacers = 4; + std::atomic go{false}; + std::atomic raced{0}; + std::promise done_p; + std::shared_future done_f(done_p.get_future()); + std::vector racers; + for (int i = 0; i < kRacers; ++i) + racers.emplace_back([&] { + while (!go.load(std::memory_order_acquire)) + std::this_thread::yield(); + EXPECT_NO_THROW(d->bind_current_thread()); + raced.fetch_add(1, std::memory_order_acq_rel); + done_f.wait(); // stay alive: our thread id must not be recycled yet + }); + go.store(true, std::memory_order_release); + while (raced.load(std::memory_order_acquire) < kRacers) + std::this_thread::yield(); + std::thread fresh([&] { + cudaqx_affinity_syscall_reset(); + auto results = d->decode_batch(kChunk); + EXPECT_GT(cudaqx_affinity_syscall_count(), 0) + << "a fresh thread that never bound must be guarded after the race"; + ASSERT_EQ(results.size(), kChunk.size()); + EXPECT_TRUE(results[0].converged); + }); + fresh.join(); + done_p.set_value(); + for (auto &t : racers) + t.join(); +} + +// ---- syscall / sysfs budget gates ------------------------------------------- + +// Perf-regression gate: the per-call guard must stay within a fixed syscall +// budget. If you legitimately add a syscall, change the constant in the same +// PR and say why in the commit message. +TEST(PinningBudget, UnboundDecodeStaysWithinSyscallBudget) { + if (!shimCountersPresent()) + GTEST_SKIP() << "affinity-counter shim not preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + (void)d->decode_batch(kChunk); // warm one-time paths + constexpr int kCalls = 10; + cudaqx_affinity_syscall_reset(); + for (int i = 0; i < kCalls; ++i) + (void)d->decode_batch(kChunk); + auto c = readShimCounts(); + long total = c.setaff + c.getaff + c.setmem + c.getmem; + std::printf( + "[pinning budget] %ld placement syscalls / %d unbound decodes " + "= %.1f per call (setaff=%ld getaff=%ld setmem=%ld getmem=%ld)\n", + total, kCalls, static_cast(total) / kCalls, c.setaff, c.getaff, + c.setmem, c.getmem); + // Measured on a mempolicy-capable host: 7 syscalls per unbound decode + // (2 sched_getaffinity + 2 sched_setaffinity + 1 get_mempolicy + + // 2 set_mempolicy). Budget = measured + one call of headroom. + constexpr long kMeasuredPerCall = 7; + EXPECT_LE(total, kMeasuredPerCall * kCalls + kMeasuredPerCall) + << "guard syscall budget exceeded: " << total << " for " << kCalls + << " calls"; + EXPECT_GT(total, 0); + }); +} + +// Sysfs budget gate: each guarded (unbound) decode currently re-reads the +// node's cpulist -- 1 open of /sys/devices/system/node/... per call. Cpuset +// caching should reduce this to amortized 0 -- tighten when it lands. +TEST(PinningBudget, UnboundDecodeStaysWithinSysfsOpenBudget) { + if (!shimCountersPresent() || !cudaqx_node_sysfs_open_count) + GTEST_SKIP() << "affinity-counter shim (with sysfs open counter) not " + "preloaded"; + run_on_worker([] { + auto d = makePinnedLut(); + (void)d->decode_batch(kChunk); // warm one-time paths + constexpr int kCalls = 10; + cudaqx_affinity_syscall_reset(); + for (int i = 0; i < kCalls; ++i) + (void)d->decode_batch(kChunk); + long opens = cudaqx_node_sysfs_open_count(); + std::printf("[pinning budget] %ld node-sysfs opens / %d unbound decodes " + "= %.1f per call\n", + opens, kCalls, static_cast(opens) / kCalls); + EXPECT_LE(opens, kCalls) << "node-sysfs open budget exceeded: " << opens + << " for " << kCalls << " calls"; + EXPECT_GT(opens, 0) << "expected the unbound guard to read the node " + "cpulist (counter wired?)"; + }); +}