# Distributed Communication Service Plugin - Technical Specification

**Parent Document**: [SPEC.md](SPEC.md)

## 1. Component Overview

### Purpose and Scope

The Distributed Communication Service Plugin provides a runtime-loadable abstraction layer for distributed communication, enabling cuDensityMat to support multiple communication backends (MPI, NCCL) without compile-time dependencies.

**Location**: `density_matrix/src/distributed_interfaces/`

### Key Responsibilities

1. **Backend Abstraction**: Uniform API across MPI and NCCL
2. **Runtime Loading**: Dynamic plugin loading via `dlopen`
3. **Version Management**: Ensure ABI compatibility
4. **Communication Services**: Point-to-point, collective, synchronization
5. **Async Operations**: Non-blocking communication with request management

### Design Philosophy

- **Plugin Architecture**: Backends as dynamically loaded shared libraries
- **Minimal Overhead**: Thin wrapper around native APIs
- **Backend Selection**: User chooses via environment variable
- **Zero Dependencies**: Main library has no MPI/NCCL compile-time dependencies

## 2. Plugin Architecture

### Interface Definition

**File**: `density_matrix/include/cudensitymat.h`

```c
// Interface version for ABI compatibility
#define CUDENSITYMAT_DISTRIBUTED_INTERFACE_VERSION 260110

// Function pointer table (all function pointers return int: 0 for success, non-zero for failure)
typedef struct {
  int version;

  // Query operations (result returned via output pointer)
  int (*getNumRanks)(const cudensitymatDistributedCommunicator_t*, int32_t*);
  int (*getNumRanksShared)(const cudensitymatDistributedCommunicator_t*, int32_t*);
  int (*getProcRank)(const cudensitymatDistributedCommunicator_t*, int32_t*);

  // Synchronization
  int (*barrier)(const cudensitymatDistributedCommunicator_t*, void* barrierBuffer);

  // Async request management
  int (*createRequest)(cudensitymatDistributedRequest_t*);
  int (*destroyRequest)(cudensitymatDistributedRequest_t);
  int (*waitRequest)(cudensitymatDistributedRequest_t);
  int (*testRequest)(cudensitymatDistributedRequest_t, int32_t*);

  // NCCL-specific grouping
  int (*groupStart)(void);
  int (*groupEnd)(void);

  // Point-to-point operations
  int (*send)(const cudensitymatDistributedCommunicator_t*,
              const void*, int32_t, cudaDataType_t, int32_t, int32_t);
  int (*sendAsync)(const cudensitymatDistributedCommunicator_t*,
                   const void*, int32_t, cudaDataType_t, int32_t, int32_t,
                   cudensitymatDistributedRequest_t);
  int (*receive)(const cudensitymatDistributedCommunicator_t*,
                 void*, int32_t, cudaDataType_t, int32_t, int32_t);
  int (*receiveAsync)(const cudensitymatDistributedCommunicator_t*,
                      void*, int32_t, cudaDataType_t, int32_t, int32_t,
                      cudensitymatDistributedRequest_t);

  // Collective operations
  int (*bcast)(const cudensitymatDistributedCommunicator_t*,
               void*, int32_t, cudaDataType_t, int32_t);
  int (*allreduce)(const cudensitymatDistributedCommunicator_t*,
                   const void*, void*, int32_t, cudaDataType_t);
  int (*allreduceInPlace)(const cudensitymatDistributedCommunicator_t*,
                          void*, int32_t, cudaDataType_t);
  int (*allreduceInPlaceMin)(const cudensitymatDistributedCommunicator_t*,
                             void*, int32_t, cudaDataType_t);
  int (*allreduceDoubleIntMinloc)(const cudensitymatDistributedCommunicator_t*,
                                  const void*, void*);
  int (*allgather)(const cudensitymatDistributedCommunicator_t*,
                   const void*, void*, int32_t, cudaDataType_t);
} cudensitymatDistributedInterface_t;
```

### Plugin Registration

Each plugin exports a global symbol:

```c
// In MPI plugin
#ifdef __cplusplus
extern "C" {
#endif

cudensitymatDistributedInterface_t cudensitymatCommInterface = {
  CUDENSITYMAT_DISTRIBUTED_INTERFACE_VERSION,
  cudensitymatMpiCommSize,
  cudensitymatMpiCommSizeShared,
  cudensitymatMpiCommRank,
  cudensitymatMpiBarrier,
  cudensitymatMpiCreateRequest,
  cudensitymatMpiDestroyRequest,
  cudensitymatMpiWaitRequest,
  cudensitymatMpiTestRequest,
  cudensitymatMpiGroupStart,
  cudensitymatMpiGroupEnd,
  cudensitymatMpiSend,
  cudensitymatMpiSendAsync,
  // ... remaining function pointers
};

#ifdef __cplusplus
}
#endif
```

