Theseus: Runtime-adaptive, Hot-swappable GPU Collective Communication

This repository implements the main Theseus runtime: a lightweight schedule selector, a resource-managing schedule executor, NCCL-compatible interfaces, and the context daemon used to publish persistent selection-context elements. The daemon lives in the context_daemon/ subdirectory and is an independent Go project; it is not built by the main CMake configuration.

Two public repositories support this runtime:

• Theseus Schedule Generator is a standalone, CPU-only C++ DSL for authoring and simulating Theseus schedule JSON files. It interacts with this repository only through the versioned schedule format; generated schedules can be published with theseus-context new sched.
• Customized MSCCL++ is the patched MSCCL++ v0.7.0 dependency used by the executor for channels, connections, registered memory, semaphores, and proxy services. CMake fetches an exact pinned commit of this dependency automatically. Users do not have to build or install it separately.

The repositories are versioned independently. For a reproducible build, use the generator release documented for the schedule contract and keep the MSCCL++ commit pinned by this repository's CMakeLists.txt.

Research Prototype Scope

Theseus is a research artifact and proof of concept developed to validate runtime schedule selection, consistency, resource sharing, and hot swapping for GPU collective communication. It is intended for research, experimentation, and reproducibility rather than production deployment. The implementation was built within a research-paper development cycle and prioritizes the mechanisms evaluated in the paper over broad API coverage, long-term compatibility, operational hardening, and support for every cluster configuration. Please review the support matrix and known issues before using it in a new environment.

Theseus does not aim to supersede NCCL. Its NCCL-compatible layer implements a documented subset of the interface and relies on an actual NCCL installation for several fallback paths. Production requirements such as complete API/ABI compatibility, automated fault recovery, hardened control-plane security, extensive platform qualification, and stable upgrade guarantees remain outside the current scope.

The current GPU implementation also does not exploit recent Hopper-era hardware features such as TMA or end-to-end NVLS schedule execution. Its kernels primarily use conventional CUDA, CUDA IPC, and IB/RoCE communication building blocks, and the evaluated schedule tuning is oriented toward the hardware used in the research prototype. Support for newer GPU architectures in the build does not imply architecture-specific optimization or performance comparable to a production collective library.

Features of the schedule selection module:

1. Support user-specified endo and exo attributes that determines the best schedule for a collective request
2. Support flexible per-schedule usability and preference conditions
3. Provide a lightweight AllReduce-based consistency protocol that ensures all ranks agree on the selected schedule for every request
4. Support addition of new schedules and attributes at runtime
5. Support update of attribute values at runtime

Features of the schedule executor:

1. Faithfully execute schedules written in the Theseus schedule JSON format (see test/schedule_files/ for examples)
2. Perform delta migration between schedules
3. Support global sharing of a single scratch buffer and resizing of scratch buffer at runtime

File Organization

