Communication Simulation: SimCCL + astra-sim + NS-3

Communication Stack Overview

SimAI's communication simulation covers the complete network stack — from high-level collective operations (AllReduce, AllGather) down to packet-level RDMA transport with DCQCN congestion control and PFC flow control. This page provides a comprehensive deep-dive organized into three layers: (1) Collective Decomposition (Sections 2-4): how NCCL algorithms decompose collectives into point-to-point flows using Ring, Tree, and NVLS patterns; (2) Transport & Network Mechanics (Sections 5-6): how flows traverse the network via RoCEv2 RDMA with DCQCN/HPCC/TIMELY congestion control, RED-based ECN marking, PFC back-pressure, and per-flow ECMP routing; (3) Topology & Configuration (Sections 7-10): the five supported datacenter topologies (Spectrum-X, AlibabaHPN, DCN+), NS-3 configuration parameters, and end-to-end usage examples.

Each layer has a well-defined interface to the one below it. Collective operations produce FlowModels (a set of SingleFlow structs). astra-sim consumes these flows and dispatches them to the network backend. The backend reports completion events back to astra-sim, which then triggers dependent flows.

NCCL Algorithm Simulation (MockNcclGroup)

The MockNcclGroup class is the core of SimCCL. It faithfully replicates NCCL's algorithm selection and flow decomposition logic, translating high-level collective operations into point-to-point data flows that can be simulated by the network backend.

2.1 Supported Algorithms

// NCCL algorithm definitions mirrored in SimCCL
#define NCCL_ALGO_TREE            0   // Tree reduction (hierarchical)
#define NCCL_ALGO_RING            1   // Ring (bandwidth-optimal)
#define NCCL_ALGO_COLLNET_DIRECT  2   // Direct CollNet (switch-assisted)
#define NCCL_ALGO_COLLNET_CHAIN   3   // Chain CollNet (pipelined)
#define NCCL_ALGO_NVLS            4   // NVLink Switching (Hopper+)
#define NCCL_ALGO_NVLS_TREE       5   // NVLS + Tree hybrid

#define NCCL_NUM_ALGORITHMS       6
#define NCCL_NUM_PROTOCOLS        3   // LL, LL128, Simple

2.2 Algorithm Selection Logic

Algorithm selection in SimCCL is GPU-type-aware. The system considers the GPU architecture (A100, H100, H800), the group type (TP, DP, PP, EP), the number of ranks, and whether NVLink Switching is available. This mirrors NCCL's real-world decision tree:

ncclInfo* MockNcclGroup::get_algo_proto_info(
    GroupType type, int rank, ComType op, uint64_t data_size) {

    ncclInfo* info = new ncclInfo();
    info->nBytes = data_size;
    info->op = op;

    if (op == All_Reduce && type == TP) {
        if (gpu_type == H100 || gpu_type == H800) {
            if (nRanks >= 8 && NVLSenable)
                info->algorithm = NCCL_ALGO_NVLS;  // Hopper favors NVLS
            else
                info->algorithm = NCCL_ALGO_RING;
        } else if (gpu_type == A100 || gpu_type == A800) {
            info->algorithm = NCCL_ALGO_RING;   // Ampere uses Ring
        }
    } else if (op == All_Reduce && type == DP) {
        info->algorithm = NCCL_ALGO_RING;       // DP always Ring
    } else if (op == All_to_All) {
        info->algorithm = NCCL_ALGO_RING;       // AllToAll uses Ring
    } else if (op == All_Gather) {
        if (type == TP && NVLSenable && nRanks >= 8)
            info->algorithm = NCCL_ALGO_NVLS_TREE;
        else
            info->algorithm = NCCL_ALGO_RING;
    }

    return info;
}

2.3 Ring AllReduce Decomposition

The Ring AllReduce is the most common collective algorithm. It operates in two phases: ReduceScatter (each GPU sends 1/N of its data around the ring, reducing as it goes) and AllGather (each GPU broadcasts its reduced chunk). For N GPUs, this requires 2*(N-1) steps total.