### Architecture Diagram

```mermaid
graph TB
    subgraph MainLibrary[cuDensityMat Main Library]
        Context[Context Class]
        DataPlans[Data Movement Plans]
    end
    
    subgraph PluginInterface[Plugin Interface]
        InterfaceStruct[cudensitymatDistributedInterface_t<br/>Function Pointer Table]
    end
    
    subgraph Plugins[Plugin Implementations]
        MPIPlugin[MPI Plugin<br/>libcudensitymat_distributed_interface_mpi.so]
        NCCLPlugin[NCCL Plugin<br/>libcudensitymat_distributed_interface_nccl.so]
    end
    
    subgraph Backends[Native Libraries]
        MPILib[MPI Library]
        NCCLLib[NCCL Library]
    end
    
    Context -->|dlopen/dlsym| InterfaceStruct
    DataPlans -->|uses| InterfaceStruct
    
    InterfaceStruct -.->|implementation| MPIPlugin
    InterfaceStruct -.->|implementation| NCCLPlugin
    
    MPIPlugin --> MPILib
    NCCLPlugin --> NCCLLib
    
    EnvVar[CUDENSITYMAT_COMM_LIB<br/>Environment Variable]
    EnvVar -.->|points to| MPIPlugin
    EnvVar -.->|points to| NCCLPlugin
```

## 3. Communication Services API

### Point-to-Point Operations

#### send / sendAsync

**Signature**:
```c
int send(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t dest,
    int32_t tag);

int sendAsync(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t dest,
    int32_t tag,
    cudensitymatDistributedRequest_t request);
```

**Semantics**:
- `send`: Blocking send (waits for completion)
- `sendAsync`: Non-blocking send (returns immediately)

**Buffer**: GPU device pointer

#### receive / receiveAsync

**Signature**: Similar to send

**Semantics**:
- `receive`: Blocking receive
- `receiveAsync`: Non-blocking receive

### Collective Operations

#### bcast

**Signature**:
```c
int bcast(
    const cudensitymatDistributedCommunicator_t* comm,
    void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t root);
```

**Semantics**: Broadcast from root to all ranks

#### allreduce / allreduceInPlace

**Signature**:
```c
int allreduce(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* sendbuf,
    void* recvbuf,
    int32_t count,
    cudaDataType_t datatype);

int allreduceInPlace(
    const cudensitymatDistributedCommunicator_t* comm,
    void* buffer,
    int32_t count,
    cudaDataType_t datatype);
```

**Reduction operation**: SUM (hardcoded; separate functions exist for other operations like MIN)

**Semantics**:
- `allreduce`: Separate send and receive buffers
- `allreduceInPlace`: In-place reduction

#### allreduceInPlaceMin

**Purpose**: Specialized in-place min reduction

#### allreduceDoubleIntMinloc

**Purpose**: Min reduction with location (for eigenvalue algorithms)

**Signature**:
```c
int allreduceDoubleIntMinloc(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* sendbuf,
    void* recvbuf);
```

#### allgather

**Signature**:
```c
int allgather(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* sendbuf,
    void* recvbuf,
    int32_t sendcount,
    cudaDataType_t datatype);
```

**Semantics**: Gather data from all ranks to all ranks

### Synchronization

#### barrier

**Signature**:
```c
int barrier(const cudensitymatDistributedCommunicator_t* comm, void* barrierBuffer);
```

**Note**: The `barrierBuffer` parameter is a device pointer. For MPI it is ignored (can be NULL). For NCCL it is used for an allreduce-based barrier implementation.

**Semantics**: Synchronize all ranks (wait until all reach barrier)

### Query Operations

#### getNumRanks

**Signature**:
```c
int getNumRanks(const cudensitymatDistributedCommunicator_t* comm, int32_t* numRanks);
```

**Output**: Total number of ranks in communicator (via `numRanks` pointer)

#### getNumRanksShared

**Signature**:
```c
int getNumRanksShared(const cudensitymatDistributedCommunicator_t* comm, int32_t* numRanksShared);
```

**Output**: Number of ranks sharing same node (via `numRanksShared` pointer)

#### getProcRank

**Signature**:
```c
int getProcRank(const cudensitymatDistributedCommunicator_t* comm, int32_t* procRank);
```

**Output**: Current rank ID, 0-indexed (via `procRank` pointer)

### Async Management

#### createRequest / destroyRequest

**Purpose**: Create/destroy request handles for async operations

```c
int createRequest(cudensitymatDistributedRequest_t* request);
int destroyRequest(cudensitymatDistributedRequest_t request);
```

#### waitRequest

**Purpose**: Wait for async operation to complete

```c
int waitRequest(cudensitymatDistributedRequest_t request);
```

#### testRequest

**Purpose**: Check if async operation is complete (non-blocking)

```c
int testRequest(cudensitymatDistributedRequest_t request, int32_t* flag);
```

**Output**: `flag = 1` if complete, `0` otherwise

### NCCL-Specific Operations

#### groupStart / groupEnd

**Purpose**: Group multiple NCCL operations for concurrent execution

```c
int groupStart(void);
int groupEnd(void);
```

**Usage**:
```c
iface->groupStart();
iface->sendAsync(...);  // Multiple sends/receives
iface->receiveAsync(...);
iface->groupEnd();  // Execute all together
```

**Note**: Required for NCCL send/recv pairing

## 4. Core Data Structures

### cudensitymatDistributedCommunicator_t

**Definition**:
```c
typedef struct {
  void* commPtr;    // Pointer to backend communicator (MPI_Comm* or ncclComm_t*)
  size_t commSize;  // Size of communicator structure
} cudensitymatDistributedCommunicator_t;
```

**Purpose**: Type-erased communicator

**Backend-Specific**:
- **MPI**: `commPtr` points to `MPI_Comm`
- **NCCL**: `commPtr` points to `ncclComm_t`

### cudensitymatDistributedRequest_t

**Definition**:
```c
typedef void* cudensitymatDistributedRequest_t;
```

**Purpose**: Type-erased async request handle

**Backend-Specific**:
- **MPI**: Points to `MPI_Request`
- **NCCL**: Points to `NcclAsyncRequest` structure (see below)

### Backend-Specific Structures

#### NcclAsyncRequest (NCCL Backend)

```c
typedef struct {
  cudaStream_t stream;   // Stream for operation
  cudaEvent_t event;     // Event for completion tracking
} NcclAsyncRequest;
```

**Purpose**: Track async NCCL operations via CUDA events

## 5. Plugin Loading and Integration

### Context Integration

**File**: `src/core/context.{hpp,cpp}`

```cpp
class Context {
public:
  // Plugin loading
  static void initDistributedInterface();
  
  // Access interface
  static const cudensitymatDistributedInterface_t* getDistributedInterface();
  
  // Query methods
  static bool isDistributed();
  static int getNumRanks();
  static int getRank();
  
  // Communicator management
  static void resetCommunicator(void* comm_ptr, size_t comm_size);
  static void syncCommunicator();
  
private:
  static cudensitymatDistributedInterface_t* distr_iface_;
  static void* comm_lib_handle_;  // dlopen handle
  static cudensitymatDistributedProvider_t distributed_provider_;
};
```

### Loading Mechanism

```cpp
void Context::initDistributedInterface() {
  // Get plugin path from environment
  const char* comm_lib_path = std::getenv("CUDENSITYMAT_COMM_LIB");
  if (!comm_lib_path) {
    distributed_provider_ = CUDENSITYMAT_DISTRIBUTED_PROVIDER_NONE;
    return;
  }
  
  // Load shared library
  comm_lib_handle_ = dlopen(comm_lib_path, RTLD_NOW | RTLD_LOCAL);
  if (!comm_lib_handle_) {
    throw std::runtime_error("Failed to load distributed interface plugin");
  }
  
  // Resolve interface symbol
  distr_iface_ = (cudensitymatDistributedInterface_t*)
      dlsym(comm_lib_handle_, "cudensitymatCommInterface");
  if (!distr_iface_) {
    throw std::runtime_error("Failed to find cudensitymatCommInterface symbol");
  }
  
  // Check version
  if (distr_iface_->version != CUDENSITYMAT_DISTRIBUTED_INTERFACE_VERSION) {
    throw std::runtime_error("Distributed interface version mismatch");
  }
  
  // Determine provider based on plugin path
  if (strstr(comm_lib_path, "mpi")) {
    distributed_provider_ = CUDENSITYMAT_DISTRIBUTED_PROVIDER_MPI;
  } else if (strstr(comm_lib_path, "nccl")) {
    distributed_provider_ = CUDENSITYMAT_DISTRIBUTED_PROVIDER_NCCL;
  }
}
```

**Cleanup**:
```cpp
Context::~Context() {
  if (comm_lib_handle_) {
    dlclose(comm_lib_handle_);
    comm_lib_handle_ = nullptr;
    distr_iface_ = nullptr;
  }
}
```

### Usage Patterns Throughout Codebase

**Data Movement Plans**:
```cpp
// In RotationDataPlanImpl
cudensitymatStatus_t RotationDataPlanImpl::submitWork(...) {
  auto* iface = Context::getDistributedInterface();
  
  int rank = iface->getProcRank(&comm_);
  int num_ranks = iface->getNumRanks(&comm_);
  
  int send_to = (rank + num_ranks - 1) % num_ranks;
  int recv_from = (rank + 1) % num_ranks;
  
  iface->sendAsync(&comm_, send_buffer, size, datatype, send_to, tag, &request);
  iface->receiveAsync(&comm_, recv_buffer, size, datatype, recv_from, tag, &request);
  
  iface->waitRequest(request);
}
```

**Context Synchronization**:
```cpp
void Context::syncCommunicator() {
  if (distr_iface_) {
    distr_iface_->barrier(&communicator_);
  }
}
```

## 6. MPI Backend Implementation

**File**: `src/distributed_interfaces/cudensitymat_distributed_interface_mpi.c`

**Build Script**: `src/distributed_interfaces/activate_mpi_cudm.sh`

### Key Features

- Wraps MPI calls (MPI_Send, MPI_Recv, MPI_Bcast, MPI_Allreduce, etc.)
- CUDA-aware MPI required for GPU buffers
- Stream ordering: Must sync before MPI operations
- Group operations: No-ops (MPI handles pairing automatically)

### MPI Wrapper Functions

#### cudensitymatMpiSend

```c
int cudensitymatMpiSend(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t dest,
    int32_t tag)
{
  // Unpack communicator
  MPI_Comm mpi_comm = unpackMpiCommunicator(comm);
  
  // Convert datatype
  MPI_Datatype mpi_type = convertCudaToMpiDataType(datatype);
  
  // Send
  return MPI_Send(buffer, count, mpi_type, dest, tag, mpi_comm);
}
```

#### cudensitymatMpiSendAsync

```c
int cudensitymatMpiSendAsync(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t dest,
    int32_t tag,
    cudensitymatDistributedRequest_t request)
{
  MPI_Comm mpi_comm = unpackMpiCommunicator(comm);
  MPI_Datatype mpi_type = convertCudaToMpiDataType(datatype);
  
  MPI_Request* mpi_request = (MPI_Request*)request;
  
  // Non-blocking send
  return MPI_Isend(buffer, count, mpi_type, dest, tag, mpi_comm, mpi_request);
}
```

#### cudensitymatMpiAllreduce

```c
int cudensitymatMpiAllreduce(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* sendbuf,
    void* recvbuf,
    int32_t count,
    cudaDataType_t datatype)
{
  MPI_Comm mpi_comm = unpackMpiCommunicator(comm);
  MPI_Datatype mpi_type = convertCudaToMpiDataType(datatype);
  
  return MPI_Allreduce(sendbuf, recvbuf, count, mpi_type, MPI_SUM, mpi_comm);
}
```

### Type Conversion

#### convertCudaToMpiDataType

```c
static MPI_Datatype convertCudaToMpiDataType(const cudaDataType_t cudaDataType) {
  switch (cudaDataType) {
    case CUDA_R_8I:  return MPI_INT8_T;
    case CUDA_R_16I: return MPI_INT16_T;
    case CUDA_R_32I: return MPI_INT32_T;
    case CUDA_R_64I: return MPI_INT64_T;
    case CUDA_R_32F: return MPI_FLOAT;
    case CUDA_R_64F: return MPI_DOUBLE;
    case CUDA_C_32F: return MPI_C_FLOAT_COMPLEX;
    case CUDA_C_64F: return MPI_C_DOUBLE_COMPLEX;
    default:         exit(EXIT_FAILURE);  // Unknown type
  }
}
```

### Communicator Unpacking

```c
static MPI_Comm unpackMpiCommunicator(const cudensitymatDistributedCommunicator_t * comm) {
  if (comm->commPtr == NULL) return MPI_COMM_NULL;
  if (sizeof(MPI_Comm) != comm->commSize) {
    exit(EXIT_FAILURE);  // Size mismatch
  }
  return *((MPI_Comm *)(comm->commPtr));
}
```

### Request Management

```c
int cudensitymatMpiWaitRequest(cudensitymatDistributedRequest_t request) {
  MPI_Request* mpi_request = (MPI_Request*)request;
  MPI_Status status;
  return MPI_Wait(mpi_request, &status);
}

int cudensitymatMpiDestroyRequest(cudensitymatDistributedRequest_t request) {
  free((MPI_Request*)request);
  return 0;
}
```

### Stream Ordering

**Issue**: MPI operations are not stream-ordered

**Solution**: Sync stream before MPI calls

```cpp
// In Context or data plans
void ensureStreamOrderingForDistributedBackend(cudaStream_t stream) {
  if (Context::getDistributedProvider() == CUDENSITYMAT_DISTRIBUTED_PROVIDER_MPI) {
    cudaStreamSynchronize(stream);
  }
  // NCCL is stream-ordered, no sync needed
}
```

### Group Operations

**MPI**: Group operations are no-ops (MPI automatically matches send/recv)

```c
int cudensitymatMpiGroupStart(void) {
  return 0;  // No-op
}

int cudensitymatMpiGroupEnd(void) {
  return 0;  // No-op
}
```

## 7. NCCL Backend Implementation (Experimental)

**File**: `src/distributed_interfaces/cudensitymat_distributed_interface_nccl.c`

**Build Script**: `src/distributed_interfaces/activate_nccl_cudm.sh`

### Key Features

- GPU-centric, stream-based operations
- Operations are stream-ordered (no sync needed)
- Requires explicit group start/end for send/recv pairing
- Barrier implemented via allreduce + stream sync
- Complex types: handled as pairs of real values
- Experimental status

### NCCL Wrapper Functions

#### cudensitymatNcclSendAsync

```c
int cudensitymatNcclSendAsync(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* buffer,
    int32_t count,
    cudaDataType_t datatype,
    int32_t dest,
    int32_t tag,
    cudensitymatDistributedRequest_t request)
{
  ncclComm_t nccl_comm = getNcclComm(comm);
  
  // Convert datatype and adjust count for complex
  ncclDataType_t nccl_type;
  int32_t nccl_count;
  convertCudaToNcclDataType(datatype, &nccl_type, &nccl_count, count);
  
  // Get stream from request
  NcclAsyncRequest* async_req = (NcclAsyncRequest*)(request);
  cudaStream_t stream = async_req->stream;
  
  // NCCL send (must be inside group)
  ncclResult_t result = ncclSend(buffer, nccl_count, nccl_type, 
                                 dest, nccl_comm, stream);
  
  // Record event for completion tracking
  cudaEventRecord(async_req->event, stream);
  
  return (int)result;
}
```

#### cudensitymatNcclAllreduce

```c
int cudensitymatNcclAllreduce(
    const cudensitymatDistributedCommunicator_t* comm,
    const void* sendbuf,
    void* recvbuf,
    int32_t count,
    cudaDataType_t datatype)
{
  ncclComm_t nccl_comm = getNcclComm(comm);
  
  ncclDataType_t nccl_type;
  int32_t nccl_count;
  convertCudaToNcclDataType(datatype, &nccl_type, &nccl_count, count);
  
  // Create stream for this operation
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  
  ncclResult_t result = ncclAllReduce(sendbuf, recvbuf, nccl_count,
                                      nccl_type, ncclSum, nccl_comm, stream);
  
  // Synchronize
  cudaStreamSynchronize(stream);
  cudaStreamDestroy(stream);
  
  return (int)result;
}
```

### Type Conversion with Count Adjustment

```c
static void convertCudaToNcclDataType(const cudaDataType_t cudaDataType,
                                       ncclDataType_t* ncclDataType,
                                       int32_t* ncclCount,
                                       int32_t cudaCount)
{
  switch (cudaDataType) {
    case CUDA_R_32F:
      *ncclDataType = ncclFloat;
      *ncclCount = cudaCount;
      break;
      
    case CUDA_R_64F:
      *ncclDataType = ncclDouble;
      *ncclCount = cudaCount;
      break;
      
    case CUDA_C_32F:
      *ncclDataType = ncclFloat;
      *ncclCount = cudaCount * 2;  // Two floats per complex
      break;
      
    case CUDA_C_64F:
      *ncclDataType = ncclDouble;
      *ncclCount = cudaCount * 2;  // Two doubles per complex
      break;
      
    default:
      exit(EXIT_FAILURE);  // Unknown type
      break;
  }
}
```

**Key Point**: Complex types handled as pairs of real values (count doubled)

### Communicator Handling

```c
static ncclComm_t getNcclComm(const cudensitymatDistributedCommunicator_t* comm) {
  if (comm->commPtr == NULL) return NULL;
  if (sizeof(ncclComm_t) != comm->commSize) {
    exit(EXIT_FAILURE);  // Size mismatch
  }
  return *((ncclComm_t*)(comm->commPtr));
}
```

### NcclAsyncRequest Structure

```c
typedef struct {
  cudaStream_t stream;
  cudaEvent_t event;
} NcclAsyncRequest;

int cudensitymatNcclCreateRequest(
    cudensitymatDistributedRequest_t* request)
{
  NcclAsyncRequest* async_req = (NcclAsyncRequest*)malloc(sizeof(NcclAsyncRequest));
  
  cudaStreamCreate(&async_req->stream);
  cudaEventCreate(&async_req->event);
  
  *request = (cudensitymatDistributedRequest_t)async_req;
  return 0;
}

int cudensitymatNcclWaitRequest(
    cudensitymatDistributedRequest_t request)
{
  NcclAsyncRequest* async_req = (NcclAsyncRequest*)request;
  
  // Wait on event
  cudaEventSynchronize(async_req->event);
  
  return 0;
}

int cudensitymatNcclDestroyRequest(
    cudensitymatDistributedRequest_t request)
{
  NcclAsyncRequest* async_req = (NcclAsyncRequest*)request;
  
  cudaEventDestroy(async_req->event);
  cudaStreamDestroy(async_req->stream);
  free(async_req);
  
  return 0;
}
```

### Group Operations

**Required for send/recv pairing**:

```c
int cudensitymatNcclGroupStart(void) {
  return (int)ncclGroupStart();
}

int cudensitymatNcclGroupEnd(void) {
  return (int)ncclGroupEnd();
}
```

**Usage**:
```c
// Group send/recv operations
iface->groupStart();

for (int i = 0; i < num_exchanges; ++i) {
  iface->sendAsync(comm, send_buf[i], ...);
  iface->receiveAsync(comm, recv_buf[i], ...);
}

iface->groupEnd();  // Launch all operations
```

### Barrier Implementation

```c
int cudensitymatNcclBarrier(const cudensitymatDistributedCommunicator_t* comm,
                            void* barrierBuffer) {
  ncclComm_t nccl_comm = getNcclComm(comm);
  
  // Create stream
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  
  // Allreduce-based barrier using the provided device buffer
  ncclResult_t result = ncclAllReduce(barrierBuffer, barrierBuffer, 1, ncclInt, 
                                      ncclSum, nccl_comm, stream);
  
  // Synchronize
  cudaStreamSynchronize(stream);
  cudaStreamDestroy(stream);
  
  return (int)result;
}
```

**Note**: NCCL has no native barrier; implemented via allreduce + sync

### Stream Ordering

**NCCL operations are stream-ordered**: No explicit sync needed before NCCL calls

```cpp
// In data plans
void submitWork(cudaStream_t stream) {
  // No sync needed with NCCL
  // Operations automatically ordered with stream
  
  iface->sendAsync(comm, buffer, ..., stream);
  // Subsequent CUDA operations on same stream wait for send
}
```

## 8. Backend Selection and Configuration

### Build-Time Configuration

**Environment Variables for Build Scripts**:
- `CUDA_PATH`: Path to CUDA installation
- `MPI_PATH`: Path to MPI installation (for MPI plugin)
- `NCCL_PATH`: Path to NCCL installation (for NCCL plugin)

**Build Scripts**:
```bash
# MPI plugin
cd density_matrix/src/distributed_interfaces/
./activate_mpi_cudm.sh

# NCCL plugin
./activate_nccl_cudm.sh
```

**Output**: Shared library in same directory
- `libcudensitymat_distributed_interface_mpi.so`
- `libcudensitymat_distributed_interface_nccl.so`

### Runtime Selection

**Environment Variable**: `CUDENSITYMAT_COMM_LIB`

```bash
# Use MPI backend
export CUDENSITYMAT_COMM_LIB=/path/to/libcudensitymat_distributed_interface_mpi.so

# Use NCCL backend
export CUDENSITYMAT_COMM_LIB=/path/to/libcudensitymat_distributed_interface_nccl.so
```

### User API for Distributed Setup

**C API**:
```c
// Activate distributed execution
cudensitymatStatus_t cudensitymatResetDistributedConfiguration(
    cudensitymatHandle_t handle,
    void* comm_ptr,
    size_t comm_size);

// Query
int cudensitymatGetNumRanks();
int cudensitymatGetProcRank();
```

**Usage** (MPI example):
```c
// Initialize MPI
MPI_Init(&argc, &argv);
MPI_Comm world = MPI_COMM_WORLD;

// Duplicate communicator
MPI_Comm cudm_comm;
MPI_Comm_dup(world, &cudm_comm);

// Pass to cuDensityMat
cudensitymatResetDistributedConfiguration(handle, &cudm_comm, sizeof(MPI_Comm));

// ... use library ...

// Cleanup
MPI_Comm_free(&cudm_comm);
MPI_Finalize();
```

**Usage** (NCCL example):
```c
// Initialize NCCL
ncclUniqueId id;
if (rank == 0) ncclGetUniqueId(&id);
MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);

ncclComm_t nccl_comm;
ncclCommInitRank(&nccl_comm, nranks, id, rank);

// Pass to cuDensityMat
cudensitymatResetDistributedConfiguration(handle, &nccl_comm, sizeof(ncclComm_t));

// ... use library ...

// Cleanup
ncclCommDestroy(nccl_comm);
```

### Restrictions

- **Power-of-Two Ranks**: cuDensityMat requires 1, 2, 4, 8, 16, 32, 64, ... ranks
- **One GPU per Process**: Each rank must be assigned exactly one GPU
- **User Responsibility**: GPU assignment via `cudaSetDevice` before context creation

## 9. Integration with Core Components

### Data Movement Plans Usage

**RotationDataPlanImpl**:
```cpp
cudensitymatStatus_t RotationDataPlanImpl::submitWork(...) {
  auto* iface = Context::getDistributedInterface();
  
  // Ring communication
  iface->sendAsync(&comm_, send_buf, size, type, prev_rank, tag, &req);
  iface->receiveAsync(&comm_, recv_buf, size, type, next_rank, tag, &req);
  iface->waitRequest(req);
}
```

**RedistributionDataPlanImpl**:
```cpp
cudensitymatStatus_t RedistributionDataPlanImpl::submitWork(...) {
  auto* iface = Context::getDistributedInterface();
  
  // Pairwise exchanges
  if (provider == NCCL) {
    iface->groupStart();
  }
  
  for (auto& exchange : schedule_) {
    if (my_rank == exchange.src) {
      iface->sendAsync(&comm_, slice_buf, size, type, exchange.dst, tag, &req);
    }
    if (my_rank == exchange.dst) {
      iface->receiveAsync(&comm_, recv_buf, size, type, exchange.src, tag, &req);
    }
  }
  
  if (provider == NCCL) {
    iface->groupEnd();
  }
  
  iface->barrier(&comm_);
}
```

### Stream Ordering Utilities

```cpp
// In context.hpp
void ensureStreamOrderingForDistributedBackend(cudaStream_t stream) {
  auto provider = Context::getDistributedProvider();
  
  if (provider == CUDENSITYMAT_DISTRIBUTED_PROVIDER_MPI) {
    // MPI requires stream sync
    cudaStreamSynchronize(stream);
  }
  // NCCL is stream-ordered, no sync needed
}
```

### Context Methods

```cpp
// Synchronize all ranks
void Context::syncCommunicator() {
  if (distr_iface_) {
    distr_iface_->barrier(&communicator_);
  }
}

// Set communicator
void Context::resetCommunicator(void* comm_ptr, size_t comm_size) {
  communicator_.commPtr = comm_ptr;
  communicator_.commSize = comm_size;
}

// Query
bool Context::isDistributed() {
  return (distr_iface_ != nullptr);
}

int Context::getNumRanks() {
  return distr_iface_ ? distr_iface_->getNumRanks(&communicator_) : 1;
}

int Context::getRank() {
  return distr_iface_ ? distr_iface_->getProcRank(&communicator_) : 0;
}
```

## 10. Type Conversion and Data Marshalling

### CUDA to MPI Datatype Mapping

| CUDA Type | MPI Type |
|-----------|----------|
| `CUDA_R_8I`  | `MPI_INT8_T` |
| `CUDA_R_16I` | `MPI_INT16_T` |
| `CUDA_R_32I` | `MPI_INT32_T` |
| `CUDA_R_64I` | `MPI_INT64_T` |
| `CUDA_R_32F` | `MPI_FLOAT` |
| `CUDA_R_64F` | `MPI_DOUBLE` |
| `CUDA_C_32F` | `MPI_C_FLOAT_COMPLEX` |
| `CUDA_C_64F` | `MPI_C_DOUBLE_COMPLEX` |

### CUDA to NCCL Datatype Mapping

| CUDA Type | NCCL Type | Count Adjustment |
|-----------|-----------|------------------|
| `CUDA_R_32F` | `ncclFloat` | x1 |
| `CUDA_R_64F` | `ncclDouble` | x1 |
| `CUDA_C_32F` | `ncclFloat` | x2 (treat as 2 floats) |
| `CUDA_C_64F` | `ncclDouble` | x2 (treat as 2 doubles) |

**Key Difference**: NCCL handles complex as pairs of reals (count doubled)

### Alignment Considerations

- **GPU Buffers**: Must be 256-byte aligned (same as other cuDensityMat requirements)
- **MPI**: CUDA-aware MPI handles GPU pointers directly
- **NCCL**: Designed for GPU buffers, handles alignment internally

## 11. Error Handling

### Return Code Convention

All plugin function pointers return `int`:

```c
// Success
return 0;

// Failure (non-zero; actual value is backend-specific, e.g., MPI error code or ncclResult_t)
return error_code;
```

### Error Reporting Patterns

**MPI Backend**:
```c
// MPI functions return int (MPI error code)
// The return value is passed through directly
int mpi_result = MPI_Send(...);
return mpi_result;  // 0 (MPI_SUCCESS) on success
```

**NCCL Backend**:
```c
// NCCL functions return ncclResult_t, cast to int
ncclResult_t nccl_result = ncclSend(...);
return (int)nccl_result;  // 0 (ncclSuccess) on success
```

### Backend-Specific Error Cases

**MPI**:
- Invalid rank (out of range)
- Tag mismatch
- Communicator invalid
- Network errors

**NCCL**:
- CUDA errors (invalid device, stream)
- Network errors
- Group mismatch (send/recv not paired)
- Unsupported operation

## 12. File Organization

### Directory Structure

```
density_matrix/src/distributed_interfaces/
├── cudensitymat_distributed_interface_mpi.c    # MPI plugin implementation
├── activate_mpi_cudm.sh                        # MPI plugin build script
├── cudensitymat_distributed_interface_nccl.c   # NCCL plugin implementation
└── activate_nccl_cudm.sh                       # NCCL plugin build script
```

### Build Scripts

**activate_mpi_cudm.sh**:
```bash
#!/bin/bash
set -e

# Check environment variables
if [ -z "$CUDA_PATH" ]; then
  echo "Error: CUDA_PATH not set"
  exit 1
fi

if [ -z "$MPI_PATH" ]; then
  echo "Error: MPI_PATH not set"
  exit 1
fi

# Compile
gcc -shared -fPIC -O3 \
    -I$CUDA_PATH/include \
    -I$MPI_PATH/include \
    -I../../include \
    cudensitymat_distributed_interface_mpi.c \
    -L$MPI_PATH/lib -lmpi \
    -o libcudensitymat_distributed_interface_mpi.so

# Set environment variable
export CUDENSITYMAT_COMM_LIB=$PWD/libcudensitymat_distributed_interface_mpi.so

echo "MPI plugin built successfully"
echo "Set CUDENSITYMAT_COMM_LIB=$CUDENSITYMAT_COMM_LIB"
```

**activate_nccl_cudm.sh**: Similar structure for NCCL

### Integration Points in Main Codebase

**Context**:
- `src/core/context.{hpp,cpp}` - Plugin loading and interface access

**Data Plans**:
- `src/core/coplan_data_state_dense_rotation.hpp` - Uses plugin for rotation
- `src/core/coplan_data_state_dense_redistribution.hpp` - Uses plugin for redistribution

**Public API**:
- `include/cudensitymat.h` - Interface definition and data structures

## 13. Extension and Customization

### Adding New Backends

To add a new backend (e.g., NVSHMEM):

1. **Implement Plugin**:
```c
// cudensitymat_distributed_interface_nvshmem.c

// Implement all interface functions (return int: 0 for success)
int cudensitymatNvshmemSend(...) {
  // NVSHMEM implementation
}

// ... all other functions ...

// Export interface
cudensitymatDistributedInterface_t cudensitymatCommInterface = {
  CUDENSITYMAT_DISTRIBUTED_INTERFACE_VERSION,
  cudensitymatNvshmemCommSize,
  cudensitymatNvshmemCommSizeShared,
  cudensitymatNvshmemCommRank,
  // ... all function pointers in struct order
};
```

2. **Create Build Script**:
```bash
#!/bin/bash
# activate_nvshmem_cudm.sh

gcc -shared -fPIC -O3 \
    -I$CUDA_PATH/include \
    -I$NVSHMEM_PATH/include \
    -I../../include \
    cudensitymat_distributed_interface_nvshmem.c \
    -L$NVSHMEM_PATH/lib -lnvshmem \
    -o libcudensitymat_distributed_interface_nvshmem.so

export CUDENSITYMAT_COMM_LIB=$PWD/libcudensitymat_distributed_interface_nvshmem.so
```

3. **Update Provider Enum** (optional):
```c
// In cudensitymat.h
typedef enum {
  CUDENSITYMAT_DISTRIBUTED_PROVIDER_NONE    = 0,
  CUDENSITYMAT_DISTRIBUTED_PROVIDER_MPI     = 1,
  CUDENSITYMAT_DISTRIBUTED_PROVIDER_NCCL    = 2,
  CUDENSITYMAT_DISTRIBUTED_PROVIDER_NVSHMEM = 3,  // New
} cudensitymatDistributedProvider_t;
```

### Interface Versioning

**Version Field**: Ensures ABI compatibility

**Version Bumps**:
- **Patch**: Bug fixes, no interface changes
- **Minor**: Add new optional functions (backward compatible)
- **Major**: Breaking changes (function signature changes, removed functions)

**Check at Load Time**:
```cpp
if (distr_iface_->version != CUDENSITYMAT_DISTRIBUTED_INTERFACE_VERSION) {
  throw std::runtime_error("Version mismatch");
}
```

### Backward Compatibility

- Keep old interface versions available
- Provide version-specific loading
- Support multiple interface versions simultaneously (if needed)

---

## Related Specifications

- **[SPEC.md](SPEC.md)**: Main technical specification
- **[SPEC_CORE.md](SPEC_CORE.md)**: Core classes using distributed communication
- **[SPEC_EXECUTOR.md](SPEC_EXECUTOR.md)**: Executor component
- **[SPEC_BASIC_OPS.md](SPEC_BASIC_OPS.md)**: Basic operations

---

**Document Version**: 1.0  
**Last Updated**: 2026-02-11  
**Part of**: cuDensityMat Technical Specifications