├── CMakeLists.txt
├── context_daemon
│   ├── DESIGN.md                     # Describe daemon architecture and guarantees
│   ├── README.md                     # Build, configuration, and operation guide
│   ├── cmd                           # Daemon and CLI entry points
│   ├── internal                      # Config, protocol, server, and persistent store
│   └── tests/e2e                     # Docker multi-node integration test
├── include
│   └── theseus
│       ├── executor.hpp               # Define schedule executor
│       ├── rank_schedule.hpp          # Represent a parsed JSON schedule of a rank
│       ├── schedule_common.hpp        # Define basic constants and structs in Theseus
│       ├── schedule_delta.hpp         # Represent the difference between two schedules
│       ├── selector_context.hpp       # Maintaining selector contexts (attrs + scheds)
│       └── selector.hpp               # Define schedule selector
├── README.md
├── src
│   ├── include                        # Inner Theseus header files
│   │   └── theseus
│   │       ├── alltoallv_execution_context.hpp # Defines alltoAllv execution context
│   │       ├── alltoallv_execution_kernel.hpp  # Defines the AlltoAllv kernel
│   │       ├── context_delta.hpp      # Represent the difference in selector context
│   │       ├── context_evaluator.hpp  # Evaluate scheds' usability & preference expressions
│   │       ├── execution_context.hpp  # Define the four resource management levels
│   │       ├── execution_kernel.hpp   # Define the CUDA kernel that executes a schedule
│   │       ├── exprtk.hpp             # Evaluate a plaintext expression (third-party)
│   │       └── lower_operation.hpp    # Lower operations in schedules to a GPU-aware format
│   │
│   ├── alltoallv_execution_context.cc # Implement AlltoAllv execution context
│   ├── alltoallv_execution_kernel.cu  # Define host-side stub for the AlltoAllv kernel
│   ├── alltoallv_executor.cu          # Implement AlltoAllv executor
│   ├── buffer_execution_context.cc    # Implement BEC (Level 2) in execution_context.hpp
│   ├── context_delta.cc               # Implement context_delta.hpp
│   ├── context_evaluator.cc           # Implement context_evaluator.hpp
│   ├── device_execution_context.cc    # Implement DEC (Level 3) in execution_context.hpp
│   ├── execution_kernel.cu            # Define host-side stub that invokes the CUDA kernel
│   ├── executor.cu                    # Implement executor.hpp
│   ├── global_execution_context.cc    # Implement GEC (Level 0) in execution_context.hpp
│   ├── lower_operation.cc             # Implement lower_operation.hpp
│   ├── rank_schedule.cc               # Implement rank_schedule.hpp
│   ├── schedule_delta.cc              # Implement schedule_delta.hpp
│   ├── schedule_execution_context.cc  # Implement SEC (Level 1) in execution_context.hpp
│   ├── selector.cc                    # Implement selector.hpp
│   └── selector_context.cc            # Implement seelctor_context.hpp
├── test
│   ├── CMakeLists.txt
│   ├── include                        # Define reader and writer for manifest/log files
│   ├── mpi_executor                   # Test executor with MPI on 2 GPUs
│   ├── mpi_nccl                       # Test Theseus NCCL interfaces with MPI on 2 GPUs
│   ├── mpi_selector                   # Test selector with MPI on numerous CPU cores
│   ├── manifest_files                 # Contain manifest files used by test/mpi_nccl
│   ├── schedule_files                 # Contain schedule files used by test/mpi_executor
│   └── unit                           # Contain single-process unit tests
├── theseus_nccl
│   ├── CMakeLists.txt                 # Theseus NCCL interfaces
│   ├── audit_shim                     # Audit shim for seamless NCCL integration
│   ├── include
│   │   └── theseus
│   │       └── nccl.h                 # Public NCCL APIs
│   └── src
│       ├── include
│       │   └── theseus                # Inner helper classes for NCCL APIs
│       ├── nccl_comm.cc               # Implement NCCL communicator APIs
│       ├── nccl.cu                    # Implement public NCCL APIs
│       └── nccl_fallback.cc           # Implement NCCL fallback
└── VERSION

Currently, we isolate the processing logic of AlltoAllv collective request from other common collectives. This is because while other collectives have a static traffic matrix, AlltoAllv traffic matrix is usually known at runtime when the request arrives. Therefore, we define a separate AlltoAllv executor, GPU kernel, and execution context for AlltoAllv collective requests. We plan to unify it and other collectives in the future.

We expose a new NCCL API (non-existent in the actual NCCL) for users to issue AlltoAllv requests. Its signature is as follows. See the ncclAlltoAllv API test for an example of its usage.

ncclResult_t ncclAlltoAllv(const void* sendbuff, void* recvbuff, const size_t *cntMatrixCpu, const size_t *cntMatrixGpu,
                           ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream);

The special part is the cntMatrixCpu and cntMatrixGpu, the traffic matrices. They should be flattened as an array of size $N\times N$, where $N$ is the world size. The value at index $i\cdot N+j$ means the number of data elements (not bytes, this is why it is called cntMatrix) sent from rank $i$ to rank $j$. In other words, each line in the matrix is a sender, and each column is a receiver. We require the user to provide two identical copies of this matrix, one on the CPU and the other on the GPU. The CPU-side traffic matrix is for the AlltoAllv executor to select the best schedule within the AlltoAllv kernel's capacity. The GPU-side traffic matrix is for the GPU runtime to read.

AlltoAllv kernel implements the hierarchical all-pairs algorithm. It ensures that all inter-node traffic never crosses the rail.

Building and Runtime Requirements

We assume the host machine has CMake >= 3.25, CUDA >= 12.0, and MPI (any MPI implementation is fine, e.g., OpenMPI and MPICH). Supporting AMD GPUs with ROCm/HIP runtime is future work.

At runtime, make sure the following requirements are met in your cluster:

1. GPUs on the same machine have peer access to each other (either via PCIe/NVLink/NVLS).
2. GPUs on different machines are connected via the RDMA network (InfiniBand and RoCEv2 are both fine) and IB verbs are correctly installed on all participating machines.
3. All participating machines should support GPU-Directed RDMA (GDR), which is true if all machines have the module nvidia_peermem or the older nv_peer_mem loaded. Otherwise, running multi-node communication will throw exceptions because we assume GDR when developing the code (which greatly simplifies cross-node data transmission).

Job Layout Assumptions

Theseus validates the following assumptions and rejects the job with an explicit error (ncclInvalidUsage through the NCCL interfaces) when they are violated, instead of failing far from the cause:

1. Same number of ranks on every node, placed in contiguous blocks. With $n$ ranks per node, node $k$ must host exactly ranks $[k\cdot n, (k+1)\cdot n)$. Do not set MSCCLPP_HOSTID to the same value on different nodes, as that makes distinct hosts indistinguishable.
2. At most 8 ranks per node in multi-node jobs. Cross-node transport uses a fixed rank-to-RNIC affinity map specified by MSCCLPP_HCA_DEVICES (local rank $k$ uses RNIC $k$ in the list). If MSCCLPP_HCA_DEVICES is not specified, we assume GPU $k$ is in affinity with the $k$-th RNIC on the node. Supporting NCCL-style topology detection is our future work.
3. One process manages exactly one GPU. All communicators created by a process must be bound to one CUDA device, and that device must remain the current device for every collective call on those communicators.
4. Send/receive buffer alignment. For any schedule-executed collective, the beginning addresses of the send and receive buffers must be congruent modulo 16 (see Usage requirement below).

The full list of enforced checks and their error semantics is in docs/support-matrix.md.

Building the Project

The project is built using CMake. Configuration downloads the customized MSCCL++ source from its public repository at the commit pinned by CMakeLists.txt. The schedule generator is not part of the runtime build and is needed only when authoring new schedule JSON files.

The context daemon is built separately from context_daemon/ using Go. Building the C++/CUDA runtime does not build or install the daemon. See the context daemon README for its build and deployment instructions.

The following script builds and installs the project.

mkdir build && cd build
cmake ..
make -j
make install # Optional

The default installation directory is the build directory. To set another installation directory, use cmake .. -DINSTALL_PREFIX=[install dir], where [install dir] should be an absolute path.

By default, tests and NCCL interfaces are built and NPKit is disabled. If you want to disable tests, use cmake .. -DTHESEUS_BUILD_TESTS=OFF. If you want to disable NCCL interfaces, use cmake .. -DTHESEUS_BUILD_NCCL=OFF. If you want to enable NPKit, use the following command.

cmake .. -DTHESEUS_NPKIT_FLAGS="-DENABLE_NPKIT \
-DENABLE_NPKIT_EVENT_TIME_SYNC_CPU \
-DENABLE_NPKIT_EVENT_TIME_SYNC_GPU \
-DENABLE_NPKIT_EVENT_EXECUTOR_INIT_ENTRY \
-DENABLE_NPKIT_EVENT_EXECUTOR_INIT_EXIT \
-DENABLE_NPKIT_EVENT_EXECUTOR_OP_BASE_ENTRY \
-DENABLE_NPKIT_EVENT_EXECUTOR_OP_BASE_EXIT"

This will instruct NPKit to capture all kinds of events during kernel execution (at some profiling overhead).

After building, you can test basic functionalities with ctest --output-on-failure. Note that some tests require 2 local GPUs to run. Such tests will not be built if the host machine has only one GPU card.

If you build the project for the first time, cmake and make will likely take several minutes each. This is because Theseus relies on many third-party projects, including MSCCL++, GoogleTest, nlohmann JSON, exprtk, etc. During make, C++ compiler spends most of the time on parsing the huge template-only exprtk header file which serves expression evaluation. NVCC compiler spends most of the time on lowering hundreds of instantiated template kernels. Make sure you add the -j flag in the make command to minimize compilation time.

Once installed, you may easily adopt Theseus in other projects like PyTorch. The header files are in [install dir]/include and the libraries are in [install dir]/lib.

Theseus Core APIs

Theseus provides three novel and useful features.

1. (F1) Flexible schedule selection policy. It ensures flexible schedule selection policy based on the schedule context. Schedule context consists of (1) user-defined endo and exo attributes and (2) per-schedule usability and preference conditions.
2. (F2) Schedule consistency. It ensure schedule consistency for each collective op with little overhead. Schedule consistency means that all ranks choose the same schedules for any collective requests.
3. (F3) Delta migration. It enables efficient migration between any two valid schedules. The migration reuses existing resources as much as possible. It triggers minimal cross-rank communication incurred by the delta between two schedules.

These three features are built in two core Theseus APIs.

• std::optional<uint32_t> theseus::Selector::selectSchedule(const CollectiveRequest &request); This API implements (F1) and (F2).
  • Args: It expects a collective request as input. A collective request packs arguments in NCCL Collective APIs, which contains send/recv buff, datatype, reduction op, root rank, etc.
  • RetVal: It returns the best matched schedule (if any) in the current context. Each schedule is identified by an uint32_t ID, which is the entry position in the Schedule Manifest File (SMF). No schedule will be returned if any of the two conditions hold: (1) there is no valid schedule that matches this request; (2) the usability conditions of all matching schedules evaluate to false. Among the usable schedules, the one with the highest preference score wins; on equal preference, the latest added schedule wins. A schedule that has been replaced remains selectable: since a schedule can only replace an earlier one, the replacement wins ties against the replaced schedule, and the selector returns to the replaced schedule once it outscores the replacement again (e.g., when a fail-slow condition recovers).
  • Feature: Under the hood, it leverages theseus::SelectorContext to maintain all the attributes and active schedules (F1). It exploits theseus::ContextEvaluator to efficiently evaluate usability and preference conditions to find the best match (F2). The evaluator implements sophisticated cache policy and expression AST optimization to avoid redundant computation.
• void theseus::Exector::execute(const CollectiveRequest &request, uint32_t schedId, std::shared_ptr<RankSchedule> schedule, std::optional<uint32_t> replacedSchedId, cudaStream_t stream);
  • Args:
    • request: Collective request is the same as above
    • schedId: The selected schedule's ID returned by the first API
    • schedule: Parsed schedule file for the current rank
    • replacedSchedId: The schedule ID that this schedule replaces (if any). Replacement will trigger delta migration from the replaced schedule to the current schedule.
    • stream: CUDA stream to launch the execution kernel on
  • Feature:
    • Under the hood, Executor manages materialized resources for all requests and selected schedules on four levels, thereby allowing for maximum resource sharing and proper isolation.
    • The workflow of Executor::execute consists of two steps: (1) CPU: materializing new resources required by the current request and schedule; (2) GPU: launching a CUDA kernel to execute the schedule.
    • In step (1), we use a 15-state FSM (Finite State Machine) to perform resource materialization and delta migration in a unified manner (F3). It walks down the four levels. Once it hits the bottom level, we have all required resources at hand. Repetitive requests will result in cache hits on all four levels, so it takes minimal time to prepare for kernel launch.
    • In step (2), we provide a general and highly-optimized CUDA kernel to execute a given schedule. We also expose many performance knobs that users can tune to achieve the best performance, e.g., num_threads_per_block.

Endogenous attribute paths. Endo attributes are declared in the Attribute Manifest File with "origin":"endo" and a path naming a field of the current collective request. Their values refresh on every request, so usability/preference expressions can depend on them. The supported paths are:

Path	Value (as integer)
request.coll	Collective type (enum theseus::CollectiveType)
request.sendbuf	Send buffer address
request.recvbuf	Receive buffer address
request.count	Cross-rank element count (the count argument of the NCCL API)
request.datatype	Data type (enum theseus::DataType)
request.datatypesize	Size of the data type in bytes
request.redop	Reduction operation (enum theseus::RedOpType)
request.packettype	LL packet format (enum theseus::PacketType)
request.root	Root rank (rooted collectives; -1 otherwise)
request.nlocalranks	Number of ranks per node
request.nranks	World size
request.messagesize	Cross-rank message size in bytes

Declaring an endo attribute with any other path fails with an explicit error when the attribute is first evaluated.

All other APIs in Theseus are the workhorse to realize these three features. Hopefully, users do not have to know about these APIs. If you do need to, please refer to the Theseus API documentation generated by doxygen.

CUDA Graph Capture

Theseus calls captured into CUDA graphs must obey two rules:

1. They form a single dependency chain (i.e., captured on one stream plus no multiple streams being captured in one graph). This is the usual way collectives are captured — collectives on one communicator must execute in the same order on every rank, so benchmarks and frameworks like PyTorch serialize them on one stream within a capture. This restriction arises because Theseus tracks each captured graph's kernels with one CUDA event, so concurrent captured branches could overwrite each other's completion state and defeat the wait that protects hot swapping and scratch growth.
2. No resource setup, scratch reallocation, or delta migration is allowed during the graph capture. This is because they trigger cudaMalloc/cudaStreamSynchronize /cudaMemcpy/cudaMemset that are not permitted when capturing.

The following table summarizes whether a Theseus call is allowed, based on what it triggers in the executor. YES + WAIT means the call is allowed but first waits until all kernels previously enqueued by this communicator have finished (kernels of other communicators and the application are not waited for).

What the call triggers	In capture	Not in capture
Cache hit	YES	YES
Normal resource setup	NO	YES
Scratch reallocation	NO	YES + WAIT
Delta migration	NO	YES + WAIT

To keep captured calls at cache hits, warm up every captured request shape (same collective, buffers, and sizes) once before capturing, with THESEUS_INITIAL_SCRATCH_SIZE set large enough for the largest request. There are examples of using Theseus with the CUDA graph in the test. Theseus calls not captured into CUDA graphs do not need to obey these rules.

Theseus NCCL Interfaces

To allow integration with modern LLM frameworks like PyTorch and Megatron, Theseus further wraps the two core APIs into NCCL interfaces. The compatibility baseline is NCCL 2.25.1: NCCL_VERSION_CODE and ncclGetVersion() both report 2.25.1. The public header also declares selected newer NCCL APIs since 2.26 or Theseus-specific extensions (for example, ncclAlltoAllv). Status of individual APIs is listed below and must not be interpreted as complete compatibility with a newer NCCL release.

Theseus provides static and shared libraries whose implementations are based on the Theseus core APIs rather than the original NCCL internals. Applications using the supported API subset can adopt Theseus by changing dynamic-library resolution via LD_AUDIT and LD_LIBRARY_PATH, or by linking the static library as described below. We recommend LD_AUDIT because Theseus NCCL library is named as libtheseus_nccl.so rather than libnccl.so. Unsupported calls are either forwarded to the configured fallback NCCL or rejected according to the status table.

Fallback mechanism. Currently, Theseus only implements the core NCCL APIs. For unimplemented NCCL APIs, we require the user to provide a system path to the actual NCCL library via the env var THESEUS_NCCL_LIB_PATH. At runtime, we will get available fallback NCCL APIs via dlsym and fall back to NCCL in such scenarios.

Environment variables. Theseus NCCL recognizes a different set of env vars compared to NCCL. We list them as follows.

Env var	Value	Explanation	Default
THESEUS_DEBUG	VERSION/WARN/INFO/ABORT/TRACE	Set the debug level	-
THESEUS_DEBUG_SUBSYS	See below	Set the enabled debug subsystems	INIT
THESEUS_DEBUG_FILE	System path	Write debug logs to a file	stdout
THESEUS_NCCL_LIB_PATH	System path	Point to the actual NCCL as fallback	-
THESEUS_SELECTION_CONTEXT	System path	Directory holding the selection context	~/.theseus/contexts/default
THESEUS_NPKIT_DUMP_DIR	System path	Path to dump NpKit trace	-
THESEUS_CONTEXT_CHECK_PERIOD	Number	Check for context update every $K$ requests	1000
THESEUS_VECTOR_ATTRIBUTE_MAX_LENGTH	Number	Maximum length for vector-valued attributes	1024
THESEUS_INITIAL_SCRATCH_SIZE	Number / numeric string ending with KB/MB/GB	Initial scratch size in Executor	16MB
THESEUS_ALTERNATE_SCRATCH_BUFFER	1/0	Use alternating scratch buffer in Executor	0
THESEUS_LL_PROTOCOL	LL8/LL16	Packet format for low-latency protocol	LL16
MSCCLPP_HCA_DEVICES	Comma-separated RNIC names	RNIC used by each local rank (local rank $k$ uses the $k$-th entry)	IB verbs enumeration order
MSCCLPP_GID_INDEX	Number	GID index of the RNICs, used in RoCE mode	0
MSCCLPP_SOCKET_IFNAME	Interface name	Network interface used by the bootstrap TCP network	Auto-detected
MSCCLPP_SOCKET_FAMILY	AF_INET/AF_INET6	Restrict the bootstrap network to IPv4/IPv6	-

Additional explanation:

• Prefixes: Every variable is recognized under both the THESEUS_ and MSCCLPP_ prefixes (e.g., THESEUS_HCA_DEVICES is equivalent to MSCCLPP_HCA_DEVICES); if both are set, the MSCCLPP_-prefixed one takes precedence. By convention, we list Theseus-specific variables under THESEUS_ and variables inherited from the customized MSCCL++ (transport/bootstrap configuration) under MSCCLPP_.
• THESEUS_DEBUG_SUBSYS: A comma-separated list among INIT, COLL, P2P, SHM, NET, GRAPH, TUNING, ENV, ALLOC, CALL, EXECUTOR, NCCL, SELECTOR, and ALL. Prefix the list with ^ to enable all subsystems except the listed ones.
• THESEUS_SELECTION_CONTEXT: The selection context directory must contain the Attribute Manifest File attribute_manifest.jsonl and the Schedule Manifest File schedule_manifest.jsonl (each may appear after startup). Both use the Newline-Delimited JSON format (NDJSON), as the .jsonl extension emphasizes. Attribute log files for exo attributes should also be named .jsonl. Only schedule files should be named .json. Each communicator binds to the selection context named by this variable at creation time. When the variable is unset, the default context $HOME/.theseus/contexts/default is used (falling back to ./theseus_context if HOME is unavailable).
• All valid_after timestamps in manifest and attribute log files are interpreted in the local time zone of the reading process (as standard, non-DST time). All participating machines and all manifest writers must therefore share one time zone and have synchronized clocks; a diverging machine only delays when new entries become visible, but the delay can be surprising. Note that manifest files are not portable across clusters configured with different time zones.
• THESEUS_NPKIT_DUMP_DIR: Only effective when Theseus is built with NpKit enabled.
• THESEUS_INITIAL_SCRATCH_SIZE: The scratch buffer grows automatically when a request demands more. Growth (like schedule migration) first waits for the kernels Theseus previously enqueued, since they still use the old resources (unrelated GPU work keeps running), and is rejected while a CUDA graph is being captured. Size the initial scratch generously to avoid growth in the steady state, and capture only steady-state requests into CUDA graphs.
• THESEUS_ALTERNATE_SCRATCH_BUFFER: If turned on, Executor will ensure two back-to-back collective requests will never use overlapping scratch area. This prevents data hazard in certain circumstances at the cost of doubling GPU scratch buffer size. See doc for more details.
• MSCCLPP_HCA_DEVICES: See job layout assumption 2 above. The list must name at least as many devices as there are ranks per node; the mapping is uniform within a node.

Theseus NCCL API status.

• ✔: Supported
• Placeholder: Provide an oversimplistic implementation
• Not supported: reasons include: (1) It is NCCL-specific call and not transferrable to Theseus architecture; (2) Implementation requires too much effort; (3) It rewrites critical states in NCCL communicator, and it is unsafe to fallback to NCCL directly. All unsupported API calls are not necessary for typical NCCL usage
• Under development: Code is on its way
• Fallback: We directly call NCCL fallback

Theseus NCCL API	Status
Management Functions
ncclMemAlloc	✔
ncclMemFree	✔
ncclGetVersion	✔
ncclGetUniqueId	✔
ncclCommInitRankConfig	Placeholder
ncclCommInitRank	✔
ncclCommInitAll	Not supported
ncclCommFinalize	✔
ncclCommDestroy	✔
ncclCommAbort	Placeholder
ncclCommSplit	✔ (config currently ignored with a warning)
ncclCommShrink	Under development
ncclCommInitRankScalable	Not supported
ncclGetErrorString	✔
ncclGetLastError	✔
ncclCommGetAsyncError	Placeholder
ncclCommCount	✔
ncclCommCuDevice	✔
ncclCommUserRank	✔
ncclCommRegister	Fallback
ncclCommDeregister	Fallback
ncclCommWindowRegister	Not supported
ncclCommWindowDeregister	Not supported
User Defined Reduction Operators
ncclRedOpCreatePreMulSum	Not supported
ncclRedOpDestroy	Not supported
Collective Communication Operations
ncclReduce	✔
ncclBcast	✔
ncclBroadcast	✔
ncclAllReduce	✔
ncclReduceScatter	✔
ncclAllGather	✔
ncclAlltoAll	✔
ncclAlltoAllv (not provided in NCCL)	✔
ncclGather	✔
ncclScatter	✔
Point to Point Communication Functions
ncclSend	Fallback
ncclRecv	Fallback
Group Calls
ncclGroupStart	Fallback
ncclGroupEnd	Fallback
ncclGroupSimulateEnd	Not supported

Usage requirement. For any collective API executed with a Theseus schedule, we require that the beginning addresses of send and receive buffers modulo 16 are the same. This ensures that the misaligned parts (relative to 16-byte words) can be copied between devices correctly. Violations are rejected with ncclInvalidUsage before any resource is materialized. Note that mainstream GPU workloads naturally meet this requirement. For example, the beginning addresses of on-device tensors in PyTorch are multiple of 16 bytes. The requirement does not apply to ncclAlltoAllv (its specialized executor supports 1-byte alignment) or to requests served by the NCCL fallback.

Integration with existing applications. Theseus acts as a drop-in replacement of the actual NCCL. Based on how the GPU application links to NCCL, there are two ways of integration.

• Dynamic library: If the application links to libnccl.so or libnccl.so.2 dynamically (you can check linkage of your application via ldd [your application]), no code changes are needed. Once you installed this project, the dynamic library is at [install dir]/lib/libtheseus_nccl.so. At runtime, specifying LD_AUDIT=[install dir]/lib/libtheseus_audit_nccl.so and LD_LIBRARY_PATH=[install dir]/lib:$LD_LIBRARY_PATH will resolve the symbols from Theseus. For example, Theseus can be integrated into NCCL tests and PyTorch in this way without changing any code.
• Static library: If the application links NCCL statically, you have to substitute Theseus's static library ([install dir]/lib/libnccl.a) in the code and build the application from source. The concrete method depends on your application.

Note 1: The only purpose of the audit shim is to intercept the name of the NCCL library that the dynamic linker is looking for and change it to libtheseus_nccl.so. Remember to unset LD_AUDIT and restore LD_LIBRARY_PATH when you want the official NCCL.

Note 2: In order to utilize ncclAlltoAllv in PyTorch, the new collective API non-existent in the official NCCL, you have to modify PyTorch code and build it from source in order to invoke this function. Following the guide at this PyTorch extension to integrate Theseus with PyTorch if you want optimized AlltoAllv communication.

To ensure success execution of application, it is recommended that you set the fallback NCCL library via THESEUS_NCCL_LIB_PATH. When the NCCL function is not implmented by Theseus or no schedule is able to serve the current request (e.g., ncclSend/ncclRecv), Theseus will execute the actual NCCL function instead. This ensures that the application runs successfully.

Persistent Context Management

Theseus separates context publication from the per-collective selection and execution path. The custom schedule generator produces schedule JSON. Users may either manually replicate and persist context elements on all nodes that may host a Theseus communicator, or run the daemon on each node to do this task. With manual replication, the additional requirement is that each file in a selection context are append-only when any communicator is using that context. Daemon meets this requirement. Either way, managing context are completely asynchronous with respect to the Theseus runtime.

When using the daemon, theseus-context new sched will publish that schedule and its policy. Attributes are published with new attr, and exogenous values are updated with set. See the damon's README.md for more information.

Supported Features and Limitations

docs/support-matrix.md is the authoritative classification of every collective, NCCL API, data type, protocol, channel type, and schedule opcode: whether it executes natively, is covered by a test, falls back to NCCL, or is rejected. It also records the exact software/hardware combinations the test suite has passed on. Consult it before relying on a feature that this README does not explicitly demonstrate.

Known Issues

1. If user frees the send or receive buffer before the communicator is torn down, future requests may run into the error thrown by cudaIpcOpenMemHandle saying that "resource already mapped". This is due to caching in the four-level Theseus execution context. According to the CUDA documentation, calling cudaFree on an exported memory region before calling cudaIpcCloseMemHandle in the importing context will result in undefined behavior. Since exported IPC handles are cached, unless user notifies the CCL that this user buffer has been deallocated, we cannot know when to evict and close the cached handles. In NCCL, this is solved by User buffer registration, namely ncclCommRegister and ncclCommDeregister. We plan to integrate this solution into Theseus. Currently, make sure all buffers passed to Theseus are alive before the communicator ends.

2. We assume that each process manages only one GPU when developing the code (the most common case). The code needs great modification to support one-process-many-GPUs mode (theoretically feasible, much engineering). Creating communicators on two different CUDA devices in one process, or switching the current device between calls, is rejected with ncclInvalidUsage.

3. Stable buffer pairing for repeated requests. The Level-2 (buffer) cache of the executor is keyed by local buffer allocations only, while a cached entry also holds the exchanged remote buffer registrations. If, for the same schedule, one rank reuses a local buffer while a peer switches to a new one, the ranks disagree on cache hit/miss and the miss side waits for a remote-buffer exchange that never happens — a deadlock, not an error message. To stay safe, satisfy one of:

   • each local send/receive allocation always communicates with the same remote allocations for equivalent requests (the natural pattern in LLM training, where the same set of tensors keeps communicating); or
   • use schedules whose channels are all scratch-to-scratch (scratch buffers are fully managed by Theseus, so their pairing is always stable).

   Varying the input/output ranges (offsets/sizes) within the same allocations is always safe; it is the allocation identity that must stay paired. This restriction will be lifted by a symmetric materialization-key design in a future release.