std::map<int, shared_ptr<FlowModels>>
MockNcclGroup::genAllReduceRingFlowModels(
    GroupType type, int rank, uint64_t data_size) {

    int nranks = gp_info.nRanks;
    int chunkcount = 2 * (nranks - 1);  // reduce + broadcast phases
    chunksize = data_size / nranks / ringchannels.size();

    // Phase 1: ReduceScatter
    for (int step = 0; step < nranks - 1; step++) {
        for (auto& ring : ringchannels) {
            int src_rank = ring[(ring_idx + nranks - step) % nranks];
            int dest_rank = ring[(ring_idx + nranks - step + 1) % nranks];

            tmp_result = SingleFlow(
                flow_id, src_rank, dest_rank, chunksize,
                {prev_flow_id},      // depends on previous step
                {},                   // no parallel deps
                {child_flow_id},      // next step depends on this
                ring_id, chunk_id,
                chunkcount, "RING");
        }
    }

    // Phase 2: AllGather
    for (int step = 0; step < nranks - 1; step++) {
        for (auto& ring : ringchannels) {
            // Similar flow generation with broadcast semantics
            tmp_result = SingleFlow(
                flow_id, src_rank, dest_rank, chunksize,
                {prev_flow_id}, {}, {child_flow_id},
                ring_id, chunk_id, chunkcount, "RING");
        }
    }

    return flow_models;
}

2.4 NVLS AllReduce (Hopper / Blackwell)

NVLS (NVLink Switching) leverages NVSwitch to provide all-to-all NVLink connectivity within a node. Instead of passing data around a ring, all GPUs can simultaneously access a shared memory region through the NVSwitch. The NVLS-Tree variant combines NVSwitch-based intra-node communication with a tree reduction for inter-node communication.

shared_ptr<FlowModels>
MockNcclGroup::genallReduceNVLSTreeFlowModels(
    GroupType type, int rank, uint64_t data_size) {

    // Step 1: NVLS Reduce (intra-node via NVSwitch)
    // All GPUs within a node reduce their data through
    // the NVSwitch shared memory multicast
    generate_flow_model_nvls_tree_allreduce_up(
        rank, data_size, flow_models);

    // Step 2: Tree Reduce (inter-node via NIC)
    // One representative GPU per node participates in
    // a tree reduction across nodes

    // Step 3: Tree Broadcast (inter-node via NIC)
    // Root broadcasts the fully reduced result down the tree

    // Step 4: NVLS Broadcast (intra-node via NVSwitch)
    // Representative GPU shares the result with all node-local GPUs
    generate_flow_model_nvls_tree_allreduce_down(
        rank, data_size, flow_models);

    return flow_models;
}

// NVLS uses multicast addressing through NVSwitch
// Each GPU writes to a shared buffer that is atomically
// reduced by the NVSwitch hardware itself
void MockNcclGroup::generate_flow_model_nvls_tree_allreduce_up(
    int rank, uint64_t data_size,
    shared_ptr<FlowModels> flow_models) {

    // NVLS multicast: each GPU sends to NVSwitch
    for (auto& nvswitch : gp_info.NVSwitchs) {
        SingleFlow flow(flow_id, rank, nvswitch,
            data_size / nRanks,
            {}, {}, {child_ids},
            0, chunk_id, chunk_count, "NVLS");
        flow_models->add(flow);
    }
}

2.5 Hardware Latency Parameters

SimCCL models hardware-specific latencies for each algorithm and protocol combination. These base latency values are added to the data transfer time to account for protocol overhead, synchronization, and memory copies:

// Base latency (microseconds) per algorithm per protocol
// Protocols: [LL, LL128, Simple]
static const float baseLat[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = {
    {6.8,  14.0, 0},       // Tree
    {6.6,  14.0, 8.4},     // Ring
    {0,    0,    0},       // CollNet Direct
    {0,    0,    0},       // CollNet Chain
    {0,    0,    23.0},    // NVLS
    {0,    0,    0},       // NVLS-Tree
};

// Per-path latencies: [NVLINK, PCI, NET]
// Each indexed by [path][algorithm][protocol]
static float hwLat[3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = {
    // NVLINK path latencies
    { {0.6, 1.25, 28.0},   // Tree over NVLINK
      {0.6, 1.9,  3.4},    // Ring over NVLINK
      {0,   0,    0},      // CollNet Direct
      {0,   0,    0},      // CollNet Chain
      {0,   0,    0},      // NVLS
      {0,   0,    0} },    // NVLS-Tree

    // PCI path latencies
    { {1.0, 1.0,  0},      // Tree over PCI
      {1.0, 1.0,  0},      // Ring over PCI
      ... },

    // Network (NIC) path latencies
    { {5.0, 8.5,  0},      // Tree over NET
      {2.7, 5.5,  0},      // Ring over NET
      ... },
};

SingleFlow Data Structure

The SingleFlow struct is the fundamental unit of communication in SimAI. Every collective operation is decomposed into a directed acyclic graph (DAG) of SingleFlows. Each flow represents a point-to-point data transfer between two ranks, with explicit dependency tracking via prev and child_flow_id fields.

struct SingleFlow {
    int flow_id;                  // Unique identifier for this flow
    int src, dest;                // Source and destination rank
    uint64_t flow_size;           // Data size in bytes

    vector<int> prev;            // Dependency flows (must complete first)
    vector<int> parallel_deps;   // Flows that run in parallel
    vector<int> child_flow_id;   // Downstream flows (triggered on completion)

    int channel_id;               // Ring/tree channel index
    int chunk_id;                 // Which chunk of the collective
    int chunk_count;              // Total chunks in this collective
    string conn_type;             // "RING", "TREE", "NVLS", "PXN_INIT"

    // Constructor
    SingleFlow(int id, int s, int d, uint64_t size,
              vector<int> p, vector<int> par,
              vector<int> child,
              int ch, int ck, int cc,
              string ct)
        : flow_id(id), src(s), dest(d), flow_size(size),
          prev(p), parallel_deps(par), child_flow_id(child),
          channel_id(ch), chunk_id(ck), chunk_count(cc),
          conn_type(ct) {}
};

The flow DAG ensures correct execution ordering. For a Ring AllReduce with 4 GPUs, the prev field chains flows so that step 2 of the ReduceScatter cannot begin until step 1 completes. The child_flow_id enables the simulation engine to eagerly dispatch dependent flows once a flow finishes.

Group Topology

SimAI organizes GPU ranks into communication groups based on parallelism strategy. Each GroupInfo struct describes the membership, topology, and NVSwitch connectivity of a group. The GroupType enum determines which collective algorithm is used:

enum GroupType { TP, DP, PP, EP, DP_EP, NONE };

struct GroupInfo {
    GroupType type;            // Parallelism type
    int nNodes;                // Number of nodes in this group
    int nRanks;                // Number of GPU ranks
    vector<int> Ranks;        // List of rank IDs
    vector<int> NVSwitchs;     // NVSwitch IDs (if applicable)
};

// Example: 8-GPU TP group within one node
// GroupInfo { TP, 1, 8, {0,1,2,3,4,5,6,7}, {8,9,10,11} }
//   - 8 ranks on 1 node
//   - 4 NVSwitches (H100 SXM)

GroupType Communication Patterns

The group topology determines how MockNcclGroup generates ring or tree channels. For a TP group confined to one node, rings are constructed using the NVLink topology (e.g., rail-optimized patterns). For DP groups spanning multiple nodes, rings traverse NIC links and the network fabric.

astra-sim Three Backends

astra-sim supports three different network backends, each offering a different trade-off between simulation speed and fidelity. Users choose the backend at build time, and the same workload description works with all three.

5.1 Analytical Mode (AnaSim)

class AnaSim {
    static queue<struct CallTask> call_list;
    static uint64_t tick;

    static void Run() {
        while (!call_list.empty()) {
            CallTask task = call_list.front();
            while (tick != task.time) tick++;
            call_list.pop();
            task.fun_ptr(task.fun_arg);  // Execute callback
        }
    }

    static void Schedule(
        uint64_t delay,              // Computed from busbw config
        void (*fun_ptr)(void*),      // Completion callback
        void* fun_arg) {
        call_list.push({tick + delay, fun_ptr, fun_arg});
    }
};

// Transfer time = data_size / busbw + base_latency
// No congestion, no contention, no queue dynamics

5.2 NS-3 Simulation Mode

5.3 Physical Mode (Beta)

void send_flow(int src, int dst, uint64_t maxPacketCount,
    void (*msg_handler)(void*), void* fun_arg,
    int tag, sim_request* request) {

    ncclFlowTag flowtag = request->flowTag;
    TransportData send_data = TransportData(
        flowtag.channel_id,
        flowtag.chunk_id,
        flowtag.chunk_count,
        flowtag.current_flow_id,
        flowtag.child_flow_id,
        flowtag.tree_flow_list,
        flowtag.data_size);

    // Serialize and send via actual RDMA verbs
    rdma_transport->post_send(
        dst, &send_data, sizeof(send_data),
        tag, msg_handler, fun_arg);
}

Backend Comparison

P2P Transport Mechanics: How KV Cache Actually Travels

When SimAI simulates a KV cache transfer from a prefill node to a decode node, the NS-3 backend doesn't just estimate latency — it simulates the full RoCEv2 RDMA transport stack at packet granularity. This section traces the exact mechanisms: how RDMA Queue Pairs are created, how DCQCN adjusts sending rates on congestion, how PFC prevents buffer overflow, and how ECMP distributes flows across equal-cost paths.

6.1 RDMA Queue Pair (QP) Lifecycle

Every P2P flow in SimAI becomes one or more RDMA Queue Pairs. The call chain starts from astra-sim's collective decomposition and ends at NS-3's packet-level simulation:

class RdmaQueuePair : public Object {
    Ipv4Address sip, dip;             // Source & destination IP
    uint16_t sport, dport;             // Source & dest port (ECMP hash input)
    uint64_t m_size;                    // Total bytes to transfer (= KV cache size)
    uint16_t m_pg;                      // Priority group (queue index for PFC)
    uint64_t snd_nxt, snd_una;          // Next seq to send, unacked seq
    uint32_t m_win;                     // Window size (BDP-based)
    DataRate m_rate;                    // Current sending rate (DCQCN-controlled)
    uint64_t m_baseRtt;                 // Base RTT for this src-dst pair
};

6.2 DCQCN Congestion Control (CC_MODE=1)

SimAI's default congestion control is DCQCN (Data Center QCN), Mellanox's adaptation of QCN for RoCEv2. It operates in three phases: rate decrease on CNP, alpha tracking via EWMA, and rate recovery via additive/hyper-additive increase.

// DCQCN rate control state (per QP)
double m_g;                    // EWMA gain = 0.00390625 (1/256)
double m_rateOnFirstCNP;       // Rate fraction on first CNP (1.0 = full)
double m_rpgTimeReset;         // Rate increase timer = 900 μs
double m_rateDecreaseInterval;  // Rate decrease check interval = 4 μs
double m_alpha_resume_interval; // Alpha EWMA update interval = 1 μs
DataRate m_rai;                // Additive increase = 50 Mb/s
DataRate m_rhai;               // Hyper-additive increase = 100 Mb/s
uint32_t m_rpgThreshold;       // Fast recovery threshold = 1 CNP

Rate Decrease (on CNP received)

When a CNP (Congestion Notification Packet) arrives at the sender, DCQCN reduces the rate multiplicatively. The key formula:

// On each rate_decrease_interval (4 μs), if CNP was received:
target_rate = rate × (1 - alpha / 2)
rate = max(rate × (1 - alpha / 2), MIN_RATE)

// Alpha is updated via EWMA every alpha_resume_interval (1 μs):
// If CNP received in this interval:
alpha = (1 - m_g) × alpha + m_g        // α increases toward 1
// If no CNP:
alpha = (1 - m_g) × alpha              // α decays toward 0

// With m_g = 1/256, alpha converges slowly — providing stability

Rate Increase (recovery)

// Every RP_TIMER (900 μs), if no CNP received:
// Phase 1: Fast Recovery (first m_rpgThreshold rounds)
rate = (rate + target_rate) / 2

// Phase 2: Additive Increase
rate = rate + m_rai                    // +50 Mb/s per interval

// Phase 3: Hyper-Additive Increase (after sustained no-congestion)
rate = rate + m_rhai                   // +100 Mb/s per interval

6.3 ECN Marking at Switches (RED-based)

Switches mark packets with ECN (setting the ECE bits in IPv4 header) based on probabilistic RED (Random Early Detection). The marking probability increases linearly between kmin and kmax:

bool SwitchMmu::ShouldSendCN(uint32_t ifindex, uint32_t qIndex) {
    if (qIndex == 0) return false;          // Queue 0 = highest priority, never marked
    if (egress_bytes[ifindex][qIndex] > kmax[ifindex])
        return true;                           // Above kmax: always mark (100%)
    if (egress_bytes[ifindex][qIndex] > kmin[ifindex]) {
        // Between kmin and kmax: linear probability
        double p = pmax[ifindex]
                 * (double)(egress_bytes[ifindex][qIndex] - kmin[ifindex])
                 / (kmax[ifindex] - kmin[ifindex]);
        if (UniformVariable(0, 1).GetValue() < p)
            return true;                       // Probabilistic mark
    }
    return false;                              // Below kmin: never mark
}

The kmin/kmax/pmax values are rate-dependent, configured per link speed in SimAI.conf:

6.4 PFC (Priority Flow Control) — Layer 2 Back-Pressure

PFC operates at Layer 2, independently from DCQCN. When a switch's ingress buffer fills beyond the threshold, it sends a PAUSE frame to the upstream sender, halting transmission on that priority queue. This prevents packet loss but can cause head-of-line blocking and PFC storms.

bool SwitchMmu::CheckShouldPause(uint32_t port, uint32_t qIndex) {
    return !paused[port][qIndex] &&
           (hdrm_bytes[port][qIndex] > 0 ||                  // Headroom occupied
            GetSharedUsed(port, qIndex) >= GetPfcThreshold(port));  // Shared buffer full
}

bool SwitchMmu::CheckShouldResume(uint32_t port, uint32_t qIndex) {
    if (!paused[port][qIndex]) return false;
    return hdrm_bytes[port][qIndex] == 0 &&                 // Headroom drained
           (GetSharedUsed(port,qIndex) == 0 ||
            GetSharedUsed(port,qIndex) + resume_offset       // 3 KB hysteresis
                <= GetPfcThreshold(port));
}

// Switch buffer: 12 MB total (static) or 32 MB (SimAI.conf BUFFER_SIZE)
// Dynamic threshold: USE_DYNAMIC_PFC_THRESHOLD = 1 (enabled by default)
// 8 priority queues (qCnt=8), queue 0 = highest (ACK/NACK/CNP bypass PFC)

6.5 ECMP Routing (Per-Flow, MurmurHash3)

SimAI uses per-flow ECMP routing at every switch. When a packet arrives at a switch with multiple equal-cost next hops, the switch hashes the flow's 5-tuple to deterministically pick a path:

int SwitchNode::GetOutDev(Ptr<Packet> p, CustomHeader &ch) {
    // Extract 5-tuple: src_ip(4B) + dst_ip(4B) + src_port(2B) + dst_port(2B)
    union { uint8_t u8[12]; uint32_t u32[3]; } buf;
    buf.u32[0] = ch.sip;    buf.u32[1] = ch.dip;
    buf.u32[2] = ch.sport | ((uint32_t)ch.dport << 16);

    // MurmurHash3 with per-switch seed (= node ID)
    uint32_t idx = EcmpHash(buf.u8, 12, m_ecmpSeed) % nexthops.size();
    return nexthops[idx];     // Deterministic path for this flow
}

6.6 Flow Completion Time (FCT) Tracking

void qp_finish(FILE *fout, Ptr<RdmaQueuePair> q) {
    uint64_t standalone_fct = base_rtt
        + total_bytes * 8000000000lu / bandwidth;  // Ideal FCT (no congestion)

    fprintf(fout, "%08x %08x %u %u %lu %lu %lu %lu\n",
        q->sip, q->dip,             // Source & dest IP (hex)
        q->sport, q->dport,          // Ports
        q->m_size,                    // Data bytes (= KV cache size)
        q->startTime,                 // Start time (ns)
        Simulator::Now() - q->startTime,  // Actual FCT (ns)
        standalone_fct);              // Ideal FCT (ns, zero-congestion)
}
// FCT ratio = actual_fct / ideal_fct → measures congestion impact
// For PD disagg: this directly determines pd_p2p_comm_time

6.7 Transport Stack Summary

6.8 NACK and Layer-2 Retransmission

When a receiver detects an out-of-order sequence number, it generates a NACK packet (protocol 0xFD). The sender retransmits from the NACKed sequence. This provides a reliable transport layer beneath congestion control.

Key parameters from SimAI.conf:

• L2_CHUNK_SIZE = 4000 — chunk size for reliability (bytes)
• L2_ACK_INTERVAL = 1 — ACK every N chunks
• L2_BACK_TO_ZERO = 0 — Go-Back-N mode (0=selective, 1=go-back-to-zero)

// rdma-hw.cc — ReceiverCheckSeq()
int RdmaHw::ReceiverCheckSeq(uint32_t seq, Ptr<RdmaRxQueuePair> q, uint32_t size) {
    uint32_t expected = q->ReceiverNextExpectedSeq;
    if (seq == expected) {
        q->ReceiverNextExpectedSeq = expected + size;
        // Check if need to send ACK
        if (m_ack_interval == 0) return 1; // ACK
        else return 5; // delayed ACK
    } else if (seq > expected) {
        // Gap detected → generate NACK
        return 2; // NACK
    } else {
        // Duplicate or retransmitted packet
        return 4; // duplicate
    }
}

6.9 NVSwitch / NVLS Dedicated Routing

SimAI has a dedicated NVSwitchNode class (separate from regular SwitchNode) for intra-node NVLink routing. This enables distinct routing paths for GPU-to-GPU traffic within the same server versus inter-server traffic through NICs.

• Separate routing table: m_rtTable_nxthop_nvswitch for next-hop routing through NVSwitch
• m_gpus_per_server determines which GPUs are in the same server
• When nvls_enable=true on a QP, packets route through NVSwitch instead of NIC→ASW
• NVLS ACKs use special ToS=4 for identification
• SwitchAsHostSend() handles NVLS reply path

// rdma-hw.h
std::unordered_map<uint32_t, std::vector<int>> m_rtTable;              // Normal routing
std::unordered_map<uint32_t, std::vector<int>> m_rtTable_nxthop_nvswitch; // NVSwitch routing
uint32_t m_gpus_per_server;  // Determines intra-node boundary

// rdma-hw.cc — route selection
if (nvls_enable && IsInSameServer(sip, dip)) {
    // Use NVSwitch routing table
    nexthops = m_rtTable_nxthop_nvswitch[dip];
} else {
    // Use normal ECMP routing
    nexthops = m_rtTable[dip];
}

6.10 Window + Rate Dual Control

SimAI uses both window-based and rate-based flow control simultaneously. A packet is sent only if both conditions are met: (1) within window, AND (2) within rate limit.

• HAS_WIN=1 — enables window-based control
• VAR_WIN=1 — enables variable window that scales with current rate
• Formula: effective_window = m_win × (current_rate / max_rate)
• IsWinBound() checks if on-the-fly bytes exceed window

// rdma-queue-pair.cc
bool RdmaQueuePair::IsWinBound() {
    uint64_t on_the_fly = snd_nxt - snd_una;
    uint64_t win;
    if (m_var_win) {
        win = m_win * m_rate.GetBitRate() / m_max_rate.GetBitRate();
        if (win == 0) win = 1;
    } else {
        win = m_win;
    }
    return on_the_fly >= win;
}

6.11 Alternative Congestion Control: HPCC, TIMELY, DCTCP

Beyond DCQCN (CC_MODE=1), SimAI implements three more congestion control algorithms, each using different network signals:

HPCC (CC_MODE=3) — Uses INT (In-Network Telemetry) headers. Each switch adds an IntHop record (timestamp, bytes, queue length, line rate). The sender uses precise utilization info to set rate:

// IntHop structure (int-header.h)
struct IntHop {
    uint32_t time : 24;     // Timestamp (ns)
    uint32_t bytes : 20;    // Bytes transmitted since last sample
    uint32_t qlen : 17;     // Queue occupancy
    uint32_t lineRate : 3;  // Encoded link rate
};
// Max 5 hops per packet (IntHeader::maxHop = 5)
// Target utilization: m_targetUtil = 0.95

TIMELY (CC_MODE=7) — RTT-based, no ECN required:

// Parameters (rdma-hw.h)
double m_tmly_alpha;    // Rate decrease factor
double m_tmly_beta;     // Rate increase factor
uint64_t m_tmly_TLow;  // Low RTT threshold
uint64_t m_tmly_THigh; // High RTT threshold
uint64_t m_tmly_minRtt; // Minimum RTT tracking

DCTCP (CC_MODE=8) — ECN-based with additive increase:

DataRate m_dctcp_rai;  // Additive increase = 1000 Mb/s (from SimAI.conf DCTCP_RATE_AI)

Comparison of all supported congestion control algorithms: