diff --git a/.gitignore b/.gitignore index e4a9790f4..91d37ad8a 100644 --- a/.gitignore +++ b/.gitignore @@ -2,9 +2,27 @@ astra-sim-alibabacloud/build/simai_analytical/build/ astra-sim-alibabacloud/build/astra_ns3/build/ astra-sim-alibabacloud/extern/ +build-optimized/ bin/ results/ test/log/ *.log +*.bak +ncclFlowModel_*.csv +send.txt .cur* .DS_Store + +# SimCCL standalone build artifacts +SimCCL/standalone/build/ +SimCCL/standalone/build_*/ +SimCCL/tools/pat_simulator + +# Generated topology files (kept local, excluded from PR) +astra-sim-alibabacloud/inputs/topo/Rail_Opti_SingleToR_* +/Rail_Opti_SingleToR_* + +# Editor temp files +*.swp +*.swo +*~ diff --git a/CHANGELOG.md b/CHANGELOG.md new file mode 100644 index 000000000..375ba53d7 --- /dev/null +++ b/CHANGELOG.md @@ -0,0 +1,23 @@ +# Changelog + +All notable changes to SimAI are documented in this file. + +## [1.7] - 2026-07-07 + +- [SimCCL](https://github.com/aliyun/SimCCL) v2.30 mock: NCCL-style collective flow decomposition with Ring, PAT, and NVLS algorithm support. +- Protocol-aware selection (LL/LL128/Simple based on message size). +- SimCCL standalone binary for independent collective operation analysis (no GPU required). +- Per-(algorithm, protocol, link_type) send latency table for higher simulation fidelity. + +## [1.6] - 2026-04-23 + +- GPU memory modeling for inference simulation (parameter counting & KV cache). +- Linear interpolation for decode time estimation (replacing nearest-neighbor). +- Prefill-Decode Disaggregation memory planning (independent budgets for Prefill/Decode). + +## [1.5] - 2025-12-30 + +- End-to-end simulation for multi-request **inference** workloads. +- Advanced Inference Simulation: Model complex scenarios with Prefill/Decode separation. +- Modern Model Support: DeepSeek, Qwen3Moe and Qwen3Next. +- Request Scheduling: Adapted from Microsoft's [Vidur](https://github.com/microsoft/vidur). diff --git a/Dockerfile b/Dockerfile index 8166a9a90..ea6c5f933 100644 --- a/Dockerfile +++ b/Dockerfile @@ -8,8 +8,8 @@ RUN apt-get update && apt-get install -y git && rm -rf /var/lib/apt/lists/* WORKDIR /workspace/SimAI # [Optional] Configure pip and uv to use Aliyun mirror for faster package downloads. -RUN pip config set global.index-url http://mirrors.aliyun.com/pypi/simple -RUN pip config set install.trusted-host mirrors.aliyun.com +# Use HTTPS to avoid MITM during dependency download (no trusted-host needed with TLS). +RUN pip config set global.index-url https://mirrors.aliyun.com/pypi/simple ENV UV_DEFAULT_INDEX="https://mirrors.aliyun.com/pypi/simple" RUN pip install --no-cache-dir uv @@ -26,7 +26,11 @@ RUN UV_TORCH_BACKEND=auto uv pip install -v --system --no-cache-dir --no-build-i # Copy the rest of the application source code into the image. COPY . . -RUN mv ./workload_generator /usr/local/lib/python3.12/dist-packages &&\ - mv ./utils /usr/local/lib/python3.12/dist-packages &&\ - mv ./log_analyzer /usr/local/lib/python3.12/dist-packages +# Move helper packages onto the Python site-packages path. Compute the path +# dynamically so this does not break if the base image's Python minor version +# or site-packages location changes. +RUN SITE_PACKAGES="$(python3 -c 'import site; print(site.getsitepackages()[0])')" &&\ + mv ./workload_generator "$SITE_PACKAGES" &&\ + mv ./utils "$SITE_PACKAGES" &&\ + mv ./log_analyzer "$SITE_PACKAGES" ENV PYTHONPATH=/workspace/SimAI:/workspace/SimAI/aicb:/workspace/SimAI/vidur:$PYTHONPATH \ No newline at end of file diff --git a/README.ja.md b/README.ja.md index 5539f64cc..cd4da1407 100644 --- a/README.ja.md +++ b/README.ja.md @@ -11,6 +11,12 @@ ### 最近のアップデート +- [2026/07] **SimAI 1.7 リリース！** 主な更新： + - [SimCCL](https://github.com/aliyun/SimCCL) v2.30 mock：Ring、PAT、NVLS アルゴリズムをサポートする NCCL スタイルの集合通信フロー分解。 + - プロトコル認識選択（メッセージサイズに基づく LL/LL128/Simple の自動選択）。 + - SimCCL スタンドアロンバイナリ（GPU不要で集合操作分析が可能）。 + - (アルゴリズム, プロトコル, リンクタイプ) 別の送信レイテンシテーブルによるシミュレーション精度向上。 + - [2026/04] **SimAI 1.6 リリース！** 主な更新： - 推論シミュレーション向け GPU メモリモデリング（パラメータカウント＆KV Cache）。 - Decode 時間推定の線形補間（最近傍探索の代替）。 @@ -44,6 +50,7 @@ | 日付 | イベント | 場所 | 内容 | 形式 | |:----------------:|:------------------------------------------------------------------------ |:----------------------- |:-------------------------------------------------------- |:-------------:| +| 2026年7月7日 | SimAI 1.7 | 🌐 オンライン | SimAI 1.7 のリリース | 💻 バーチャル | | 2026年4月23日 | SimAI 1.6 | 🌐 オンライン | SimAI 1.6 のリリース | 💻 バーチャル | | 2025年12月30日 | SimAI 1.5 | 🌐 オンライン | SimAI 1.5 のリリース | 💻 バーチャル | | 2025年6月4日 | SimAIコミュニティ第1回ワークショップ | 📍 北京大学 | コミュニティ貢献者による3つの講演 | 🎓 現地 | @@ -77,6 +84,7 @@ - [セットアップ](#セットアップ) - [SimAI-Analyticalの使い方](#simai-analyticalの使い方) - [SimAI-Simulationの使い方](#simai-simulationの使い方) + - [SimCCL の使い方](#simccl-の使い方) - [マルチリクエスト推論シミュレーションの使い方](#マルチリクエスト推論シミュレーションの使い方) # SimAI 概要 @@ -118,7 +126,7 @@ astra-sim-alibabacloudは[astra-sim](https://github.com/astra-sim/astra-sim/tree SimAIは、さまざまなシミュレーション要件を満たすために、3つの主要な動作モードをサポートしています： -**SimAI-Analytical**は、バス帯域幅（busbw）を使用して集合通信時間を見積もることにより、ネットワーク通信の詳細を抽象化し、高速なシミュレーションを提供します。現在、ユーザー定義のbusbwをサポートしていますが、自動busbw計算機能はまもなく登場予定です。 +**SimAI-Analytical**は、バス帯域幅（busbw）を使用して集合通信時間を見積もることにより、ネットワーク通信の詳細を抽象化し、高速なシミュレーションを提供します。現在のオープンソース版バイナリは `-nv`/`-nic`/`-n_p_s` から busbw を自動計算します（デフォルト、下記参照）。注意：ユーザー定義の `busbw.yaml`（`-busbw`）は現在のオープンソース版 analytical バイナリでは未接続です。この方法は以前の SimAI バージョンを参照してください。 **SimAI-Simulation**は、きめ細かいネットワーク通信モデリングを備えたフルスタックシミュレーションを提供します。NS3や他のネットワークシミュレータ（現在はNS3がオープンソース化されています）を活用して、すべての通信動作を詳細にシミュレーションし、実際のトレーニング環境を高忠実に再現することを目指しています。 @@ -146,7 +154,7 @@ SimAIを基盤とした革新的な研究や拡張を奨励します。ディス # クイックスタート -以下に簡単な例を示します。SimAIの完全なチュートリアルはこちらにあります：[**SimAI@Tutorial**](./docs/Tutorial.md)、[**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md)、[SimCCL@Tutorial]、[ns-3-alibabacloud@Tutorial] +以下に簡単な例を示します。SimAIの完全なチュートリアルはこちらにあります：[**SimAI@Tutorial**](./docs/Tutorial.md)、[**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md)、[**SimCCL@Tutorial**](./SimCCL/README.md)、[**ns-3-alibabacloud@Tutorial**](https://github.com/aliyun/ns-3-alibabacloud) ## セットアップ @@ -179,14 +187,16 @@ $ ./scripts/build.sh -c ns3 ## SimAI-Analyticalの使い方 +デフォルトでは、SimAI-Analytical はバス帯域幅を自動的に計算します： + ```bash -$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml +$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- ``` -バス帯域幅を自動で計算するには、次のコマンドを試してください： +> 注意：以下の `-busbw example/busbw.yaml`（ユーザー定義 busbw）コマンドは、現在のオープンソース版 analytical バイナリでは**サポートされていません**。`-busbw` は解析されず、コマンドは使用法を表示して終了します。参考のために残しています。手動の `busbw.yaml` の使い方は以前の SimAI バージョンを参照してください。 ```bash -$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- +$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml ``` ## SimAI-Simulationの使い方 @@ -199,6 +209,28 @@ $ python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo Spectr $ AS_SEND_LAT=3 AS_NVLS_ENABLE=1 ./bin/SimAI_simulator -t 16 -w ./example/microAllReduce.txt -n ./Spectrum-X_128g_8gps_100Gbps_A100 -c astra-sim-alibabacloud/inputs/config/SimAI.conf ``` +## SimCCL の使い方 + +SimCCL は、フルネットワークシミュレーションを実行せずに、集合通信フロー分解を独立して分析できます： + +```bash +# SimCCL スタンドアロンバイナリのビルド +$ cd SimCCL/standalone +$ bash build.sh v2.30 + +# 単一集合操作の分析 +$ ./build/simccl-standalone --op AllReduce --size 4194304 \ + --nRanks 8 --nNodes 1 --gpus_per_node 8 --gpu_type H20 + +# ワークロードファイルモード +$ ./build/simccl-standalone -w ../../example/microAllReduce.txt \ + --nRanks 16 --nNodes 2 --gpus_per_node 8 --gpu_type H20 +``` + +出力：`ncclFlowModel_detailed_flows.csv` — 集合操作のポイントツーポイントフロー分解。 + +詳細なドキュメントについては [SimCCL README](./SimCCL/README.md) を参照してください。 + ## マルチリクエスト推論シミュレーションの使い方 詳細については、`vidur-alibabacloud` ディレクトリ内の [README](./vidur-alibabacloud/README.md) ファイルを参照してください。このモジュールは AICB を活用して**推論**ワークロードの計算時間をプロファイリングします。DeepGEMM や FlashMLA などの特定のハードウェアアクセラレーションライブラリに依存するため、**Hopper (SM90)** および **Blackwell (SM100)** アーキテクチャベースの NVIDIA GPU のみに対応しています。 diff --git a/README.md b/README.md index 9b8e898c5..8df100c5a 100755 --- a/README.md +++ b/README.md @@ -11,6 +11,12 @@ ### Recent Updates +- [2026/07] **SimAI 1.7 Released!** Key updates: + - [SimCCL](https://github.com/aliyun/SimCCL) v2.30 mock: NCCL-style collective flow decomposition with Ring, PAT, and NVLS algorithm support. + - Protocol-aware selection (LL/LL128/Simple based on message size). + - SimCCL standalone binary for independent collective operation analysis (no GPU required). + - Per-(algorithm, protocol, link_type) send latency table for higher simulation fidelity. + - [2026/04] **SimAI 1.6 Released!** Key updates: - GPU memory modeling for inference simulation (parameter counting & KV cache). - Linear interpolation for decode time estimation (replacing nearest-neighbor). @@ -44,6 +50,7 @@ | Date | Event | Location | Content | Type | |:----------------:|:------------------------------------------------------------------------ |:----------------------- |:-------------------------------------------------------- |:-------------:| +| Jul 7, 2026 | SimAI 1.7 | 🌐 Online | The release of SimAI 1.7 | 💻 Virtual | | Apr 23, 2026 | SimAI 1.6 | 🌐 Online | The release of SimAI 1.6 | 💻 Virtual | | Dec 30, 2025 | SimAI 1.5 | 🌐 Online | The release of SimAI 1.5 | 💻 Virtual | | Jun 4, 2025 | The first workshop of the SimAI community | 📍 Peking University | Three talks from community contributors | 🎓 On-site | @@ -77,6 +84,7 @@ See [Tutorial](./docs/Tutorial.md) for full documentation. - [Setup](#setup) - [Use SimAI-Analytical](#use-simai-analytical) - [Use SimAI-Simulation](#use-simai-simulation) + - [Use SimCCL](#use-simccl) - [Use Multi-requests Inference Simulation](#use-multi-requests-inference-simulation) # SimAI Overview @@ -118,7 +126,7 @@ astra-sim-alibabacloud is extended from [astra-sim](https://github.com/astra-sim SimAI supports three major operation modes to meet different simulation requirements: -**SimAI-Analytical** offers fast simulation by abstracting network communication details using bus bandwidth (busbw) to estimate collective communication time. While it currently supports user-defined busbw, automatic busbw calculation feature is coming soon. +**SimAI-Analytical** offers fast simulation by abstracting network communication details using bus bandwidth (busbw) to estimate collective communication time. The current open-source binary computes busbw automatically from `-nv`/`-nic`/`-n_p_s` (the default, shown below). Note: the user-defined `busbw.yaml` path (`-busbw`) is not wired in the current open-source analytical binary; refer to earlier SimAI versions for that workflow. **SimAI-Simulation** provides full-stack simulation with fine-grained network communication modeling. It leverages NS3 or other network simulators (NS3 currently open-sourced) to achieve detailed simulation of all communication behaviors, aiming for high-fidelity reproduction of actual training environments. @@ -146,7 +154,7 @@ We encourage innovative research and extensions based on SimAI. Welcome to join # Quick Start -Here are some simple examples. SimAI full tutorials can be found here: [**SimAI@Tutorial**](./docs/Tutorial.md), [**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md), [SimCCL@Tutorial], [ns-3-alibabacloud@Tutorial] +Here are some simple examples. SimAI full tutorials can be found here: [**SimAI@Tutorial**](./docs/Tutorial.md), [**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md), [**SimCCL@Tutorial**](./SimCCL/README.md), [**ns-3-alibabacloud@Tutorial**](https://github.com/aliyun/ns-3-alibabacloud) ## Setup @@ -179,14 +187,16 @@ $ ./scripts/build.sh -c ns3 ## Use SimAI-Analytical +By default, SimAI-Analytical calculates bus bandwidth automatically: + ```bash -$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml +$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- ``` -For calculating bus bandwidth automatically, please try the following command: +> Note: The `-busbw example/busbw.yaml` (user-defined busbw) command below is **not supported** by the current open-source analytical binary — the `-busbw` flag is not parsed, so the command prints usage and exits. It is kept here for reference; see earlier SimAI versions for the manual `busbw.yaml` workflow. ```bash -$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- +$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml ``` ## Use SimAI-Simulation @@ -199,6 +209,28 @@ $ python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo Spectr $ AS_SEND_LAT=3 AS_NVLS_ENABLE=1 ./bin/SimAI_simulator -t 16 -w ./example/microAllReduce.txt -n ./Spectrum-X_128g_8gps_100Gbps_A100 -c astra-sim-alibabacloud/inputs/config/SimAI.conf ``` +## Use SimCCL + +SimCCL can be used independently to analyze collective communication flow decomposition without running full network simulation: + +```bash +# Build SimCCL standalone binary +$ cd SimCCL/standalone +$ bash build.sh v2.30 + +# Single collective operation analysis +$ ./build/simccl-standalone --op AllReduce --size 4194304 \ + --nRanks 8 --nNodes 1 --gpus_per_node 8 --gpu_type H20 + +# Workload file mode +$ ./build/simccl-standalone -w ../../example/microAllReduce.txt \ + --nRanks 16 --nNodes 2 --gpus_per_node 8 --gpu_type H20 +``` + +Output: `ncclFlowModel_detailed_flows.csv` — point-to-point flow decomposition of collective operations. + +For detailed documentation, see [SimCCL README](./SimCCL/README.md). + ## Use Multi-requests Inference Simulation For detailed information, please refer to the [README](./vidur-alibabacloud/README.md) file in the `vidur-alibabacloud` directory. This module leverages AICB to profile the computation time of **inference** workloads. Due to its reliance on specific hardware-accelerated libraries like DeepGEMM and FlashMLA, it is exclusively compatible with NVIDIA GPUs based on the **Hopper (SM90)** and **Blackwell (SM100)** architectures. diff --git a/README_CN.md b/README_CN.md index 357afc4d2..bb6dfeb1d 100644 --- a/README_CN.md +++ b/README_CN.md @@ -11,6 +11,12 @@ ### 近期更新 +- [2026/07] **SimAI 1.7 正式发布！** 主要更新： + - [SimCCL](https://github.com/aliyun/SimCCL) v2.30 mock：NCCL 风格的集合通信流分解，支持 Ring、PAT、NVLS 算法。 + - 协议感知选择（基于消息大小自动选择 LL/LL128/Simple）。 + - SimCCL 独立二进制，无需 GPU 即可进行集合操作分析。 + - 按 (算法, 协议, 链路类型) 的 send latency 分桶表，提升仿真保真度。 + - [2026/04] **SimAI 1.6 正式发布！** 主要更新： - 推理仿真 GPU 显存建模（参数计数与 KV Cache 管理）。 - Decode 耗时线性插值估算（替代最近邻查找）。 @@ -43,6 +49,7 @@ | 日期 | 活动 | 地点 | 内容 | 形式 | |:----------------:|:------------------------------------------------------------------------ |:----------------------- |:-------------------------------------------------------- |:-------------:| +| Jul 7, 2026 | SimAI 1.7 | 🌐 线上 | SimAI 1.7 正式发布 | 💻 线上直播 | | Apr 23, 2026 | SimAI 1.6 | 🌐 线上 | SimAI 1.6 正式发布 | 💻 线上直播 | | Dec 30, 2025 | SimAI 1.5 | 🌐 线上 | SimAI 1.5 正式发布 | 💻 线上直播 | | Jun 4, 2025 | SimAI 社区第一届研讨会 | 📍 北京大学 | 三场社区贡献者演讲 | 🎓 线下 | @@ -76,6 +83,7 @@ - [环境搭建](#环境搭建) - [使用 SimAI-Analytical](#使用-simai-analytical) - [使用 SimAI-Simulation](#使用-simai-simulation) + - [使用 SimCCL](#使用-simccl) - [使用多请求推理仿真](#使用多请求推理仿真) # SimAI 概述 @@ -117,7 +125,7 @@ astra-sim-alibabacloud 基于 [astra-sim](https://github.com/astra-sim/astra-sim SimAI 支持三种主要运行模式： -**SimAI-Analytical** 通过使用总线带宽（busbw）抽象网络通信细节来估算集合通信时间，实现快速仿真。目前支持用户自定义 busbw，自动计算 busbw 功能即将推出。 +**SimAI-Analytical** 通过使用总线带宽（busbw）抽象网络通信细节来估算集合通信时间，实现快速仿真。当前开源二进制通过 `-nv`/`-nic`/`-n_p_s` 自动计算 busbw（默认方式，见下文）。注意：用户自定义 `busbw.yaml`（`-busbw`）在当前开源 analytical 二进制中尚未接线，该用法请参考早期 SimAI 版本。 **SimAI-Simulation** 提供基于细粒度网络通信建模的全栈仿真。利用 NS-3 或其他网络模拟器（当前 NS-3 已开源）实现对所有通信行为的详细仿真，力求高保真还原真实训练环境。 @@ -145,7 +153,7 @@ SimAI 论文已被 NSDI'25 Spring 接收，详情请参阅： # 快速开始 -以下为简单示例。完整教程请参见：[**SimAI@Tutorial**](./docs/Tutorial.md)、[**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md)、[SimCCL@Tutorial]、[ns-3-alibabacloud@Tutorial] +以下为简单示例。完整教程请参见：[**SimAI@Tutorial**](./docs/Tutorial.md)、[**aicb@Tutorial**](https://github.com/aliyun/aicb/blob/master/training/tutorial.md)、[**SimCCL@Tutorial**](./SimCCL/README.md)、[**ns-3-alibabacloud@Tutorial**](https://github.com/aliyun/ns-3-alibabacloud) ## 环境搭建 @@ -178,14 +186,16 @@ $ ./scripts/build.sh -c ns3 ## 使用 SimAI-Analytical +默认情况下，SimAI-Analytical 自动计算总线带宽： + ```bash -$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml +$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- ``` -若需自动计算总线带宽，请尝试： +> 注意：下面的 `-busbw example/busbw.yaml`（用户自定义 busbw）命令在当前开源 analytical 二进制中**不受支持**——`-busbw` 参数不会被解析，命令会打印用法并退出。此处仅作保留参考；手动 `busbw.yaml` 用法请参考早期 SimAI 版本。 ```bash -$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- +$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml ``` ## 使用 SimAI-Simulation @@ -198,6 +208,28 @@ $ python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo Spectr $ AS_SEND_LAT=3 AS_NVLS_ENABLE=1 ./bin/SimAI_simulator -t 16 -w ./example/microAllReduce.txt -n ./Spectrum-X_128g_8gps_100Gbps_A100 -c astra-sim-alibabacloud/inputs/config/SimAI.conf ``` +## 使用 SimCCL + +SimCCL 可独立使用，无需运行完整网络仿真即可分析集合通信流分解： + +```bash +# 编译 SimCCL 独立二进制 +$ cd SimCCL/standalone +$ bash build.sh v2.30 + +# 单个集合操作分析 +$ ./build/simccl-standalone --op AllReduce --size 4194304 \ + --nRanks 8 --nNodes 1 --gpus_per_node 8 --gpu_type H20 + +# 工作负载文件模式 +$ ./build/simccl-standalone -w ../../example/microAllReduce.txt \ + --nRanks 16 --nNodes 2 --gpus_per_node 8 --gpu_type H20 +``` + +输出：`ncclFlowModel_detailed_flows.csv` — 集合操作的点对点流分解。 + +详细文档请参见 [SimCCL README](./SimCCL/README.md)。 + ## 使用多请求推理仿真 详情请参见 `vidur-alibabacloud` 目录下的 [README](./vidur-alibabacloud/README_CN.md)。该模块利用 AICB 对**推理**工作负载的计算时间进行 profiling。由于依赖 DeepGEMM 和 FlashMLA 等特定硬件加速库，目前仅兼容基于 **Hopper（SM90）** 和 **Blackwell（SM100）** 架构的 NVIDIA GPU。 diff --git a/astra-sim-alibabacloud/CMakeLists.txt b/astra-sim-alibabacloud/CMakeLists.txt index 1873afa99..145d4d290 100755 --- a/astra-sim-alibabacloud/CMakeLists.txt +++ b/astra-sim-alibabacloud/CMakeLists.txt @@ -1,5 +1,12 @@ set(use_rdma ${USE_RDMA}) set(use_analytical ${USE_ANALYTICAL}) +# SimCCL mock version selection (matches ns3 build default) +if(DEFINED ENV{SIMAI_NCCL_VERSION}) + set(MOCK_VERSION $ENV{SIMAI_NCCL_VERSION}) +else() + set(MOCK_VERSION "v2.30") +endif() + file(GLOB astra_SRC "${PROJECT_SOURCE_DIR}/../../astra-sim/system/collective/*.cc" "${PROJECT_SOURCE_DIR}/../../astra-sim/system/fast-backend/*.cc" @@ -9,6 +16,7 @@ file(GLOB astra_SRC "${PROJECT_SOURCE_DIR}/../../astra-sim/system/*.cc" "${PROJECT_SOURCE_DIR}/../../astra-sim/system/*.cc" "${PROJECT_SOURCE_DIR}/../../astra-sim/workload/*.cc" + "${PROJECT_SOURCE_DIR}/../../../SimCCL/src/mock/${MOCK_VERSION}/*.cc" ) file(GLOB HEADERS "${PROJECT_SOURCE_DIR}/../../astra-sim/system/collective/*.hh" diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/analytical/AnalyticalAstra.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/analytical/AnalyticalAstra.cc index 48fa0da1b..035726358 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/analytical/AnalyticalAstra.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/analytical/AnalyticalAstra.cc @@ -19,7 +19,7 @@ #include #include "astra-sim/system/Sys.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "astra-sim/system/AstraComputeAPI.hh" #include "astra-sim/system/AstraParamParse.hh" diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/AstraSimNetwork.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/AstraSimNetwork.cc index 9ad3ff938..c5d850d58 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/AstraSimNetwork.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/AstraSimNetwork.cc @@ -17,7 +17,7 @@ #include "astra-sim/system/Sys.hh" #include "astra-sim/system/RecvPacketEventHadndlerData.hh" #include "astra-sim/system/Common.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "ns3/applications-module.h" #include "ns3/core-module.h" #include "ns3/csma-module.h" diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h b/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h index cb43eb8c7..8537176aa 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h @@ -45,8 +45,8 @@ #include "ns3/mtp-interface.h" #endif #include -#include"astra-sim/system/MockNcclQps.h" -#include "astra-sim/system/MockNcclLog.h" +#include"SimCCL/mock/MockNcclQps.h" +#include "SimCCL/mock/MockNcclLog.h" using namespace ns3; using namespace std; @@ -125,7 +125,47 @@ void SendFlow(int src, int dst, uint64_t maxPacketCount, int flow_id = request->flowTag.current_flow_id; bool nvls_on = request->flowTag.nvls_on; int pg = 3, dport = 100; - int send_lat = 6000; + + // send_lat bucketing: select latency based on (algorithm, protocol) from NCCL v2.30 tuning. + // Formula: send_lat_ns = (baseLatencies[algo][proto] + hwLatencies[hw][algo][proto]) * 1000 + // Source: nccl-2.30/src/graph/tuning.cc L150-174 + // Table value 0 = unsupported (algo, proto) combo → use default fallback. + // NVLINK table: for same-node (intra-node) communication + static const int send_lat_table_nvlink[7][3] = { + /* Tree: baseLat(6.8,14,8.4)+hwLat(0.6,1.25,4) */ {7400, 15250, 12400}, + /* Ring: baseLat(6.6,14,8.4)+hwLat(0.6,1.9,3.4) */ {7200, 15900, 11800}, + /* CollNetDir: baseLat(0)+hwLat(0,0,3.7) */ {0, 0, 3700}, + /* CollNetChn: baseLat(0)+hwLat(0,0,2.8) */ {0, 0, 2800}, + /* NVLS: baseLat(0)+hwLat(0,0,25) */ {0, 0, 25000}, + /* NVLS_TREE: baseLat(0)+hwLat(0,0,25) */ {0, 0, 25000}, + /* PAT: baseLat(8)+hwLat(0,0,4), Simple-only */ {0, 0, 12000}, + }; + // NET table: for cross-node (inter-node) communication + static const int send_lat_table_net[7][3] = { + /* Tree: baseLat(6.8,14,8.4)+hwLat_NET(5,8.5,14) */ {11800, 22500, 22400}, + /* Ring: baseLat(6.6,14,8.4)+hwLat_NET(2.7,4,14) */ {9300, 18000, 22400}, + /* CollNetDir: baseLat(0)+hwLat_NET(0,0,31) */ {0, 0, 31000}, + /* CollNetChn: baseLat(0)+hwLat_NET(0,0,30) */ {0, 0, 30000}, + /* NVLS: baseLat(0)+hwLat_NET(0,0,18) */ {0, 0, 18000}, + /* NVLS_TREE: baseLat(0)+hwLat_NET(0,0,20.9) */ {0, 0, 20900}, + /* PAT: baseLat(8)+hwLat_NET(0,0,14), Simple-only */ {0, 0, 22000}, + }; + + int send_lat = 6000; // default fallback (6 μs) + int algo = request->flowTag.algorithm; + int proto = request->flowTag.protocol; + if (algo >= 0 && algo < 7 && proto >= 0 && proto < 3) { + // Dynamic link type detection: use gpus_per_node from flowTag (set by NcclFlowModel) + int gpus_per_node = request->flowTag.gpus_per_node; + if (gpus_per_node <= 0) gpus_per_node = 8; // fallback for v2.20 / backward compat + bool same_node = (src / gpus_per_node) == (dst / gpus_per_node); + const int (*table)[3] = same_node ? send_lat_table_nvlink : send_lat_table_net; + int table_val = table[algo][proto]; + if (table_val > 0) { // 0 = unsupported combo, keep default + send_lat = table_val; + } + } + // AS_SEND_LAT env var overrides everything (highest priority, for A/B experiments) const char* send_lat_env = std::getenv("AS_SEND_LAT"); if (send_lat_env) { try { @@ -135,7 +175,7 @@ void SendFlow(int src, int dst, uint64_t maxPacketCount, exit(-1); } } - send_lat *= 1000; + send_lat *= 1000; // ns → ps (ns3 Time unit) flow_input.idx++; if(real_PacketCount == 0) real_PacketCount = 1; MockNcclLog* NcclLog = MockNcclLog::getInstance(); diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/PhySimAi.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/PhySimAi.cc index fd9afb891..f46f240b5 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/PhySimAi.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/PhySimAi.cc @@ -15,7 +15,7 @@ #include #include"PhySimAi.h" -#include"astra-sim/system/MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" using namespace std; queue PhyNetSim::call_list = {}; diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.cc index 20033dba3..64e491818 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.cc @@ -13,7 +13,7 @@ *limitations under the License. */ -#include"astra-sim/system/MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" #ifdef PHY_RDMA #include"astra-sim/system/SimAiFlowModelRdma.hh" #endif diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.h b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.h index 7cd1b07e2..2d1346f60 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.h +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiEntry.h @@ -31,7 +31,7 @@ #include"astra-sim/system/RecvPacketEventHadndlerData.hh" #include"astra-sim/system/AstraNetworkAPI.hh" -#include"astra-sim/system/MockNcclQps.h" +#include"SimCCL/mock/MockNcclQps.h" #include"astra-sim/system/SimAiPhyCommon.hh" using namespace std; diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiMain.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiMain.cc index 6641078ac..8e17c71dc 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiMain.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiMain.cc @@ -23,7 +23,7 @@ #include "astra-sim/system/AstraComputeAPI.hh" #include "astra-sim/system/Sys.hh" #include "astra-sim/system/RecvPacketEventHadndlerData.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "astra-sim/system/BootStrapnet.hh" #include "astra-sim/system/PhyMultiThread.hh" #include "astra-sim/system/Common.hh" diff --git a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiPhyNetwork.cc b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiPhyNetwork.cc index e5e147130..93035059d 100644 --- a/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiPhyNetwork.cc +++ b/astra-sim-alibabacloud/astra-sim/network_frontend/phynet/SimAiPhyNetwork.cc @@ -13,7 +13,7 @@ *limitations under the License. */ -#include"astra-sim/system/MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" #include"astra-sim/system/PhyMultiThread.hh" #include"SimAiPhyNetwork.h" diff --git a/astra-sim-alibabacloud/astra-sim/system/AstraNetworkAPI.hh b/astra-sim-alibabacloud/astra-sim/system/AstraNetworkAPI.hh index 324f9d7a8..2013b4488 100644 --- a/astra-sim-alibabacloud/astra-sim/system/AstraNetworkAPI.hh +++ b/astra-sim-alibabacloud/astra-sim/system/AstraNetworkAPI.hh @@ -38,6 +38,15 @@ struct ncclFlowTag { int tag_id; std::vector tree_flow_list; bool nvls_on; + // B-Task1 (calib round2, 260611): algorithm/protocol carried alongside the flow. + // Default values (-1 = UNDEF) preserve all existing constructor calls (PhyMultiThread, + // SimAiFlowModelRdma, NcclFlowModel) without changes; populated only when + // AS_PROTO_AWARE=1 (see MockNcclGroup.cc B-Task2). Enables future bucketed + // AS_SEND_LAT (B-Task3) and proto-aware busbw saturation (B-Task4); ns3 datapath + // does NOT yet consume these fields (planned, see report). + int algorithm = -1; // NCCL_ALGO_UNDEF + int protocol = -1; // NCCL_PROTO_UNDEF + int gpus_per_node = -1; // for link type detection (NVLINK vs NET); -1 = use fallback 8 ncclFlowTag(): channel_id(-1), chunk_id(-1), diff --git a/astra-sim-alibabacloud/astra-sim/system/Common.hh b/astra-sim-alibabacloud/astra-sim/system/Common.hh index 35ec6d00e..df2762ac6 100644 --- a/astra-sim-alibabacloud/astra-sim/system/Common.hh +++ b/astra-sim-alibabacloud/astra-sim/system/Common.hh @@ -25,7 +25,8 @@ enum class ComType { All_Reduce, All_to_All, All_Reduce_All_to_All, - All_Reduce_NVLS + All_Reduce_NVLS, + Broadcast }; enum class CollectiveOptimization { Baseline, LocalBWAware }; enum class CollectiveImplementationType { @@ -41,7 +42,6 @@ enum class CollectiveImplementationType { HalvingDoubling, OneHalvingDoubling, NcclFlowModel, - NcclTreeFlowModel, }; enum class CollectiveBarrier { Blocking, Non_Blocking }; enum class SchedulingPolicy { LIFO, FIFO, HIGHEST, None }; diff --git a/astra-sim-alibabacloud/astra-sim/system/DataSet.cc b/astra-sim-alibabacloud/astra-sim/system/DataSet.cc index 09dd66164..945acfe15 100644 --- a/astra-sim-alibabacloud/astra-sim/system/DataSet.cc +++ b/astra-sim-alibabacloud/astra-sim/system/DataSet.cc @@ -6,7 +6,7 @@ LICENSE file in the root directory of this source tree. #include "DataSet.hh" #include "IntData.hh" #include "Sys.hh" -#include "MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" namespace AstraSim { int DataSet::id_auto_increment = 0; DataSet::DataSet(int total_streams) { diff --git a/astra-sim-alibabacloud/astra-sim/system/MemBus.cc b/astra-sim-alibabacloud/astra-sim/system/MemBus.cc index f5971e41b..805b72b75 100644 --- a/astra-sim-alibabacloud/astra-sim/system/MemBus.cc +++ b/astra-sim-alibabacloud/astra-sim/system/MemBus.cc @@ -6,7 +6,7 @@ LICENSE file in the root directory of this source tree. #include "MemBus.hh" #include "LogGP.hh" #include "Sys.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" namespace AstraSim { MemBus::~MemBus() { delete NPU_side; diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNccl.h b/astra-sim-alibabacloud/astra-sim/system/MockNccl.h deleted file mode 100644 index 84f369b40..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNccl.h +++ /dev/null @@ -1,225 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ - -#ifndef ASTRA_SIM_MOCKNCCL_MOCKNCCL_H -#define ASTRA_SIM_MOCKNCCL_MOCKNCCL_H - -#define NCCL_NUM_ALGORITHMS 6 // Tree/Ring/CollNet* -#define NCCL_ALGO_UNDEF -1 -#define NCCL_ALGO_TREE 0 -#define NCCL_ALGO_RING 1 -#define NCCL_ALGO_COLLNET_DIRECT 2 -#define NCCL_ALGO_COLLNET_CHAIN 3 -#define NCCL_ALGO_NVLS 4 -#define NCCL_ALGO_NVLS_TREE 5 - -#define NCCL_NUM_PROTOCOLS 3 // Simple/LL/LL128 -#define NCCL_PROTO_UNDEF -1 -#define NCCL_PROTO_LL 0 -#define NCCL_PROTO_LL128 1 -#define NCCL_PROTO_SIMPLE 2 -#define NCCL_WORK_SIZE 512 - -/* Array indexes used below */ -#define VOLTA_COMPCAP_IDX 0 -#define AMPERE_COMPCAP_IDX 1 -#define HOPPER_COMPCAP_IDX 2 -#define NCCL_TOPO_CPU_VENDOR_AMD 2 - - -#define NCCL_TOPO_CPU_ARCH_X86 1 -#define NCCL_TOPO_CPU_ARCH_POWER 2 -#define NCCL_TOPO_CPU_VENDOR_INTEL 1 -#define NCCL_TOPO_CPU_VENDOR_AMD 2 - -#define NCCL_TOPO_PATTERN_BALANCED_TREE 1 // 目前正在使用的 Spread NIC traffic between two GPUs (Tree parent + one child on first GPU, second child on second GPU) -#define NCCL_TOPO_PATTERN_SPLIT_TREE 2 // Spread NIC traffic between two GPUs (Tree parent on first GPU, tree children on the second GPU) -#define NCCL_TOPO_PATTERN_TREE 3 // All NIC traffic going to/from the same GPU -#define NCCL_TOPO_PATTERN_RING 4 // Ring -#define NCCL_TOPO_PATTERN_NVLS 5 // NVLS+SHARP and NVLS+Tree -// Latencies in us, Bandwidths in GB/s - -#define NCCL_NUM_FUNCTIONS 5 // Send/Recv not included for now -typedef enum { ncclFuncBroadcast, ncclFuncReduce, ncclFuncAllGather, ncclFuncReduceScatter, ncclFuncAllReduce, ncclFuncSendRecv, ncclFuncSend, ncclFuncRecv, ncclNumFuncs} ncclFunc_t; - -// LL128 max BW per channel -static const double llMaxBws[3][3] = { - /* Volta-N1/Intel-N2/Intel-N4) */ {39.0, 39.0, 20.4}, - /* Ampere-N1/AMD-N2/AMD-N4) */ {87.7, 22.5 /*avg of ring & tree*/, 19.0}, - /* Hopper-N1/AMD-N2/AMD-N4) */ {87.7, 22.5 /*avg of ring & tree*/, 19.0} -}; -// Latencies in us, Bandwidths in GB/s -// Tree { LL, LL128, Simple } , Ring { LL, LL128, Simple } -static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { - { 6.8, 14.0, 0 }, { 6.6, 14.0, 8.4 }, // Tree, Ring - { 6.8, 14.0, 0 }, { 6.8, 14.0, 0 }, // Collnet Direct, Chain - { 0, 0, 23.0 }, { 0, 0, 23.0 }}; // NVLS, NVLS Tree - -static const double perChMaxRingLL128Bws[3][3] = { - /* Volta (N1/N2/N4) */ {20.0, 20.0, 20.0}, - /* Ampere (N1/N2/N4) */ {20.0, 20.0, 20.0}, - /* Hopper (N1/N2/N4) */ {36.7, 36.7, 36.7}, -}; -static const double perChMaxTreeLL128Bws[3][3] = { - /* Volta (N1/N2/N4) */ {20.0, 20.0, 20.0}, - /* Ampere (N1/N2/N4) */ {20.0, 20.0, 20.0}, - /* Hopper (N1/N2/N4) */ {36.7, 36.7, 29.0}, -}; -static const double perChMaxTreeBws[3][3] = { - /* Volta (N1/N2/N4) */ {26.5, 18.5, 10.0}, - /* Ampere (N1/N2/N4) */ {24.0, 23.6, 17.8}, - /* Hopper (N1/N2/N4) */ {38.7, 41.4, 36.0}, -}; - -// NVLink, PCI, Network 硬件时延 -#define NCCL_HW_NVLINK 0 -#define NCCL_HW_PCI 1 -#define NCCL_HW_NET 2 -static float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = - { /* NVLINK */ - { /* Tree (LL/LL128/Simple)*/ { .6, 1.25, 4 }, /* Ring (LL/LL128/Simple)*/ { .6, 1.9, 3.4 }, - /* CollNetDirect (Simple)*/ { 0, 0, 8.0 }, /* CollNetChain (Simple)*/ { 0, 0, 4.75 }, - /* NVLS */ { 0, 0, 0 }, /* NVLSTree */ { 0, 0, 0 } }, - /* PCI */ - { /* Tree (LL/LL128/Simple)*/ { 1.0, 1.9, 6 }, /* Ring (LL/LL128/Simple)*/ { 1.0, 2.5, 5.7 }, - /* CollNetDirect (Simple)*/ { 0, 0, 8.0 }, /* CollNetChain (Simple)*/ { 0, 0, 8.0 }, - /* NVLS */ { 0, 0, 0 }, /* NVLSTree */ { 0, 0, 0 } }, - /* NET */ - { /* Tree (LL/LL128/Simple)*/ { 5.0, 8.5, 14 }, /* Ring (LL/LL128/Simple)*/ { 2.7, 4.0, 14.0 }, - /* CollNetDirect (Simple)*/ { 0, 0, 10.7 }, /* CollNetChain (Simple)*/ { 0, 0, 14 }, - /* NVLS */ { 0, 0, 18 }, /* NVLSTree */ { 0, 0, 19 } } -}; - -// We want link types and path types to match as much as possible -#define LINK_LOC 0 -#define LINK_NVL 1 -// Skipping 2 for PATH_NVB -#define LINK_PCI 3 -// Skipping 4 for PATH_PXB -// Skipping 5 for PATH_PXN -// Skipping 6 for PATH_PHB -#define LINK_SYS 7 -#define LINK_NET 8 -#define PCI_BW 12.0 // PCI Gen3 x16 - -// Local (myself) -#define PATH_LOC 0 - -// Connection traversing NVLink -#define PATH_NVL 1 - -// Connection through NVLink using an intermediate GPU -#define PATH_NVB 2 - -// Connection traversing at most a single PCIe bridge -#define PATH_PIX 3 - -// Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge) -#define PATH_PXB 4 - -// Connection between a GPU and a NIC using an intermediate GPU. Used to enable rail-local, aggregated network send/recv operations. -#define PATH_PXN 5 - -// Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU) -#define PATH_PHB 6 - -// Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI) -#define PATH_SYS 7 - -// Connection through the network -#define PATH_NET 8 - -// Disconnected -#define PATH_DIS 9 - -#define DIVUP(x, y) \ - (((x)+(y)-1)/(y)) -template -constexpr Z alignUp(X x, int a) { - return (x+a-1) & Z(-a); -} - -enum ncclWorkType : uint8_t { - ncclWorkTypeUnused=0, - ncclWorkTypeColl=1, - ncclWorkTypeP2p=2, - ncclWorkTypeRegColl=3 -}; - -struct ncclWorkHeader { - union { - int32_t workNext; // when isLast=0: Offset from kernel argument workHead - uint32_t doneAcks; // when isLast=1: Monotonic (mod 1<<32) ack value to send back. - }; - uint16_t funcIndex; - uint8_t isLast:1; // last work for this kernel - uint8_t inFifo:1; // is this work in the fifo - enum ncclWorkType type; -}; - -struct ncclWorkElem { - union { - uint8_t flagBits; - struct { - uint8_t isUsed:1, redOpArgIsPtr:1, regUsed:1; - }; - }; - uint8_t nWarps; - uint8_t direct; - - const void * sendbuff; - void * recvbuff; - - size_t count; - size_t lastChunkSize; - uint32_t root; - uint8_t bid; - uint8_t nChannels; - uint64_t redOpArg; -}; - -#define NCCL_MAX_WORK_ELEMENTS ((NCCL_WORK_SIZE - alignUp(sizeof(ncclWorkHeader), alignof(ncclWorkElem)))/sizeof(ncclWorkElem)) -// Trees are not perfectly sticking to the model for medium sizes. Applying a static correction -// factor is not ideal but works quite well. Powers of two, 64 B to 256MB. -static float treeCorrectionFactor[NCCL_NUM_PROTOCOLS][23] = { - { 1.0, 1.0, 1.0, 1.0, .9, .8, .7, .7, .7, .7, .6, .5, .4, .4, .5, .6, .7, .8, .9, 1.0, 1.0, 1.0, 1.0 }, - { 1.0, 1.0, 1.0, 1.0, 1.0, .9, .8, .8, .8, .7, .6, .6, .6, .6, .6, .6, .8, .9, .9, .9, .9, 1.0, 1.0 }, - { .9, .9, .9, .9, .9, .9, .9, .8, .7, .6, .6, .5, .5, .5, .5, .6, .7, .8, .7, .7, .8, .9, .9 } -}; -#define MAXCHANNELS 10 -#define NCCL_TOPO_MAX_NODES 10 -#define NCCL_MAX_TREE_ARITY 2 - struct ncclTopoGraph { - // Input / output - int id; // ring : 0, tree : 1, collnet : 2 - int pattern; - int crossNic; - int collNet; - int minChannels; - int maxChannels; - // Output - int nChannels; //搜索到的Chanel数 - float bwIntra; //节点内单个chnnel带宽 - float bwInter; //节点间单个channel带宽 - float latencyInter; - int typeIntra; //节点内的channel路径类型 - int typeInter; //节点间的channel路径类型 - int sameChannels; //channel是否一样 - int nHops; - int intra[MAXCHANNELS*NCCL_TOPO_MAX_NODES]; //节点内每个channel路径 - int inter[MAXCHANNELS*2]; // 节点间每个channel路径 - }; - -#endif // ASTRA_SIM_MOCKNCCL_MOCKNCCL_H diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.cc b/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.cc deleted file mode 100644 index b85e91038..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.cc +++ /dev/null @@ -1,70 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ - -#include "MockNcclChannel.h" -#include - -namespace MockNccl { - MockNcclComm::MockNcclComm(int _rank,GroupType _type,MockNcclGroup* _GlobalGroup) :rank(_rank),type(_type),GlobalGroup(_GlobalGroup){ - this->ringchannels = this->GlobalGroup->genringchannels(rank,type); - this->treechannels = this->GlobalGroup->gettreechannels(rank,type); - this->nvlschannels = this->GlobalGroup->get_nvls_channels(rank,type); - // this->nvlstreechannels = this->GlobalGroup->get_nvls_tree_channels(rank,type); - } - - MockNcclComm::~MockNcclComm(){}; - - std::map>> MockNcclComm::get_rings() { - std::map>> result; - for(auto it = ringchannels.begin(); it !=ringchannels.end(); it++) { - auto ring = it->second; - auto ring_id = it->first; - for(auto rank_it = ring.begin();rank_it != ring.end(); rank_it++) { - result[rank_it->first][ring_id]= rank_it->second; - } - } - return result; - } - - MockNccl::TreeChannels MockNcclComm::get_treechannels(){ - TreeChannels nvlschannel ; - nvlschannel[0][0]=ncclTree(-1,0,8,{}); - nvlschannel[0][1]=ncclTree(-1,1,8,{}); - nvlschannel[0][2]=ncclTree(-1,2,8,{}); - nvlschannel[0][3]=ncclTree(-1,3,8,{}); - nvlschannel[0][4]=ncclTree(-1,4,8,{}); - nvlschannel[0][5]=ncclTree(-1,5,8,{}); - nvlschannel[0][6]=ncclTree(-1,6,8,{}); - nvlschannel[0][7]=ncclTree(-1,7,8,{}); - nvlschannel[0][8]=ncclTree(-1,8,-1,{0,1,2,3,4,5,6,7}); - return nvlschannel; - } - - MockNccl::TreeChannels MockNcclComm::get_nvls_channels(){ - return this->nvlschannels; - } - - MockNccl::NVLStreechannels MockNcclComm::get_nvls_tree_channels(){ - return this->nvlstreechannels; - } - - std::shared_ptr MockNcclComm::get_flow_model(uint64_t data_size,AstraSim::ComType collective_type,int layer_num,State loopstate) { - return this->GlobalGroup->getFlowModels(type,rank,collective_type,data_size,layer_num,loopstate); - } - - struct ncclInfo* MockNcclComm::get_algo_proto_info(uint64_t data_size,AstraSim::ComType collective_type){ - return this->GlobalGroup->get_algo_proto_info(type,rank,collective_type,data_size); - } -} \ No newline at end of file diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.h b/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.h deleted file mode 100644 index eca573340..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclChannel.h +++ /dev/null @@ -1,122 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#ifndef __MOCKNCCLCHANNEL_HH__ -#define __MOCKNCCLCHANNEL_HH__ - -#include -#include -#include -#include "astra-sim/system/Common.hh" -#include "MockNcclGroup.h" - -namespace MockNccl { - struct SingleFlow{ - int flow_id; - int src; - int dest; - uint64_t flow_size; - std::vectorprev; - std::vector parent_flow_id; - std::vector child_flow_id; - int channel_id; - int chunk_id; - int chunk_count; - std::string conn_type; - SingleFlow(){}; - SingleFlow( - int _flow_id, - int _src, - int _dest, - uint64_t _flow_size, - std::vector_prev, - std::vector _parent_flow_id, - std::vector _child_flow_id, - int _channel_id, - int _chunk_id, - int _chunk_count, - std::string _conn_type) - : flow_id(_flow_id), - src(_src), - dest(_dest), - flow_size(_flow_size), - prev(_prev), - parent_flow_id(_parent_flow_id), - child_flow_id(_child_flow_id), - channel_id(_channel_id), - chunk_id(_chunk_id), - chunk_count(_chunk_count), - conn_type(_conn_type) {} - ~SingleFlow(){}; - }; - - enum class State{ - Forward_Pass, - Weight_Gradient, - Input_Gradient, - }; - - enum class ComType { - None, - Reduce_Scatter, - All_Gather, - All_Reduce, - All_to_All, - All_Reduce_All_to_All - }; - - struct ncclTree { - int depth; - int rank; - int up; - std::vector down; - ncclTree(){}; - ncclTree(int _depth, int _rank, int _up, std::vector _down) - : depth(_depth), rank(_rank), up(_up), down(_down) {}; - ~ncclTree(){}; - }; - - struct ncclChannelNode{ - int depth; - int rank; - ncclChannelNode* up; - std::vector down; - ncclChannelNode(){}; - ncclChannelNode(int _depth,int _rank,ncclChannelNode* _up,std::vector_down):depth(_depth),rank(_rank),up(_up),down(_down){}; - ~ncclChannelNode(){}; - }; - - class MockNcclComm{ - public: - MockNcclComm(int _rank,GroupType _type,MockNcclGroup* _GlobalGroup); - ~MockNcclComm(); - - MockNccl::MockNcclGroup* GlobalGroup; - GroupType type; - int rank; - std::map>> ringchannels; - TreeChannels treechannels; - TreeChannels nvlschannels; - NVLStreechannels nvlstreechannels; - - std::map>> get_rings(); - MockNccl::TreeChannels get_treechannels(); - MockNccl::TreeChannels get_nvls_channels(); - MockNccl::NVLStreechannels get_nvls_tree_channels(); - std::shared_ptr get_flow_model(uint64_t data_size,AstraSim::ComType collective_type,int layer_num,State loopstate); - struct ncclInfo* get_algo_proto_info(uint64_t data_size,AstraSim::ComType collective_type); - }; -} - -#endif \ No newline at end of file diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.cc b/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.cc deleted file mode 100644 index 4daf68d53..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.cc +++ /dev/null @@ -1,2103 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#include "MockNcclGroup.h" -#include "MockNcclChannel.h" -#include -#include -#include -#include -#include -#include -#include "astra-sim/system/MockNcclLog.h" -using namespace std; -namespace MockNccl { - MockNcclGroup::MockNcclGroup(int _ngpus,int _gpus_per_nodes,int _TP_size,int _DP_size,int _PP_size,int _EP_size,int _DP_EP_size,std::vector_NVSwitch,GPUType _gpu_type):g_flow_id(0),gpu_type(_gpu_type){ - /*init groups - */ - MockNcclLog *NcclLog = MockNcclLog::getInstance(); - if (_ngpus % _gpus_per_nodes != 0 || _ngpus / _gpus_per_nodes <= 0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"The number of GPUs used is not a multiple of the number of GPUs per node."); - return; - } - int all_group_idx = 0; - int nNodes = _ngpus/_gpus_per_nodes; - int nlocalranks = _gpus_per_nodes; - int TP_nums = _ngpus/_TP_size; - int DP_nums = _ngpus/_DP_size; - int PP_nums = _ngpus/_PP_size; - int EP_nums = _ngpus/_EP_size; - int DP_EP_nums = _ngpus/_DP_EP_size; - if (TP_nums <= 0 || DP_nums <= 0 || PP_nums <= 0 || EP_nums <= 0 || DP_EP_nums <= 0 || (_TP_size * _DP_size * _PP_size != _ngpus) || (_EP_size * _DP_EP_size != _DP_size)){ - NcclLog->writeLog(NcclLogLevel::ERROR,"The group division method is incorrect."); - return; - } - int nNodesPerTPGroup = _TP_size / nlocalranks + (_TP_size % nlocalranks > 0 ? 1 : 0); - std::vectorranks; - std::vectorNVSwitchs; - // init TP group - if(_TP_size>1){ - std::setTPnodes; - for(int i =0;i1){ - std::setDPnodes; - for(int i =0;i 1){ - - } - // init EP - std::map AllTPGroups; - for(auto it = AllGroups.begin();it!=AllGroups.end();it++){ - if(it->second.type==TP){ - AllTPGroups[it->second.group_index]=it->second; - } - } - if(_EP_size>1){ - int TP_idx=0; - std::set EPnodes; - for (int i = 0; i < TP_nums / _EP_size; i++){ - TP_idx = i*_EP_size; - for(int j =0;j<_EP_size;j++){ - for(int k = 0;k 1){ - int TP_idx = 0; - std::set DP_EP_nodes; - for (int i = 0; i < TP_nums / _DP_EP_size; i++){ - TP_idx = i; - for (int j = 0; j < _DP_EP_size; j++){ - for (int k = 0; k < AllTPGroups[TP_idx].Ranks.size(); k++){ - ranks.clear(); - DP_EP_nodes.clear(); - for (int l = TP_idx; l < TP_idx + _DP_EP_size * _EP_size; l += _EP_size){ - int tmp_rank = AllTPGroups[l].Ranks[k]; - int node_idx = tmp_rank / _gpus_per_nodes; - ranks.push_back(tmp_rank); - GroupIndex[std::make_pair(tmp_rank, DP_EP)] = all_group_idx; - DP_EP_nodes.insert(node_idx); - } - NVSwitchs.clear(); - for (int idx : DP_EP_nodes){ - NVSwitchs.push_back(_NVSwitch[idx]); - GroupIndex[std::make_pair(_NVSwitch[idx], DP_EP)] = all_group_idx; - } - AllGroups[all_group_idx] = GroupInfo(all_group_idx, DP_EP, DP_EP_nodes.size(), _DP_EP_size, ranks, NVSwitchs); - all_group_idx++; - } - } - } - } - return; - } - - void MockNcclGroup::generateringchannels(std::map> localrings, MockNccl::GroupInfo* groupInfo, std::map>>& ringchannels) { - std::map>::iterator ring_it; - int current; - int prev; - int next; - int end_rank; - int nNodes = groupInfo->nNodes; - int nlocalRanks = groupInfo->nRanks/nNodes; - int delta = nNodes > 1 ? groupInfo->Ranks[nlocalRanks]-groupInfo->Ranks[0] : 0; - for(ring_it = localrings.begin();ring_it != localrings.end();ring_it++) { - prev = -1; - next = -1; - for(int i = 0; i < nNodes; i++) { - int node_send; - int node_recv; - node_recv = ring_it->second[0] + i * delta; - node_send = ring_it->second[nlocalRanks-1] + i * delta; - for(int j = 0; j < nlocalRanks; j++) { - current = ring_it->second[j] + i * delta; - if (j == nlocalRanks-1) { - next = ring_it->second[0] + (i + 1) * delta; - } else { - next = ring_it->second[j+1] + i * delta; - } - ringchannels[ring_it->first][current] = {prev,next,node_recv,node_send}; - prev = current; - } - } - end_rank = ring_it->second[nlocalRanks-1] + (nNodes - 1) * delta; - ringchannels[ring_it->first][ring_it->second[0]][0] = end_rank; - ringchannels[ring_it->first][end_rank][1] = ring_it->second[0]; - - } - } - - std::map> MockNcclGroup::gen_local_ring(int rank, GroupType type){ - GroupInfo gp_info; - int gp_idx; - std::vectorranks; - std::vectorlocalranks; - std::map>localrings; - int nNodes; - int nlocalranks; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type)) == 0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no relevant group info, resulting in an error in gen_local_ring"); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - ranks = gp_info.Ranks; - nNodes = gp_info.nNodes; - nlocalranks = ranks.size()/nNodes; - std::sort(ranks.begin(), ranks.end()); - for (int i = 0; i < nlocalranks; i++){ - localranks.push_back(ranks[i]); - } - for(int i =0;i vec; - for (int j = 0; j < nlocalranks; ++j) { - vec.push_back(localranks[(i + j) % nlocalranks]); - } - localrings[i] = vec; - } - return localrings; - } - - RingChannels MockNcclGroup::genringchannels(int rank, MockNccl::GroupType type) { - std::map>>ringchannels; - std::map>localrings; - std::map>::iterator ring_it; - GroupInfo gp_info; - int gp_idx; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - - int current; - int prev; - int next; - int end_rank; - int nNodes; - int nlocalRanks; - int delta; - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"No corresponding group information is generated, and there is an error in creating the ring channel."); - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[GroupIndex[std::make_pair(rank,type)]]; - nNodes = gp_info.nNodes; - nlocalRanks = gp_info.nRanks/nNodes; - localrings = gen_local_ring(rank,type); - - delta = nNodes > 1 ? gp_info.Ranks[nlocalRanks]-gp_info.Ranks[0] : 0; - for(ring_it = localrings.begin();ring_it != localrings.end();ring_it++) { - prev = -1; - next = -1; - for(int i = 0; i < nNodes; i++) { - int node_send; - int node_recv; - node_recv = ring_it->second[0] + i * delta; - node_send = ring_it->second[nlocalRanks-1] + i * delta; - for(int j = 0; j < nlocalRanks; j++) { - current = ring_it->second[j] + i * delta; - if (j == nlocalRanks-1) { - next = ring_it->second[0] + (i + 1) * delta; - } else { - next = ring_it->second[j+1] + i * delta; - } - ringchannels[ring_it->first][current] = {prev,next,node_recv,node_send}; - prev = current; - } - } - end_rank = ring_it->second[nlocalRanks-1] + (nNodes - 1) * delta; - ringchannels[ring_it->first][ring_it->second[0]][0] = end_rank; - ringchannels[ring_it->first][end_rank][1] = ring_it->second[0]; - } - Allringchannels[gp_idx]=ringchannels; - return ringchannels; - } - - std::shared_ptr MockNcclGroup::getFlowModels(GroupType type , int rank, AstraSim::ComType op,uint64_t data_size,int layer_num,State loopstate){ - std::string flow_model_name; - GroupInfo gp_info; - int gp_idx; - int end_rank; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in generating the flow model."); - return nullptr; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - switch (type){ - case TP: - flow_model_name = "TP"; - break; - case DP: - flow_model_name = "DP"; - break; - case EP: - flow_model_name = "EP"; - break; - case DP_EP: - flow_model_name = "DP_EP"; - break; - default: - break; - } - flow_model_name = flow_model_name + "_" + std::to_string(gp_idx) + "_" + std::to_string(layer_num) + "_" + std::to_string(static_cast(loopstate)) + "_" + std::to_string(static_cast(op)) + "_" + std::to_string(data_size); - if(flow_models.count(flow_model_name)){ - FlowName2nums[flow_model_name] ++; - std::shared_ptr presult; - if(flow_models[flow_model_name].count(rank)!=0) - presult = flow_models[flow_model_name][rank]; - else{ - presult = nullptr; - } - return presult; - } else { - flow_models[flow_model_name] = genFlowModels(type,rank,op,data_size); - FlowName2nums[flow_model_name]= 1; - return flow_models[flow_model_name][rank]; - } - } - - std::map> MockNcclGroup::genFlowModels(GroupType type , int rank, AstraSim::ComType op,uint64_t data_size){ - switch (op) { - case AstraSim::ComType::All_Reduce: - return genAllReduceFlowModels(type,rank,data_size); - case AstraSim::ComType::All_Gather: - return genAllGatherFlowModels(type,rank,data_size); - case AstraSim::ComType::Reduce_Scatter: - return genReduceScatterFlowModels(type,rank,data_size); - case AstraSim::ComType::All_to_All: - return genAlltoAllFlowModels(type,rank,data_size); - default: - break; - } - return {}; - } - - std::map> MockNcclGroup::genAlltoAllFlowModels(GroupType type, int rank, uint64_t data_size){ - FlowModels result = {}; - std::maprank2flowmodels; - std::map>rank2pflowmodels; - SingleFlow tmp_result; - uint64_t chunksize; - uint64_t send_size; - int nranks; - int chunkcount; - int chunkid; - GroupInfo gp_info; - int gp_idx; - RingChannels ringchannels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in generating the flow model."); - return {}; - } else { - gp_idx = GroupIndex[std::make_pair(rank,type)]; - ringchannels = Allringchannels[gp_idx]; - gp_info = AllGroups[gp_idx]; - } - nranks = gp_info.nRanks; - chunkcount = nranks - 1; - chunksize = data_size / nranks; - data_size = data_size / nranks; - for (int i = 0; i < gp_info.Ranks.size(); i++) { - std::vector prev; - for(int j = 0;jsecond.src; - int dst = flow_models_it->second.dest; - rank2flowmodels[src][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - rank2flowmodels[dst][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - } - for(auto it = rank2flowmodels.begin();it!=rank2flowmodels.end();it++){ - rank2pflowmodels[it->first] = std::make_shared(it->second); - } - return rank2pflowmodels; - } - - std::map> MockNcclGroup::genReduceScatterFlowModels( - GroupType type, - int rank, - uint64_t data_size) { - FlowModels result = {}; - std::maprank2flowmodels; - std::map>rank2pflowmodels; - std::map task_list = {}; - std::map task_list2 = {}; - SingleFlow tmp_result; - uint64_t chunksize; - uint64_t send_size; - int nranks; - int chunkcount; - int chunkid; - GroupInfo gp_info; - int gp_idx; - RingChannels ringchannels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in generating the flow model."); - return {}; - } else { - gp_idx = GroupIndex[std::make_pair(rank,type)]; - ringchannels = Allringchannels[gp_idx]; - gp_info = AllGroups[gp_idx]; - } - bool PXN_ENABLE = false; - const char* PXN_ENV = std::getenv("AS_PXN_ENABLE"); - if (PXN_ENV && strcmp(PXN_ENV, "1") == 0) { - PXN_ENABLE = true; - } else { - PXN_ENABLE = false; - } - nranks = gp_info.nRanks; - chunkcount = nranks - 1; - chunksize = data_size / nranks / ringchannels.size(); - data_size = data_size / nranks / ringchannels.size(); - for (auto it = ringchannels.begin(); it != ringchannels.end(); it++) { - auto ring = it->second; - auto ring_id = it->first; - task_list = {}; - send_size = 0; - chunkid = 0; - while (send_size < data_size) { - uint64_t real_chunksize = std::min(chunksize, data_size - send_size); - int prenoderecvrank = ring.rbegin()->second[2]; - int prenodesendrank = ring.rbegin()->second[3]; - int curnoderecvrank = ring.begin()->second[2]; - int curnodesendrank = ring.begin()->second[3]; - std::vector prevranks = {}; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - int cur_rank = rank_it->first; - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if (rank_it->second[0] != -1) - prevranks = {rank_it->second[0]}; - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {}, - {g_flow_id + 1}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - if (rank_it->first != -1) { - prevranks = {rank_it->first}; - } else { - prevranks = {}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunkid, - chunkcount, - "PXN_INIT"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(prenoderecvrank!=-1){ - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } - } - chunkid++; - for (int i = 0; i < nranks - 2; i++) { - task_list2 = {}; - prenoderecvrank = ring.rbegin()->second[2]; - prenodesendrank = ring.rbegin()->second[3]; - curnoderecvrank = ring.begin()->second[2]; - curnodesendrank = ring.begin()->second[3]; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - int cur_rank = rank_it->first; - int partner_flow_id = task_list[rank_it->second[0]].flow_id; - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if (rank_it->second[0] != -1) { - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {partner_flow_id}, - {g_flow_id + 1}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - if(rank_it->first!=-1){ - prevranks={rank_it->first}; - }else{ - prevranks ={}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if (prenoderecvrank != -1) { - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id .push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks= {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } - } - task_list = task_list2; - chunkid++; - } - send_size += real_chunksize; - } - } - for(auto flow_models_it = result.begin();flow_models_it!=result.end();flow_models_it++){ - int src = flow_models_it->second.src; - int dst = flow_models_it->second.dest; - rank2flowmodels[src][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - rank2flowmodels[dst][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - } - for(auto it = rank2flowmodels.begin();it!=rank2flowmodels.end();it++){ - rank2pflowmodels[it->first] = std::make_shared(it->second); - } - return rank2pflowmodels; - } - - std::map> MockNcclGroup::genAllReduceFlowModels(GroupType type , int rank,uint64_t data_size){ - ncclInfo* ncc_info = get_algo_proto_info(type,rank,AstraSim::ComType::All_Reduce,data_size); - switch (ncc_info->algorithm) { - case NCCL_ALGO_TREE: - case NCCL_ALGO_RING: - return genAllReduceRingFlowModels(type, rank, data_size); - case NCCL_ALGO_NVLS: - return genAllreduceNVLSFlowModels(type,rank,data_size); - case NCCL_ALGO_NVLS_TREE: - return {}; - default: - break; - } - } - - std::shared_ptr MockNcclGroup::genallReduceNVLSTreeFlowModels( - GroupType type, - int rank, - uint64_t data_size) { - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - GroupInfo gp_info; - int gp_idx; - int chunk_count = 1; - int chunk_size; - NVLStreechannels nvlstreechannels; - NVLStreechannels::iterator nvlstree; - FlowModels result = {}; - if(GroupIndex.count(std::make_pair(rank,type)) == 0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no relevant group info, resulting in an error in generating genallReduceNVLSTreeFlowModels."); - return nullptr; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - nvlstreechannels = AllNVLStreechannels[gp_idx]; - NcclLog->writeLog(NcclLogLevel::DEBUG," nvlstreechannels.size() %d",nvlstreechannels.size()); - chunk_size = data_size / nvlstreechannels.size() / chunk_count; - for (nvlstree = nvlstreechannels.begin(); - nvlstree != nvlstreechannels.end(); - nvlstree++) { - std::map>::iterator nvlstreenodes_it; - if (rank == 0) { - for (nvlstreenodes_it = nvlstree->second.begin(); - nvlstreenodes_it != nvlstree->second.end(); - nvlstreenodes_it++) { - NcclLog->writeLog(NcclLogLevel::DEBUG," rank %d nvls tree nodes ",nvlstreenodes_it->first); - int i = 0; - for (auto nvlstreenode : nvlstreenodes_it->second) { - NcclLog->writeLog(NcclLogLevel::DEBUG," node %d rank %d",i,nvlstreenode->rank); - if(nvlstreenode->up!=nullptr){ - NcclLog->writeLog(NcclLogLevel::DEBUG," up %d",nvlstreenode->up->rank); - } - NcclLog->writeLog(NcclLogLevel::DEBUG," down "); - for (auto down : nvlstreenode->down) { - NcclLog->writeLog(NcclLogLevel::DEBUG,"%d ",down->rank); - } - } - } - } - std::unordered_map upinDegree; - std::unordered_map downinDegree; - std::unordered_map> nodeprevs; - for (int ck = 0; ck < chunk_count; ck++) { - nodeprevs = {}; - std::vector ncclchannelnodes; - for (auto nvlstreenodes : nvlstree->second) { - for (auto nvlstreenode : nvlstreenodes.second) { - ncclchannelnodes.push_back(nvlstreenode); - upinDegree[nvlstreenode] = nvlstreenode->down.size(); - if (nvlstreenode->up == nullptr) - downinDegree[nvlstreenode] = 0; - else - downinDegree[nvlstreenode] = 1; - } - } - generate_flow_model_nvls_tree_allreduce_up( - ncclchannelnodes, - upinDegree, - nodeprevs, - chunk_size, - ck, - chunk_count, - nvlstree->first, - result); - generate_flow_model_nvls_tree_allreduce_down( - ncclchannelnodes, - downinDegree, - nodeprevs, - chunk_size, - ck, - chunk_count, - nvlstree->first, - result); - } - } - std::shared_ptr ptr_result = - std::make_shared(result); - return ptr_result; - } - - FlowModels MockNcclGroup::generate_flow_model_nvls_tree_allreduce_up( - std::vector nvlstreenodes, - std::unordered_map upinDegree, - std::unordered_map>& nodeprevs, - int chunk_size, - int chunk_id, - int chunk_count, - int channle_id, - FlowModels& result) { - std::queue q; - SingleFlow tmp_result; - for (auto entry : upinDegree) { - if (entry.second == 0) { - q.push(entry.first); - nodeprevs[entry.first] = {}; - } - } - std::string conn_tag = "NVLS_TREE"; - while (!q.empty()) { - ncclChannelNode* current = q.front(); - q.pop(); - if (current->up != nullptr) { - upinDegree[current->up]--; - std::vector _prev; - if (current->down.size() == 0) - _prev = {current->up->rank}; - else { - for (auto down : current->down) { - _prev.push_back(down->rank); - } - } - tmp_result = SingleFlow( - g_flow_id, - current->rank, - current->up->rank, - chunk_size, - _prev, - nodeprevs[current], - {}, - channle_id, - chunk_id, - chunk_count, - conn_tag); - for (int parent_flow_id : nodeprevs[current]) { - result[std::make_pair(channle_id, parent_flow_id)] - .child_flow_id.push_back(g_flow_id); - } - result[std::make_pair(channle_id, g_flow_id)] = tmp_result; - g_flow_id++; - nodeprevs[current->up].push_back(tmp_result.flow_id); - nodeprevs.erase(current); - if (upinDegree[current->up] == 0) - q.push(current->up); - } - } - return result; - } - - FlowModels MockNcclGroup::generate_flow_model_nvls_tree_allreduce_down( - std::vector nvlstreenodes, - std::unordered_map downinDegree, - std::unordered_map>& nodeprevs, - int chunk_size, - int chunk_id, - int chunk_count, - int channle_id, - FlowModels& result) { - std::queue q; - SingleFlow tmp_result; - for (auto entry : downinDegree) { - if (entry.second == 0) { - q.push(entry.first); - } - } - std::string conn_tag = "NVLS_TREE"; - while (!q.empty()) { - ncclChannelNode* current = q.front(); - q.pop(); - - if (current->down.size() > 0) { - for (ncclChannelNode* down : current->down) { - downinDegree[down]--; - std::vector _prev; - if (current->up == nullptr) { - for (ncclChannelNode* down1 : current->down) { - _prev.push_back(down1->rank); - } - } else { - _prev = {current->up->rank}; - } - tmp_result = SingleFlow( - g_flow_id, - current->rank, - down->rank, - chunk_size, - _prev, - nodeprevs[current], - {}, - channle_id, - chunk_id, - chunk_count, - conn_tag); - for (int parent_flow_id : nodeprevs[current]) { - result[std::make_pair(channle_id, parent_flow_id)] - .child_flow_id.push_back(g_flow_id); - } - result[std::make_pair(channle_id, g_flow_id)] = tmp_result; - g_flow_id++; - nodeprevs[down].push_back(tmp_result.flow_id); - if (downinDegree[down] == 0) - q.push(down); - } - } - } - return result; - } - - std::map> MockNcclGroup::genAllreduceNVLSFlowModels(GroupType type,int rank,uint64_t data_size){ - GroupInfo gp_info; - int gp_idx; - int chunk_count = 4; - std::maprank2flowmodels; - std::map>rank2pflowmodels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info , resulting in an error in genAllreduceNVLSFlowModels."); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - FlowModels result={}; - SingleFlow treeflow; - if(gp_info.nNodes == 1){ - std::vectorNVswitchs = gp_info.NVSwitchs; - std::vectorranks = gp_info.Ranks; - int chunk_size = data_size / chunk_count; - for(int ck =0;ckprevs; - std::vectorparents; - for(int k = 0;ksecond.src; - int dst = flow_models_it->second.dest; - rank2flowmodels[src][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - rank2flowmodels[dst][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - } - for(auto it = rank2flowmodels.begin();it!=rank2flowmodels.end();it++){ - rank2pflowmodels[it->first] = std::make_shared(it->second); - } - return rank2pflowmodels; - } - - std::shared_ptr MockNcclGroup::genAllReduceTreeFlowModels(GroupType type , int rank,uint64_t data_size){ - int chunk_count = 64; - int chunk_size; - SingleFlow tmp_result; - FlowModels result1 = {}; - FlowModels result = {}; - std::map task_list = {}; - std::map>::iterator tree; - GroupInfo gp_info; - int gp_idx; - TreeChannels treechannels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - - if(GroupIndex.count(std::make_pair(rank,type))==0||Alltreechannels.count(gp_idx)==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info , resulting in an error in genAllreduceNVLSFlowModels."); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - treechannels = Alltreechannels[gp_idx]; - chunk_size = data_size / treechannels.size() / chunk_count; - for(tree = treechannels.begin(); tree !=treechannels.end(); tree++) { - std::unordered_map upinDegree; - std::unordered_map downinDegree; - std::unordered_map> nodeprevs; - for(int ck = 0; ck < chunk_count; ck++){ - nodeprevs = {}; - for(auto treenode:tree->second){ - upinDegree[treenode.first] = treenode.second.down.size(); - if(treenode.second.up == -1) - downinDegree[treenode.first] = 0; - else - downinDegree[treenode.first] = 1; - } - generate_flow_model_tree_allreduce_up(tree->second,upinDegree,nodeprevs,chunk_size,ck,chunk_count,tree->first,result); - generate_flow_model_tree_allreduce_down(tree->second,downinDegree,nodeprevs,chunk_size,ck,chunk_count,tree->first,result); - } - } - std::shared_ptr ptr_result = std::make_shared(result); - return ptr_result; - } - - FlowModels MockNcclGroup::generate_flow_model_tree_allreduce_up(std::map &nodes,std::unordered_map upinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result){ - std::queue q; - std::map task_list2={}; - SingleFlow tmp_result; - for (auto entry : upinDegree) { - if (entry.second == 0) { - q.push(nodes[entry.first]); - nodeprevs[entry.first]={}; - } - } - std::string conn_tag = "TREE_INIT"; - while (!q.empty()) { - ncclTree current = q.front(); - q.pop(); - if(current.up != -1) { - upinDegree[current.up]--; - std::vector _prev; - if (current.down.size() == 0) - _prev = {current.up}; - else - _prev = current.down; - tmp_result = SingleFlow(g_flow_id,current.rank,current.up,chunk_size,_prev,nodeprevs[current.rank],{},channle_id,chunk_id,chunk_count,conn_tag); - for(int parent_flow_id:nodeprevs[current.rank]) - result[std::make_pair(channle_id,parent_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(channle_id,g_flow_id)] = tmp_result; - g_flow_id++; - nodeprevs[current.up].push_back(tmp_result.flow_id); - nodeprevs.erase(current.rank); - if(upinDegree[current.up] == 0) - q.push(nodes[current.up]); - } - } - return result; - } - - FlowModels MockNcclGroup::generate_flow_model_tree_allreduce_down(std::map &nodes,std::unordered_map downinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result){ - std::queue q; - std::map task_list2={}; - SingleFlow tmp_result; - for (auto entry : downinDegree) { - if (entry.second == 0) { - q.push(nodes[entry.first]); - } - } - std::string conn_tag = "TREE_INIT"; - while (!q.empty()) { - ncclTree current = q.front(); - q.pop(); - if(current.down.size() >0 ) { - for(int down:current.down) { - downinDegree[down] --; - std::vector _prev; - if (current.up == -1) { - _prev = current.down; - } else { - _prev = {current.up}; - } - tmp_result = SingleFlow(g_flow_id,current.rank,down,chunk_size,_prev,nodeprevs[current.rank],{},channle_id,chunk_id,chunk_count,conn_tag); - for(int parent_flow_id:nodeprevs[current.rank]) - result[std::make_pair(channle_id,parent_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(channle_id,g_flow_id)] = tmp_result; - g_flow_id++; - nodeprevs[down].push_back(tmp_result.flow_id); - if(downinDegree[down] == 0) - q.push(nodes[down]); - } - } - } - return result; - } - - std::map> MockNcclGroup::genAllReduceRingFlowModels(GroupType type , int rank,uint64_t data_size){ - FlowModels result = {}; - std::maprank2flowmodels; - std::map>rank2pflowmodels; - std::map task_list = {}; - std::map task_list2 = {}; - SingleFlow tmp_result; - uint64_t chunksize; - uint64_t send_size; - int nranks; - GroupInfo gp_info; - int gp_idx; - RingChannels ringchannels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in generating the flow model."); - return {}; - } else { - gp_idx = GroupIndex[std::make_pair(rank,type)]; - ringchannels = Allringchannels[gp_idx]; - gp_info = AllGroups[gp_idx]; - } - nranks = gp_info.nRanks; - bool PXN_ENABLE = false; - const char* PXN_ENV = std::getenv("AS_PXN_ENABLE"); - if (PXN_ENV && strcmp(PXN_ENV, "1") == 0) { - PXN_ENABLE = true; - } else { - PXN_ENABLE = false; - } - chunksize = data_size / nranks / ringchannels.size(); - data_size = data_size / nranks / ringchannels.size(); - int chunkcout = 2*(gp_info.nRanks-1); - - for(auto it = ringchannels.begin(); it !=ringchannels.end(); it++) { - auto ring = it->second; - auto ring_id = it->first; - task_list = {}; - send_size = 0; - int chunk_id = 0; - while (send_size < data_size) - { - uint64_t real_chunksize = std::min(chunksize, data_size - send_size); - int prenoderecvrank = ring.rbegin()->second[2]; - int prenodesendrank = ring.rbegin()->second[3]; - int curnoderecvrank = ring.begin()->second[2]; - int curnodesendrank = ring.begin()->second[3]; - std::vector prevranks = {}; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - int cur_rank = rank_it->first; - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks={rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {}, - {g_flow_id + 1}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - prevranks.clear(); - prevranks = {rank_it->first}; - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunk_id, - chunkcout, - "PXN_INIT"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if (prenoderecvrank != -1) { - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks={rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } - } - chunk_id++; - for(int i =0; i < nranks -1; i++) { - task_list2 = {}; - prenoderecvrank = ring.rbegin()->second[2]; - prenodesendrank = ring.rbegin()->second[3]; - curnoderecvrank = ring.begin()->second[2]; - curnodesendrank = ring.begin()->second[3]; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - int cur_rank = rank_it->first; - int partner_flow_id = task_list[rank_it->second[0]].flow_id; - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if (rank_it->second[0] != -1) { - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {partner_flow_id}, - {g_flow_id + 1}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - prevranks.clear(); - prevranks={rank_it->first}; - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunk_id, - chunkcout, - "PXN_INIT"); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(prenoderecvrank!=-1){ - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1) - { - prevranks ={rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } - } - task_list = task_list2; - chunk_id++; - } - for (int i = 0; i < nranks - 2; i++) { - task_list2 = {}; - prenoderecvrank = ring.rbegin()->second[2]; - prenodesendrank = ring.rbegin()->second[3]; - curnoderecvrank = ring.begin()->second[2]; - curnodesendrank = ring.begin()->second[3]; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - int cur_rank = rank_it->first; - int partner_flow_id = task_list[rank_it->second[0]].flow_id; - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {partner_flow_id}, - {g_flow_id + 1}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - prevranks.clear(); - if (rank_it->first != -1) { - prevranks = {rank_it->first}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunk_id, - chunkcout, - "PXN_INIT"); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(prenoderecvrank!=-1){ - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunk_id, - chunkcout, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } - } - task_list = task_list2; - chunk_id++; - } - send_size += real_chunksize; - } - } - rank2flowmodels.clear(); - for(auto flow_models_it = result.begin();flow_models_it!=result.end();flow_models_it++){ - int src = flow_models_it->second.src; - int dst = flow_models_it->second.dest; - rank2flowmodels[src][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - rank2flowmodels[dst][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - } - for(auto it = rank2flowmodels.begin();it!=rank2flowmodels.end();it++){ - rank2pflowmodels[it->first] = std::make_shared(it->second); - } - return rank2pflowmodels; - } - - std::map> MockNcclGroup::genAllGatherFlowModels(GroupType type , int rank,uint64_t data_size){ - FlowModels result = {}; - std::maprank2flowmodels; - std::map>rank2pflowmodels; - std::map task_list = {}; - std::map task_list2 = {}; - SingleFlow tmp_result; - uint64_t chunksize; - uint64_t send_size; - int nranks; - int chunkcount; - int chunkid; - GroupInfo gp_info; - int gp_idx; - RingChannels ringchannels; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in generating the flow model."); - return {}; - } else { - gp_idx = GroupIndex[std::make_pair(rank,type)]; - ringchannels = Allringchannels[gp_idx]; - gp_info = AllGroups[gp_idx]; - } - - nranks = gp_info.nRanks; - chunkcount = gp_info.nRanks-1; - chunksize = data_size / nranks / ringchannels.size(); - data_size = data_size / nranks / ringchannels.size(); - bool PXN_ENABLE = false; - const char* PXN_ENV = std::getenv("AS_PXN_ENABLE"); - if (PXN_ENV == "1") { - PXN_ENABLE = true; - } else { - PXN_ENABLE = false; - } - for(auto it = ringchannels.begin(); it !=ringchannels.end(); it++) { - auto ring = it->second; - auto ring_id = it->first; - task_list = {}; - send_size = 0; - chunkid = 0; - while (send_size < data_size) { - uint64_t real_chunksize = std::min(chunksize, data_size - send_size); - int prenoderecvrank = ring.rbegin()->second[2]; - int prenodesendrank = ring.rbegin()->second[3]; - int curnoderecvrank = ring.begin()->second[2]; - int curnodesendrank = ring.begin()->second[3]; - std::vector prevranks = {}; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - int cur_rank = rank_it->first; - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {}, - {g_flow_id + 1}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - prevranks.clear(); - if(rank_it->first!=-1){ - prevranks = {rank_it->first}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunkid, - chunkcount, - "PXN_INIT"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(prenoderecvrank!=-1){ - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if (rank_it->second[0] != -1) { - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - task_list[rank_it->first] = tmp_result; - g_flow_id++; - } - } - chunkid++; - for (int i = 0; i < nranks - 2; i++) { - task_list2 = {}; - prenoderecvrank = ring.rbegin()->second[2]; - prenodesendrank = ring.rbegin()->second[3]; - curnoderecvrank = ring.begin()->second[2]; - curnodesendrank = ring.begin()->second[3]; - for (auto rank_it = ring.begin(); rank_it != ring.end(); rank_it++) { - if (curnoderecvrank != rank_it->second[2] && - curnodesendrank != rank_it->second[3]) { - prenoderecvrank = curnoderecvrank; - prenodesendrank = curnodesendrank; - curnoderecvrank = rank_it->second[2]; - curnodesendrank = rank_it->second[3]; - } - int cur_rank = rank_it->first; - int partner_flow_id = task_list[rank_it->second[0]].flow_id; - if (rank_it->second[3] == cur_rank && - rank_it->second[2] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[2], - data_size, - prevranks, - {partner_flow_id}, - {g_flow_id + 1}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - prevranks.clear(); - if(rank_it->first){ - prevranks = {rank_it->first}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->second[2], - rank_it->second[1], - data_size, - prevranks, - {g_flow_id - 1}, - {}, - ring_id, - chunkid, - chunkcount, - "PXN"); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else if ( - rank_it->second[2] == cur_rank && - rank_it->second[3] != cur_rank && gp_info.nNodes > 1 && - PXN_ENABLE) { - prevranks.clear(); - if(prenoderecvrank!=-1){ - prevranks = {prenoderecvrank}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } else { - prevranks.clear(); - if(rank_it->second[0]!=-1){ - prevranks = {rank_it->second[0]}; - } - tmp_result = SingleFlow( - g_flow_id, - rank_it->first, - rank_it->second[1], - data_size, - prevranks, - {partner_flow_id}, - {}, - ring_id, - chunkid, - chunkcount, - "RING"); - result[std::make_pair(ring_id, partner_flow_id)].child_flow_id.push_back(g_flow_id); - task_list2[rank_it->first] = tmp_result; - result[std::make_pair(ring_id, g_flow_id)] = tmp_result; - g_flow_id++; - } - } - task_list = task_list2; - chunkid++; - } - send_size += real_chunksize; - } - } - for(auto flow_models_it = result.begin();flow_models_it!=result.end();flow_models_it++){ - int src = flow_models_it->second.src; - int dst = flow_models_it->second.dest; - rank2flowmodels[src][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - rank2flowmodels[dst][std::make_pair(flow_models_it->first.first,flow_models_it->first.second)]=flow_models_it->second; - } - for(auto it = rank2flowmodels.begin();it!=rank2flowmodels.end();it++){ - rank2pflowmodels[it->first] = std::make_shared(it->second); - } - return rank2pflowmodels; - } - - ncclChannelNode* MockNcclGroup::gen_nvls_tree_intra_channels(std::vector intra_topo,std::map> &nvlstreechannel){ - ncclChannelNode* root = new ncclChannelNode(-1,intra_topo[0],nullptr,{}); - nvlstreechannel[root->rank].push_back(root); - ncclChannelNode* nvswitch = new ncclChannelNode(-1,intra_topo[1],root,{}); - nvlstreechannel[nvswitch->rank].push_back(nvswitch); - root->down.push_back(nvswitch); - for(int i =2;idown.push_back(leaf); - nvlstreechannel[leaf->rank].push_back(leaf); - } - return root; - } - - TreeChannels MockNcclGroup::get_nvls_channels(int rank,GroupType type){ - GroupInfo gp_info; - int gp_idx; - TreeChannels nvlschannel; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in get_nvls_channels."); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - if (gp_info.nNodes > 1) { - NcclLog->writeLog(NcclLogLevel::DEBUG," %d","error NVLS ALGO dont"); - return {}; - } else { - std::vector ranks = gp_info.Ranks; - int NVswitch = gp_info.NVSwitchs[0]; - for (int i = 0; i < ranks.size(); i++) { - nvlschannel[0][ranks[i]] = ncclTree(-1, ranks[i], NVswitch, {}); - } - nvlschannel[0][ranks.size()] = ncclTree(-1, NVswitch, -1, ranks); - } - AllNVLSchannels[gp_idx] = nvlschannel; - return nvlschannel; - } - - NVLStreechannels MockNcclGroup::get_nvls_tree_channels(int rank,GroupType type){ - std::map>> nvlstreechannels; - std::map>localrings; - std::map>::iterator ring_it; - GroupInfo gp_info; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - int current; - int nNodes; - int nlocalRanks; - int delta; - int gp_idx; - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info , resulting in an error in get_nvls_tree_channels."); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - if(AllNVLStreechannels.count(gp_idx)){ - return AllNVLStreechannels[gp_idx]; - } - std::vectorroots; - roots = genInterDouBinTree(gp_info); - - nNodes = gp_info.nNodes; - nlocalRanks = gp_info.nRanks/nNodes; - localrings = gen_local_ring(rank,type); - delta = nNodes > 1 ? gp_info.Ranks[nlocalRanks]-gp_info.Ranks[0] : 0; - std::map>rings; - for(ring_it = localrings.begin();ring_it != localrings.end();ring_it++) { - for(int i = 0; i < nNodes; i++) { - for(int j = 0; j < nlocalRanks; j++) { - current = ring_it->second[j] + i * delta; - rings[ring_it->first].push_back(current); - } - } - } - std::map>> - allnode2ranks; - for (ring_it = rings.begin(); ring_it != rings.end(); ring_it++) { - int nrankspernode = gp_info.nRanks / nNodes; - for (int i = 0; i < gp_info.nNodes; i++) { - for (int j = 0; j < nrankspernode; j++) { - allnode2ranks[ring_it->first][i].push_back( - ring_it->second[i * nrankspernode + j]); - } - } - } - - std::map>>::iterator allnode2ranks_it; - int channel_id = 0; - std::map> node2ranks = allnode2ranks[0]; - for (DoubleBinaryTreeNode* root : roots) { - for (int index = 0; index < nlocalRanks; index++) { - std::map> nvlstreechannel; - std::map nodencclchannlenodes; - for (int i = 0; i < nNodes; i++) { - std::vector noderanks = node2ranks[i]; - std::vector intra_topo; - intra_topo.push_back(noderanks[index]); - intra_topo.push_back(gp_info.NVSwitchs[i]); - intra_topo.insert( - intra_topo.end(), noderanks.begin(), noderanks.end()); - NcclLog->writeLog(NcclLogLevel::DEBUG," node %d intra_topo",i); - for(auto num:intra_topo){ - NcclLog->writeLog(NcclLogLevel::DEBUG," %d",num); - } - ncclChannelNode* root = - gen_nvls_tree_intra_channels(intra_topo, nvlstreechannel); - nodencclchannlenodes[i] = root; - } - - std::map>::iterator nvlstreenodes_it; - if (rank == 0) { - for (nvlstreenodes_it = nvlstreechannel.begin(); - nvlstreenodes_it != nvlstreechannel.end(); - nvlstreenodes_it++) { - NcclLog->writeLog(NcclLogLevel::DEBUG," rank %d nvls tree nodes ",nvlstreenodes_it->first); - int i = 0; - for (auto nvlstreenode : nvlstreenodes_it->second) { - NcclLog->writeLog(NcclLogLevel::DEBUG," node %d rank %d",i,nvlstreenode->rank); - if(nvlstreenode->up!=nullptr) - NcclLog->writeLog(NcclLogLevel::DEBUG," up %d",nvlstreenode->up->rank); - NcclLog->writeLog(NcclLogLevel::DEBUG," down "); - for (auto down : nvlstreenode->down) { - NcclLog->writeLog(NcclLogLevel::DEBUG," %d ",down->rank); - } - } - } - } - - gen_nvls_tree_inter_channels( - root, nodencclchannlenodes, nvlstreechannel); - - nvlstreechannels[channel_id] = nvlstreechannel; - channel_id++; - } - } - AllNVLStreechannels[gp_idx] = nvlstreechannels; - return nvlstreechannels; - } - - ncclChannelNode* MockNcclGroup::gen_nvls_tree_inter_channels( - DoubleBinaryTreeNode* root, - std::map nodencclchannlenodes, - std::map>& nvlstreechannel) { - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if (root == nullptr) - return nullptr; - else { - NcclLog->writeLog(NcclLogLevel::DEBUG,"before root->right: %d",root->right); - NcclLog->writeLog(NcclLogLevel::DEBUG,"before root->left: %d",root->left); - if (root->left != nullptr) { - NcclLog->writeLog(NcclLogLevel::DEBUG,"after root->left: %d",root->left); - ncclChannelNode* cur = nodencclchannlenodes[root->node]; - ncclChannelNode* left = nodencclchannlenodes[root->left->node]; - cur->down.push_back(left); - left->up = cur; - gen_nvls_tree_inter_channels(root->left,nodencclchannlenodes,nvlstreechannel); - } - if (root->right != nullptr) { - NcclLog->writeLog(NcclLogLevel::DEBUG,"after root->right: %d",root->right); - ncclChannelNode* cur = nodencclchannlenodes[root->node]; - ncclChannelNode* right = nodencclchannlenodes[root->right->node]; - cur->down.push_back(right); - right->up = cur; - gen_nvls_tree_inter_channels(root->right,nodencclchannlenodes,nvlstreechannel); - } - } - } - - TreeChannels MockNcclGroup::gettreechannels(int rank, GroupType type){ - TreeChannels treechannels; - std::map>localrings; - std::map>::iterator ring_it; - GroupInfo gp_info; - int gp_idx; - int current; - int nNodes; - int nlocalRanks; - int delta; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info and group ring channel, resulting in an error in gettreechannels."); - return {}; - } - gp_idx = GroupIndex[std::make_pair(rank,type)]; - gp_info = AllGroups[gp_idx]; - if(Alltreechannels.count(gp_idx)){ - return Alltreechannels[gp_idx]; - } - - nNodes = gp_info.nNodes; - nlocalRanks = gp_info.nRanks/nNodes; - localrings = gen_local_ring(rank,type); - delta = nNodes > 1 ? gp_info.Ranks[nlocalRanks]-gp_info.Ranks[0] : 0; - std::map>rings; - for(ring_it = localrings.begin();ring_it != localrings.end();ring_it++) { - for(int i = 0; i < nNodes; i++) { - for(int j = 0; j < nlocalRanks; j++) { - current = ring_it->second[j] + i * delta; - rings[ring_it->first].push_back(current); - } - } - } - std::vector roots; - roots = genInterDouBinTree(gp_info); - std::map>> - allnode2ranks; - for (ring_it = rings.begin(); ring_it != rings.end(); ring_it++) { - int nrankspernode = gp_info.nRanks / nNodes; - for (int i = 0; i < gp_info.nNodes; i++) { - for (int j = 0; j < nrankspernode; j++) { - allnode2ranks[ring_it->first][i].push_back( - ring_it->second[i * nrankspernode + j]); - } - } - } - std::map>>::iterator allnode2ranks_it; - int channel_id = 0; - for (allnode2ranks_it = allnode2ranks.begin(); - allnode2ranks_it != allnode2ranks.end(); - allnode2ranks_it++) { - std::map> node2ranks = allnode2ranks_it->second; - for (DoubleBinaryTreeNode* root : roots) { - std::map treechannel; - for (int rank : gp_info.Ranks) { - ncclTree cur = ncclTree(-1, rank, -1, {}); - treechannel[rank] = cur; - } - ConnInterIntraTree(root, node2ranks, treechannel); - treechannels[channel_id] = treechannel; - channel_id++; - } - Alltreechannels[gp_idx] = treechannels; - } - return treechannels; - } - - void MockNcclGroup::ConnInterIntraTree(DoubleBinaryTreeNode*root,std::map>node2ranks,std::map&treechannel) { - if(root == nullptr) return; - std::vectorranks = node2ranks[root->node]; - for(int i=0;idown.push_back(ranks[i+1]); - down->up=ranks[i]; - } - - if(root->left!=nullptr){ - ncclTree *current = &treechannel[ranks[0]]; - int downrank = node2ranks[root->left->node][0]; - ncclTree *down = &treechannel[downrank]; - current->down.push_back(downrank); - down->up = ranks[0]; - ConnInterIntraTree(root->left,node2ranks,treechannel); - } - if(root->right!=nullptr){ - ncclTree *current = &treechannel[ranks[0]]; - int downrank = node2ranks[root->right->node][0]; - ncclTree *down = &treechannel[downrank]; - current->down.push_back(downrank); - down->up = ranks[0]; - ConnInterIntraTree(root->right,node2ranks,treechannel); - } - } - - std::vector MockNcclGroup::genInterDouBinTree(GroupInfo gp_info){ - vector q; - vector tmp_q; - vector result; - int nNodes = gp_info.nNodes; - std::vector nodes; - for(int i = 0;i < nNodes; i++) - nodes.push_back(i); - for(int i = 0;i < nodes.size();i++){ - q.push_back(new DoubleBinaryTreeNode(nodes[i])); - } - while (q.size() > 1){ - tmp_q = {}; - int i = 0; - for(i = 0;(i + 2) < q.size();i +=4){ - DoubleBinaryTreeNode* node0 = q[i]; - DoubleBinaryTreeNode* node1 = q[i+1]; - DoubleBinaryTreeNode* node2 = q[i+2]; - node1->left = node0; - node1->right = node2; - tmp_q.push_back(node1); - if(i+3 < q.size()) { - DoubleBinaryTreeNode* node3 = q[i+3]; - tmp_q.push_back((node3)); - } - } - if(q.size() - i == 1) { - DoubleBinaryTreeNode* node0 = q[i]; - tmp_q.push_back(node0); - } else if(q.size() - i == 2){ - DoubleBinaryTreeNode* node0 = q[i]; - DoubleBinaryTreeNode* node1 = q[i+1]; - node1->left = node0; - tmp_q.push_back(node1); - } - q = tmp_q; - } - DoubleBinaryTreeNode* root1 = InterDouBinTreeShift(q[0],nodes); - int chunk_count = 1; - for(int i =0;inodes){ - std::mapnode2treenode; - std::maprank2index; - std::queueq; - for(int i =0 ;inode; - int nodeshift = nodes[(rank2index[node] + 1) % nodes.size()]; - DoubleBinaryTreeNode* currentshift = node2treenode[nodeshift]; - if(current->left != nullptr) { - int leftnode = current->left->node; - int leftnodeshift = nodes[(rank2index[leftnode] + 1) % nodes.size()]; - currentshift->left = node2treenode[leftnodeshift]; - q.push(current->left); - } - if(current->right != nullptr) { - int rightnode = current->right->node; - int rightnodeshift = nodes[(rank2index[rightnode] + 1) % nodes.size()]; - currentshift->right = node2treenode[rightnodeshift]; - q.push(current->right); - } - } - return node2treenode[nodes[(rank2index[root->node] + 1) % nodes.size()]]; - } - - ncclInfo* MockNcclGroup::get_algo_proto_info( - GroupType type, - int rank, - AstraSim::ComType op, - uint64_t data_size) { - std::string ncclInfoName ; - GroupInfo gp_info; - MockNcclLog* NcclLog = MockNcclLog::getInstance(); - if(GroupIndex.count(std::make_pair(rank,type))==0){ - NcclLog->writeLog(NcclLogLevel::ERROR,"There is no corresponding group info, resulting in an error with get_algo_proto_info."); - return nullptr; - } - gp_info = AllGroups[GroupIndex[std::make_pair(rank,type)]]; - switch (type) - { - case TP: - ncclInfoName = "TP"; - break; - case DP: - ncclInfoName = "DP"; - break; - case EP: - ncclInfoName = "EP"; - break; - case DP_EP: - ncclInfoName = "DP_EP"; - break; - default: - break; - } - ncclInfoName+= "_"+std::to_string(static_cast(op))+"_"+std::to_string(data_size); - if(nccl_infos.count(ncclInfoName)){ - return nccl_infos[ncclInfoName]; - }else{ - bool NVLSenable = false; - const char* NVLSEnv = std::getenv("AS_NVLS_ENABLE"); - if (NVLSEnv && strcmp(NVLSEnv, "1")==0) { - NVLSenable = true; - } else { - NVLSenable = false; - } - struct ncclInfo* info = new ncclInfo(); - info->nBytes = data_size; - info->nChannels = 0; - info->coll = static_cast(op); - switch (op) { - case AstraSim::ComType::All_Reduce: - if(type==TP){ - if(gpu_type==GPUType::A100||gpu_type==GPUType::A800){ - info->algorithm = NCCL_ALGO_RING; - }else if(gpu_type==GPUType::H100||gpu_type==GPUType::H800){ - if (gp_info.nRanks >= 8 && NVLSenable) { - info->algorithm = NCCL_ALGO_NVLS; - } else { - info->algorithm = NCCL_ALGO_RING; - } - } else{ - info->algorithm = NCCL_ALGO_RING; - } - } else { - info->algorithm = NCCL_ALGO_RING; - } - break; - case AstraSim::ComType::All_Gather: - case AstraSim::ComType::Reduce_Scatter: - case AstraSim::ComType::All_to_All: - default: - info->algorithm = NCCL_ALGO_RING; - break; - } - info->protocol = NCCL_PROTO_UNDEF; - nccl_infos[ncclInfoName] = info; - return info; - } - } -} \ No newline at end of file diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.h b/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.h deleted file mode 100644 index 54327f815..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclGroup.h +++ /dev/null @@ -1,180 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#ifndef __MOCKNCCLGROUP_H__ -#define __MOCKNCCLGROUP_H__ - -#include -#include -#include -#include -#include -#include -#include -#include "astra-sim/system/Common.hh" -#include"astra-sim/system/MockNccl.h" -using namespace std; - -namespace MockNccl { - enum class State; - struct SingleFlow; - struct SingleFlow; - enum class ComType; - struct ncclTree; - struct TuneInfo; - typedef struct TuneInfo* TuneInfo_t; - struct ncclChannelNode; - typedef std::map,SingleFlow> FlowModels; - typedef std::map>> RingChannels; - typedef std::map>> NVLStreechannels; - typedef std::map> TreeChannels; - enum GroupType { - TP, - DP, - PP, - EP, - DP_EP, - NONE - }; - struct ncclInfo { - ncclFunc_t coll; - TuneInfo_t tuneinfo; - int algorithm; - int protocol; - int nChannels; - int nThreads; - size_t nBytes; - ncclInfo(){}; - ~ncclInfo(){}; - }; - struct TuneInfo{ - int nNodes; - int nRanks; - int nChannels; - int collNetSupport; - int nvlsSupport; - int minCompCap; - int maxCompCap; - std::vector graphs; - std::vector>> latencies; - std::vector>> bandwidths; - TuneInfo(){}; - ~TuneInfo(){}; - TuneInfo( - int _nNodes, - int _nRanks, - int _nChannels, - int _collNetSupport, - int _nvlsSupport, - int _minCompCap, - int _maxCompCap) - : nNodes(_nNodes), - nRanks(_nRanks), - nChannels(_nChannels), - collNetSupport(_collNetSupport), - nvlsSupport(_nvlsSupport), - minCompCap(_minCompCap), - maxCompCap(_maxCompCap) { - graphs = std::vector(NCCL_NUM_ALGORITHMS, nullptr); - latencies = std::vector>>( - NCCL_NUM_FUNCTIONS, - std::vector>( - NCCL_NUM_ALGORITHMS, std::vector(NCCL_NUM_PROTOCOLS, 0))); - bandwidths = std::vector>>( - NCCL_NUM_FUNCTIONS, - std::vector>( - NCCL_NUM_ALGORITHMS, std::vector(NCCL_NUM_PROTOCOLS, 0))); - } - }; - class GroupInfo { - public: - int group_index; - GroupType type; - int nNodes; - int nRanks; - std::vector Ranks; - std::vector NVSwitchs; - GroupInfo(){} - GroupInfo(int _group_index, GroupType _type, int _nNodes, int _nRanks, std::vector _Ranks,std::vector_NVSwitchs) - : group_index(_group_index),type(_type), nNodes(_nNodes), nRanks(_nRanks), Ranks(_Ranks),NVSwitchs(_NVSwitchs) {} - ~GroupInfo(){} - }; - class MockNcclGroup { - struct DoubleBinaryTreeNode { - int node; - DoubleBinaryTreeNode* left; - DoubleBinaryTreeNode* right; - DoubleBinaryTreeNode(int _node) : node(_node), left(nullptr), right(nullptr) {} - }; - public: - MockNcclGroup(){} - MockNcclGroup(int _ngpus,int _gpus_per_nodes, int _TP_size,int _DP_size,int _PP_size,int _EP_size,int _DP_EP_size,std::vector_NVSwitch,GPUType _gpu_type); - ~MockNcclGroup(){}; - - std::map,int> GroupIndex; - std::map AllGroups; - - std::map Allringchannels; - std::map AllNVLStreechannels; - std::map Alltreechannels; - std::map AllNVLSchannels; - - int g_flow_id; - GPUType gpu_type; - std::map FlowName2nums; - std::map >> flow_models; - std::map nccl_infos; - std::shared_ptr getFlowModels(GroupType type , int rank, AstraSim::ComType op,uint64_t data_size,int layer_num,State loopstate); - private: - std::map> genFlowModels(GroupType type , int rank, AstraSim::ComType op,uint64_t data_size); - std::map> genReduceScatterFlowModels(GroupType type , int rank, uint64_t data_size); - std::map> genAlltoAllFlowModels(GroupType type, int rank, uint64_t data_size); - std::map> genAllReduceFlowModels(GroupType type , int rank,uint64_t data_size); - std::map> genAllReduceRingFlowModels(GroupType type , int rank,uint64_t data_size); - std::map> genAllreduceNVLSFlowModels( - GroupType type, - int rank, - uint64_t data_size); - std::shared_ptrgenallReduceNVLSTreeFlowModels(GroupType type,int rank,uint64_t data_size); - FlowModels generate_flow_model_nvls_tree_allreduce_up(std::vectornvlstreenodes,std::unordered_map upinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result); - FlowModels generate_flow_model_nvls_tree_allreduce_down(std::vectornvlstreenodes,std::unordered_map downinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result); - std::shared_ptr genAllReduceTreeFlowModels(GroupType type , int rank,uint64_t data_size); - FlowModels generate_flow_model_tree_allreduce_up(std::map &nodes,std::unordered_map upinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result); - FlowModels generate_flow_model_tree_allreduce_down(std::map &nodes,std::unordered_map downinDegree,std::unordered_map>& nodeprevs,int chunk_size,int chunk_id,int chunk_count,int channle_id,FlowModels& result); - std::map> genAllGatherFlowModels(GroupType type , int rank,uint64_t data_size); - std::vector genInterDouBinTree(GroupInfo pgroupinfo); - DoubleBinaryTreeNode* InterDouBinTreeShift(DoubleBinaryTreeNode* root,std::vectornodes); - void ConnInterIntraTree(DoubleBinaryTreeNode*root,std::map>node2ranks,std::map&TreeChannel); - public: - void generateringchannels( - std::map> localrings, - MockNccl::GroupInfo* groupInfo, - std::map>>& ringchannels); - std::map> gen_local_ring(int rank, GroupType type); - RingChannels genringchannels( - int rank, - GroupType type); - TreeChannels gettreechannels(int rank, GroupType type); - TreeChannels get_nvls_channels(int rank,GroupType type); - NVLStreechannels get_nvls_tree_channels(int rank,GroupType type); - ncclChannelNode* gen_nvls_tree_intra_channels(std::vectorintra_topo,std::map> &nvlstreechannel); - ncclChannelNode* gen_nvls_tree_inter_channels(DoubleBinaryTreeNode* root,std::map nodencclchannlenodes,std::map> &nvlstreechannel); - ncclInfo* get_algo_proto_info( - GroupType type, - int rank, - AstraSim::ComType op, - uint64_t data_size); - }; -} -#endif \ No newline at end of file diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.cc b/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.cc deleted file mode 100644 index ae3646c65..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.cc +++ /dev/null @@ -1,20 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#include "MockNcclLog.h" - -MockNcclLog* MockNcclLog::instance = nullptr; -NcclLogLevel MockNcclLog::logLevel = NcclLogLevel::INFO; -std::mutex MockNcclLog::mtx; -std::string MockNcclLog::LogName = ""; \ No newline at end of file diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.h b/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.h deleted file mode 100644 index 3f08da30e..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclLog.h +++ /dev/null @@ -1,103 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#ifndef _MOCKNCCL_MOCKNCCLLOG_H_ -#define _MOCKNCCL_MOCKNCCLLOG_H_ -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#define LOG_PATH "/etc/astra-sim/" - -enum class NcclLogLevel { DEBUG, INFO, WARNING,ERROR}; - -class MockNcclLog { - private: - static MockNcclLog* instance; - static NcclLogLevel logLevel; - static std::mutex mtx; - static std::string LogName; - std::ofstream logfile; - MockNcclLog() { - const char* logLevelEnv = std::getenv("AS_LOG_LEVEL"); - logLevel = logLevelEnv ? static_cast(std::atoi(logLevelEnv)) - : NcclLogLevel::INFO; - logfile.open(LogName, std::ios::app); - } - std::string getCurrentTime() { - auto now = std::chrono::system_clock::now(); - std::time_t now_c = std::chrono::system_clock::to_time_t(now); - struct tm timeinfo; - localtime_r(&now_c, &timeinfo); - - std::ostringstream oss; - char buffer[80]; - std::strftime(buffer, 80, "%Y-%m-%d %X", &timeinfo); - oss << buffer; - return oss.str(); - } - public: - static MockNcclLog* getInstance() { - std::lock_guard lock(mtx); - if (instance == nullptr) { - instance = new MockNcclLog(); - } - return instance; - } - static void set_log_name(std::string log_name){ - LogName = LOG_PATH + log_name; - } - void writeLog(NcclLogLevel level, const char* format,...) { - if (level >= logLevel) { - std::string levelStr; - switch (level) { - case NcclLogLevel::DEBUG: - levelStr = "DEBUG"; - break; - case NcclLogLevel::INFO: - levelStr = "INFO"; - break; - case NcclLogLevel::WARNING: - levelStr = "WARNING"; - break; - case NcclLogLevel::ERROR: - levelStr = "ERROR"; - break; - default: - levelStr = "UNKNOWN"; - } - char buffer[256]; - va_list args; - va_start(args, format); - vsnprintf(buffer, sizeof(buffer), format, args); - va_end(args); - std::thread::id this_id = std::this_thread::get_id(); - std::lock_guard lock(mtx); - logfile << "[" << getCurrentTime() << "]" - << "[" << levelStr << "] " << "["<< std::hex << this_id <<"]"<< buffer << std::endl; - } - } - ~MockNcclLog() { - logfile.close(); - } -}; -#endif diff --git a/astra-sim-alibabacloud/astra-sim/system/MockNcclQps.h b/astra-sim-alibabacloud/astra-sim/system/MockNcclQps.h deleted file mode 100644 index 73df61656..000000000 --- a/astra-sim-alibabacloud/astra-sim/system/MockNcclQps.h +++ /dev/null @@ -1,38 +0,0 @@ -/* -*Copyright (c) 2024, Alibaba Group; -*Licensed under the Apache License, Version 2.0 (the "License"); -*you may not use this file except in compliance with the License. -*You may obtain a copy of the License at - -* http://www.apache.org/licenses/LICENSE-2.0 - -*Unless required by applicable law or agreed to in writing, software -*distributed under the License is distributed on an "AS IS" BASIS, -*WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -*See the License for the specific language governing permissions and -*limitations under the License. -*/ -#ifndef __MOCKNCCLQPS_H__ -#define __MOCKNCCLQPS_H__ - -#include -#include - -namespace MockNccl { -typedef std::map>, std::queue> - PeerWatingTasks; -typedef std::map>, int> - PeerQps; - -class NcclQps { - public: - NcclQps(){}; - ~NcclQps(){}; - public: - PeerQps peer_qps; - PeerWatingTasks peer_wating_tasks; - -}; -} // namespace MockNccl - -#endif diff --git a/astra-sim-alibabacloud/astra-sim/system/PacketBundle.cc b/astra-sim-alibabacloud/astra-sim/system/PacketBundle.cc index abea0ddf0..e85f1751b 100644 --- a/astra-sim-alibabacloud/astra-sim/system/PacketBundle.cc +++ b/astra-sim-alibabacloud/astra-sim/system/PacketBundle.cc @@ -4,7 +4,7 @@ LICENSE file in the root directory of this source tree. *******************************************************************************/ #include "PacketBundle.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "PhyMultiThread.hh" namespace AstraSim { PacketBundle::PacketBundle( diff --git a/astra-sim-alibabacloud/astra-sim/system/PhyMultiThread.hh b/astra-sim-alibabacloud/astra-sim/system/PhyMultiThread.hh index 077546aa0..ab57c93b3 100644 --- a/astra-sim-alibabacloud/astra-sim/system/PhyMultiThread.hh +++ b/astra-sim-alibabacloud/astra-sim/system/PhyMultiThread.hh @@ -20,7 +20,7 @@ #include #include -#include"MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" #include"AstraNetworkAPI.hh" #include"SimAiPhyCommon.hh" #ifdef PHY_RDMA diff --git a/astra-sim-alibabacloud/astra-sim/system/SimAiFlowModelRdma.cc b/astra-sim-alibabacloud/astra-sim/system/SimAiFlowModelRdma.cc index 125f483bb..59ca4b0f6 100644 --- a/astra-sim-alibabacloud/astra-sim/system/SimAiFlowModelRdma.cc +++ b/astra-sim-alibabacloud/astra-sim/system/SimAiFlowModelRdma.cc @@ -20,10 +20,10 @@ #include"SimAiFlowModelRdma.hh" #include"PhyMultiThread.hh" -#include"MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" #include"BootStrapnet.hh" #include"AstraNetworkAPI.hh" -#include"MockNcclLog.h" +#include"SimCCL/mock/MockNcclLog.h" #define IB_PORT 1 diff --git a/astra-sim-alibabacloud/astra-sim/system/StreamBaseline.cc b/astra-sim-alibabacloud/astra-sim/system/StreamBaseline.cc index 3e79da3b8..7d1865fe5 100644 --- a/astra-sim-alibabacloud/astra-sim/system/StreamBaseline.cc +++ b/astra-sim-alibabacloud/astra-sim/system/StreamBaseline.cc @@ -4,7 +4,7 @@ LICENSE file in the root directory of this source tree. *******************************************************************************/ #include "StreamBaseline.hh" -#include "MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "astra-sim/system/collective/Algorithm.hh" namespace AstraSim { StreamBaseline::StreamBaseline( diff --git a/astra-sim-alibabacloud/astra-sim/system/Sys.cc b/astra-sim-alibabacloud/astra-sim/system/Sys.cc index 4f1ab16c1..fc94627b6 100644 --- a/astra-sim-alibabacloud/astra-sim/system/Sys.cc +++ b/astra-sim-alibabacloud/astra-sim/system/Sys.cc @@ -19,7 +19,7 @@ LICENSE file in the root directory of this source tree. #include "astra-sim/system/collective/DoubleBinaryTreeAllReduce.hh" #include "astra-sim/system/collective/HalvingDoubling.hh" #include "astra-sim/system/collective/Ring.hh" -#include "astra-sim/system/collective/NcclTreeFlowModel.hh" +#include "astra-sim/system/collective/NcclFlowModel.hh" #include "astra-sim/system/scheduling/OfflineGreedy.hh" #include "astra-sim/system/topology/BasicLogicalTopology.hh" #include "astra-sim/system/topology/DoubleBinaryTreeTopology.hh" @@ -27,7 +27,7 @@ LICENSE file in the root directory of this source tree. #include "astra-sim/system/topology/LocalRingGlobalBinaryTree.hh" #include "astra-sim/system/topology/LocalRingNodeA2AGlobalDBT.hh" #include "astra-sim/system/topology/Torus3D.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "astra-sim/workload/Layer.hh" #include @@ -697,7 +697,7 @@ std::vector Sys:: CollectiveImplementationType::NcclFlowModel)); } else if(dimension_input == "ncclRingTreeModel") { result.push_back(new CollectiveImplementation( - CollectiveImplementationType::NcclTreeFlowModel)); + CollectiveImplementationType::NcclFlowModel)); } else { sys_panic( "Cannot interpret collective implementations. Please check the collective implementations in the sys" @@ -1224,7 +1224,7 @@ CollectivePhase Sys::generate_collective_phase( CollectivePhase vn( this, queue_id, - new NcclTreeFlowModel( + new NcclFlowModel( collective_type, id, layer_num, @@ -1234,7 +1234,9 @@ CollectivePhase Sys::generate_collective_phase( injection_policy, boost_mode, RingFlowModels, - channels.size())); + channels.size(), + nccl_info->algorithm, + nccl_info->protocol)); return vn; } else if(nccl_info->algorithm == NCCL_ALGO_TREE) { std::shared_ptr TreeFlowModels; @@ -1248,7 +1250,7 @@ CollectivePhase Sys::generate_collective_phase( CollectivePhase vn( this, queue_id, - new NcclTreeFlowModel( + new NcclFlowModel( collective_type, id, layer_num, @@ -1258,7 +1260,9 @@ CollectivePhase Sys::generate_collective_phase( injection_policy, boost_mode, TreeFlowModels, - treechannels.size())); + treechannels.size(), + nccl_info->algorithm, + nccl_info->protocol)); return vn; } else if(nccl_info->algorithm == NCCL_ALGO_NVLS) { @@ -1298,7 +1302,7 @@ CollectivePhase Sys::generate_collective_phase( CollectivePhase vn( this, queue_id, - new NcclTreeFlowModel( + new NcclFlowModel( collective_type, id, layer_num, @@ -1308,9 +1312,37 @@ CollectivePhase Sys::generate_collective_phase( injection_policy, boost_mode, RingFlowModels, - treechannels.size())); + treechannels.size(), + nccl_info->algorithm, + nccl_info->protocol)); + return vn; + } else if(nccl_info->algorithm == NCCL_ALGO_PAT) { + // PAT uses Ring-like topology (genPATFlowModels delegates to Ring) + std::shared_ptr PATFlowModels = std::static_pointer_cast(ptr_FlowModels); + std::map>> pat_channels; + { + Sys::sysCriticalSection cs; + pat_channels = mock_nccl_comms[comm_ps]->get_rings(); + cs.ExitSection(); + } + CollectivePhase vn( + this, + queue_id, + new NcclFlowModel( + collective_type, + id, + layer_num, + (RingTopology*)topology, + data_size, + direction, + injection_policy, + boost_mode, + PATFlowModels, + pat_channels.size(), + nccl_info->algorithm, + nccl_info->protocol)); return vn; - } + } } else { std::cerr diff --git a/astra-sim-alibabacloud/astra-sim/system/Sys.hh b/astra-sim-alibabacloud/astra-sim/system/Sys.hh index 2bbe3fda8..251091d91 100644 --- a/astra-sim-alibabacloud/astra-sim/system/Sys.hh +++ b/astra-sim-alibabacloud/astra-sim/system/Sys.hh @@ -25,14 +25,14 @@ LICENSE file in the root directory of this source tree. #include "Common.hh" #include "SendPacketEventHandlerData.hh" #include "UsageTracker.hh" -#include "astra-sim/system/MockNcclChannel.h" +#include "SimCCL/mock/MockNcclChannel.h" #include "astra-sim/system/topology/RingTopology.hh" #include "astra-sim/workload/Workload.hh" #ifdef NS3_MTP #include "ns3/mtp-interface.h" #endif #include -#include "astra-sim/system/MockNcclGroup.h" +#include "SimCCL/mock/MockNcclGroup.h" namespace AstraSim { class MemBus; diff --git a/astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.cc b/astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.cc similarity index 91% rename from astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.cc rename to astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.cc index 4c84a873a..39d41178d 100644 --- a/astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.cc +++ b/astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.cc @@ -19,10 +19,11 @@ #endif #include -#include "NcclTreeFlowModel.hh" +#include "NcclFlowModel.hh" +#include "astra-sim/system/Sys.hh" #include "astra-sim/system/PacketBundle.hh" #include "astra-sim/system/RecvPacketEventHadndlerData.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #ifdef PHY_RDMA #include "astra-sim/system/SimAiFlowModelRdma.hh" extern FlowPhyRdma flow_rdma; @@ -30,8 +31,8 @@ extern FlowPhyRdma flow_rdma; namespace AstraSim { -std::atomic NcclTreeFlowModel::g_flow_inCriticalSection(false); -NcclTreeFlowModel::NcclTreeFlowModel( +std::atomic NcclFlowModel::g_flow_inCriticalSection(false); +NcclFlowModel::NcclFlowModel( ComType type, int id, int layer_num, @@ -41,7 +42,9 @@ NcclTreeFlowModel::NcclTreeFlowModel( InjectionPolicy injection_policy, bool boost_mode, std::shared_ptr ptr_flow_models, - int treechannels) + int treechannels, + int algorithm, + int protocol) : Algorithm(layer_num){ this->start_time = std::chrono::high_resolution_clock::now(); this->end_time = std::chrono::high_resolution_clock::now(); @@ -55,6 +58,8 @@ NcclTreeFlowModel::NcclTreeFlowModel( this->name = Name::Ring; this->enabled = true; this->m_channels = treechannels; + this->m_algorithm = algorithm; + this->m_protocol = protocol; this->judge_exit_flag.store(false); this->judge_exit_mutex.unlock(); this->judge_mutex.unlock(); @@ -86,7 +91,7 @@ NcclTreeFlowModel::NcclTreeFlowModel( if(pQps->peer_qps.count(std::make_pair(f.second.channel_id,std::make_pair(f.second.src,f.second.dest)))==0){ pQps->peer_qps[std::make_pair(f.second.channel_id,std::make_pair(f.second.src,f.second.dest))]=1; } - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; this->_stream_count[f.second.channel_id] += 1; cs.ExitSection(); assert(this->_flow_models.count(f.first) == 0); @@ -119,7 +124,7 @@ NcclTreeFlowModel::NcclTreeFlowModel( } } -void NcclTreeFlowModel::init_indegree_mapping(){ +void NcclFlowModel::init_indegree_mapping(){ MockNccl::FlowModels::iterator tree_it; for(tree_it = _flow_models.begin();tree_it != _flow_models.end();tree_it++) { if(tree_it->second.src!=id) continue; @@ -127,11 +132,11 @@ void NcclTreeFlowModel::init_indegree_mapping(){ } } -int NcclTreeFlowModel::get_non_zero_latency_packets() { +int NcclFlowModel::get_non_zero_latency_packets() { return (nodes_in_ring - 1) * parallel_reduce * 1; } -void NcclTreeFlowModel::run(EventType event, CallData* data) { +void NcclFlowModel::run(EventType event, CallData* data) { BasicEventHandlerData* ehd = (BasicEventHandlerData*)data; MockNcclLog* NcclLog = MockNcclLog::getInstance(); if (event == EventType::General) { @@ -165,7 +170,7 @@ void NcclTreeFlowModel::run(EventType event, CallData* data) { } } assert(flow_exist == true); - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; free_packets[std::make_pair(channel_id, flowTag.sender_node)]--; bool tag = true; for (int i = 0; i < m_channels; i++) { @@ -200,7 +205,7 @@ void NcclTreeFlowModel::run(EventType event, CallData* data) { } NcclLog->writeLog(NcclLogLevel::DEBUG,"next_flow_list.size %d",next_flow_list.size()); for (int next_flow_id : next_flow_list) { - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; if (indegree_mapping.count(next_flow_id) == 0) { flow_exist = false; cs.ExitSection(); @@ -265,7 +270,7 @@ void NcclTreeFlowModel::run(EventType event, CallData* data) { } #ifdef PHY_MTP waiting_to_exit(); - NcclLog->writeLog(NcclLogLevel::DEBUG, "NcclTreeFlowModel::waiting_to_exit end "); + NcclLog->writeLog(NcclLogLevel::DEBUG, "NcclFlowModel::waiting_to_exit end "); #endif } } else if(event == EventType::PacketSentFinshed){ @@ -278,7 +283,7 @@ void NcclTreeFlowModel::run(EventType event, CallData* data) { reduce(channel_id,sent_flow_id); bool flow_exist = next_flow_list.size() == 0 ? true : false; #ifndef PHY_MTP - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; pQps->peer_qps[std::make_pair(flowTag.channel_id,std::make_pair(flowTag.sender_node,flowTag.receiver_node))]=1; cs.ExitSection(); if(pQps->peer_wating_tasks[std::make_pair(flowTag.channel_id,std::make_pair(flowTag.sender_node,flowTag.receiver_node))].size()>0){ @@ -294,7 +299,7 @@ void NcclTreeFlowModel::run(EventType event, CallData* data) { } } -bool NcclTreeFlowModel::init_recv_ready() { +bool NcclFlowModel::init_recv_ready() { std::map>,std::vector> recv_ready_flows; for(auto flow : _flow_models){ if(flow.second.src!=id) continue; @@ -330,7 +335,7 @@ bool NcclTreeFlowModel::init_recv_ready() { return true; } -bool NcclTreeFlowModel::recv_ready(int channel_id, int flow_id) { +bool NcclFlowModel::recv_ready(int channel_id, int flow_id) { std::vectorrecv_prevs; auto flow_model = _flow_models[std::make_pair(channel_id,flow_id)]; recv_prevs = flow_model.prev; @@ -365,7 +370,7 @@ bool NcclTreeFlowModel::recv_ready(int channel_id, int flow_id) { return true; } -void NcclTreeFlowModel::release_packets(int channel_id, int flow_id, uint64_t message_size) { +void NcclFlowModel::release_packets(int channel_id, int flow_id, uint64_t message_size) { MockNcclLog* NcclLog = MockNcclLog::getInstance(); if (NPU_to_MA == true) { (new PacketBundle( @@ -395,23 +400,23 @@ void NcclTreeFlowModel::release_packets(int channel_id, int flow_id, uint64_t me NcclLog->writeLog(NcclLogLevel::DEBUG,"id: %d finish release_packets",id); } -void NcclTreeFlowModel::process_stream_count(int channel_id) { +void NcclFlowModel::process_stream_count(int channel_id) { MockNcclLog*NcclLog = MockNcclLog::getInstance(); #ifdef PHY_MTP send_packets--; #else - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; if (_stream_count[channel_id] > 0) { _stream_count[channel_id]--; } - NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclTreeFlowModel::process_stream_count channel_id %d _stream_count %d",channel_id,_stream_count[channel_id]); + NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclFlowModel::process_stream_count channel_id %d _stream_count %d",channel_id,_stream_count[channel_id]); if (_stream_count[channel_id] == 0 && stream->state != StreamState::Dead) stream->changeState(StreamState::Zombie); cs.ExitSection(); #endif } -void NcclTreeFlowModel::reduce(int channel_id, int flow_id) { +void NcclFlowModel::reduce(int channel_id, int flow_id) { process_stream_count(channel_id); #ifndef PHY_MTP if(!packets[std::make_pair(channel_id, flow_id)].empty()){ @@ -420,10 +425,10 @@ void NcclTreeFlowModel::reduce(int channel_id, int flow_id) { #endif } -bool NcclTreeFlowModel::iteratable(int channel_id) { +bool NcclFlowModel::iteratable(int channel_id) { MockNcclLog* NcclLog = MockNcclLog::getInstance(); bool all_channel_finished = true, all_packets_freed = true; - NcclTreeFlowModel::FlowCriticalSection cs; + NcclFlowModel::FlowCriticalSection cs; for(int i = 0; i < m_channels; ++ i) { if(_stream_count.count(i) != 0 && _stream_count[i] != 0) all_channel_finished = false; } @@ -442,7 +447,7 @@ bool NcclTreeFlowModel::iteratable(int channel_id) { return true; } -void NcclTreeFlowModel::insert_packets(int channel_id, int flow_id) { +void NcclFlowModel::insert_packets(int channel_id, int flow_id) { MockNcclLog* NcclLog = MockNcclLog::getInstance(); assert(channel_id < m_channels); if (!enabled) { @@ -510,7 +515,7 @@ void NcclTreeFlowModel::insert_packets(int channel_id, int flow_id) { Sys::sys_panic("should not inject nothing!"); } -bool NcclTreeFlowModel::ready(int channel_id, int flow_id) { +bool NcclFlowModel::ready(int channel_id, int flow_id) { MockNcclLog* NcclLog = MockNcclLog::getInstance(); MyPacket packet; { @@ -519,7 +524,7 @@ bool NcclTreeFlowModel::ready(int channel_id, int flow_id) { stream->changeState(StreamState::Executing); } if (!enabled || packets[std::make_pair(channel_id, flow_id)].size() == 0 || _stream_count[channel_id] == 0) { - NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclTreeFlowModel not ready!"); + NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclFlowModel not ready!"); return false; } packet = packets[std::make_pair(channel_id, flow_id)].front(); @@ -582,6 +587,9 @@ bool NcclTreeFlowModel::ready(int channel_id, int flow_id) { snd_req.flowTag.sender_node = id; snd_req.flowTag.receiver_node = packet.preferred_dest; snd_req.flowTag.pQps = this->pQps; + snd_req.flowTag.algorithm = m_algorithm; + snd_req.flowTag.protocol = m_protocol; + snd_req.flowTag.gpus_per_node = Sys::all_generators[id] ? Sys::all_generators[id]->ngpus_per_node : -1; if (this->comType == ComType::All_Reduce_NVLS) snd_req.flowTag.nvls_on = true; else @@ -605,7 +613,7 @@ bool NcclTreeFlowModel::ready(int channel_id, int flow_id) { return true; } -void NcclTreeFlowModel::exit() { +void NcclFlowModel::exit() { MockNcclLog* NcclLog = MockNcclLog::getInstance(); #ifdef PHY_MTP auto now = std::chrono::system_clock::now(); @@ -615,7 +623,7 @@ void NcclTreeFlowModel::exit() { .count(); NcclLog->writeLog( NcclLogLevel::DEBUG, - "NcclTreeFlowModel exit time %lld", + "NcclFlowModel exit time %lld", now_us); end_time = std::chrono::high_resolution_clock::now(); auto duration = std::chrono::duration_cast(end_time - start_time); @@ -629,12 +637,12 @@ void NcclTreeFlowModel::exit() { } #endif stream->owner->proceed_to_next_vnet_baseline((StreamBaseline*)stream); - NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclTreeFlowModel exit"); + NcclLog->writeLog(NcclLogLevel::DEBUG,"NcclFlowModel exit"); return; } #ifdef PHY_RDMA -bool NcclTreeFlowModel::phy_iteratable(int channel_id){ +bool NcclFlowModel::phy_iteratable(int channel_id){ MockNcclLog* NcclLog = MockNcclLog::getInstance(); bool all_send_finished = true, all_recv_finished = true; bool exit_flag = true; @@ -649,7 +657,7 @@ bool NcclTreeFlowModel::phy_iteratable(int channel_id){ } } -bool NcclTreeFlowModel::phy_ready(int channel_id,int flow_id) { +bool NcclFlowModel::phy_ready(int channel_id,int flow_id) { MockNcclLog* NcclLog = MockNcclLog::getInstance(); if (stream->state == StreamState::Created || stream->state == StreamState::Ready) { @@ -714,6 +722,9 @@ bool NcclTreeFlowModel::phy_ready(int channel_id,int flow_id) { snd_req.flowTag.sender_node = id; snd_req.flowTag.receiver_node = flow.dest; snd_req.flowTag.pQps = this->pQps; + snd_req.flowTag.algorithm = m_algorithm; + snd_req.flowTag.protocol = m_protocol; + snd_req.flowTag.gpus_per_node = Sys::all_generators[id] ? Sys::all_generators[id]->ngpus_per_node : -1; if (this->comType == ComType::All_Reduce_NVLS) snd_req.flowTag.nvls_on = true; else @@ -737,10 +748,10 @@ bool NcclTreeFlowModel::phy_ready(int channel_id,int flow_id) { return true; } -void NcclTreeFlowModel::waiting_to_exit() { +void NcclFlowModel::waiting_to_exit() { MockNcclLog* NcclLog = MockNcclLog::getInstance(); NcclLog->writeLog( - NcclLogLevel::DEBUG, "NcclTreeFlowModel::waiting_to_exit begin "); + NcclLogLevel::DEBUG, "NcclFlowModel::waiting_to_exit begin "); while (!judge_exit_flag) { }; exit(); diff --git a/astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.hh b/astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.hh similarity index 87% rename from astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.hh rename to astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.hh index ff49102f1..e02208c29 100644 --- a/astra-sim-alibabacloud/astra-sim/system/collective/NcclTreeFlowModel.hh +++ b/astra-sim-alibabacloud/astra-sim/system/collective/NcclFlowModel.hh @@ -13,8 +13,8 @@ *limitations under the License. */ -#ifndef __NCCL_TREE_FLOW_MODEL_HH__ -#define __NCCL_TREE_FLOW_MODEL_HH__ +#ifndef __NCCL_FLOW_MODEL_HH__ +#define __NCCL_FLOW_MODEL_HH__ #include #include @@ -35,10 +35,10 @@ #include "astra-sim/system/MemBus.hh" #include "astra-sim/system/MyPacket.hh" #include "astra-sim/system/topology/RingTopology.hh" -#include "astra-sim/system/MockNcclQps.h" +#include "SimCCL/mock/MockNcclQps.h" namespace AstraSim { -class NcclTreeFlowModel : public Algorithm { +class NcclFlowModel : public Algorithm { public: std::chrono::time_point start_time; std::chrono::time_point end_time; @@ -64,14 +64,16 @@ class NcclTreeFlowModel : public Algorithm { uint32_t m_channels; uint32_t len_channel; MockNccl::NcclQps* pQps; + int m_algorithm; // NCCL algorithm (RING/TREE/NVLS/PAT), carried to flowTag for send_lat bucketing + int m_protocol; // NCCL protocol (LL/LL128/Simple), carried to flowTag for send_lat bucketing std::condition_variable judge_exit_cv; std::mutex judge_exit_mutex; std::mutex judge_mutex; std::atomic judge_exit_flag; - NcclTreeFlowModel(){}; - ~NcclTreeFlowModel(){}; + NcclFlowModel(){}; + ~NcclFlowModel(){}; - NcclTreeFlowModel( + NcclFlowModel( ComType type, int id, int layer_num, @@ -81,7 +83,9 @@ class NcclTreeFlowModel : public Algorithm { InjectionPolicy injection_policy, bool boost_mode, std::shared_ptr ptr_flow_models, - int treechannels); + int treechannels, + int algorithm = -1, + int protocol = -1); virtual void run(EventType event, CallData* data); void process_stream_count(int channel_id); void release_packets(int channel_id, int flow_id, uint64_t message_size); diff --git a/astra-sim-alibabacloud/astra-sim/system/topology/GeneralComplexTopology.cc b/astra-sim-alibabacloud/astra-sim/system/topology/GeneralComplexTopology.cc index ef68a82d5..9e76c01ee 100644 --- a/astra-sim-alibabacloud/astra-sim/system/topology/GeneralComplexTopology.cc +++ b/astra-sim-alibabacloud/astra-sim/system/topology/GeneralComplexTopology.cc @@ -47,7 +47,7 @@ GeneralComplexTopology::GeneralComplexTopology( collective_implementation[dim]->type == CollectiveImplementationType::NcclFlowModel || collective_implementation[dim]->type == - CollectiveImplementationType::NcclTreeFlowModel) { + CollectiveImplementationType::NcclFlowModel) { RingTopology* ring = new RingTopology( RingTopology::Dimension::NA, id, diff --git a/astra-sim-alibabacloud/astra-sim/workload/Layer.cc b/astra-sim-alibabacloud/astra-sim/workload/Layer.cc index fe2d862b7..9a78ae156 100644 --- a/astra-sim-alibabacloud/astra-sim/workload/Layer.cc +++ b/astra-sim-alibabacloud/astra-sim/workload/Layer.cc @@ -6,7 +6,7 @@ LICENSE file in the root directory of this source tree. #include "Layer.hh" #include "astra-sim/system/DataSet.hh" #include "astra-sim/system/IntData.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" #include "astra-sim/system/AstraParamParse.hh" // #ifdef ANALYTI #include "astra-sim/system/calbusbw.h" @@ -417,7 +417,7 @@ LayerData Layer::report( for (auto& ml : net_message_latency) { layerData.avg_network_message_dealy.push_back(std::make_pair(i, ml / FREQ)); } - if (seprate_log) + if (seprate_log && EndToEnd) { std::string data; std::pair total_bw; @@ -638,9 +638,9 @@ LayerData Layer::report( layerData.avg_network_message_dealy.push_back(std::make_pair(i, ml / FREQ)); } #ifdef NS3_MPI - if (seprate_log) + if (seprate_log && EndToEnd) #else - if (seprate_log) + if (seprate_log && EndToEnd) #endif { std::string data; diff --git a/astra-sim-alibabacloud/astra-sim/workload/Layer.hh b/astra-sim-alibabacloud/astra-sim/workload/Layer.hh index 14de96750..52e94d063 100644 --- a/astra-sim-alibabacloud/astra-sim/workload/Layer.hh +++ b/astra-sim-alibabacloud/astra-sim/workload/Layer.hh @@ -21,7 +21,7 @@ LICENSE file in the root directory of this source tree. #include "Workload.hh" #include "astra-sim/system/StreamStat.hh" #include "astra-sim/system/Sys.hh" -#include"astra-sim/system/MockNcclGroup.h" +#include"SimCCL/mock/MockNcclGroup.h" namespace AstraSim { class DataSet; diff --git a/astra-sim-alibabacloud/astra-sim/workload/Workload.cc b/astra-sim-alibabacloud/astra-sim/workload/Workload.cc index af4db2259..3de336bea 100755 --- a/astra-sim-alibabacloud/astra-sim/workload/Workload.cc +++ b/astra-sim-alibabacloud/astra-sim/workload/Workload.cc @@ -6,7 +6,9 @@ LICENSE file in the root directory of this source tree. #include "Workload.hh" #include "CSVWriter.hh" #include "Layer.hh" -#include "astra-sim/system/MockNcclLog.h" +#include "SimCCL/mock/MockNcclLog.h" +#include +#include namespace AstraSim { Workload::~Workload() { @@ -65,13 +67,13 @@ Workload::Workload( if (generator->id == 0 && seprate_log) { std::cout << "stat path: " << path << " ,total rows: " << total_rows << " ,stat row: " << stat_row << std::endl; - detailed = new CSVWriter(path, "detailed_"+std::to_string(generator->total_nodes)+".csv"); - end_to_end = new CSVWriter(path, "EndToEnd.csv"); - dimension_utilization = - new CSVWriter(path, run_name + "_dimension_utilization_"+std::to_string(generator->npu_offset)+".csv"); - if (stat_row == 0) { - initialize_stat_files(); - } + detailed = new CSVWriter(path, "detailed_"+std::to_string(generator->total_nodes)+".csv"); + end_to_end = new CSVWriter(path, "EndToEnd.csv"); + dimension_utilization = + new CSVWriter(path, run_name + "_dimension_utilization_"+std::to_string(generator->npu_offset)+".csv"); + if (stat_row == 0) { + initialize_stat_files(); + } } #endif } @@ -192,7 +194,7 @@ void Workload::report() { dims.push_back( generator->scheduler_unit->usage[i].report_percentage(10000)); } - dimension_utilization->finalize_csv(dims); + if (dimension_utilization) dimension_utilization->finalize_csv(dims); } #endif #ifdef NS3_MPI @@ -202,7 +204,7 @@ void Workload::report() { dims.push_back( generator->scheduler_unit->usage[i].report_percentage(10000)); } - dimension_utilization->finalize_csv(dims); + if (dimension_utilization) dimension_utilization->finalize_csv(dims); } #endif } diff --git a/astra-sim-alibabacloud/build/astra_ns3/build.sh b/astra-sim-alibabacloud/build/astra_ns3/build.sh index a80138123..689eefd74 100755 --- a/astra-sim-alibabacloud/build/astra_ns3/build.sh +++ b/astra-sim-alibabacloud/build/astra_ns3/build.sh @@ -28,6 +28,7 @@ function cleanup { rm -rf "${NS3_DIR}"/simulation/build rm -rf "${NS3_DIR}"/simulation/cmake-cache rm -rf "${NS3_APPLICATION}"/astra-sim + rm -rf "${NS3_APPLICATION}"/SimCCL cd "${SCRIPT_DIR:?}" } @@ -51,6 +52,18 @@ function compile { cp "${ASTRA_SIM_DIR}"/network_frontend/ns3/*.h "${NS3_DIR}"/simulation/scratch/ rm -rf "${NS3_APPLICATION}"/astra-sim cp -r "${ASTRA_SIM_DIR}" "${NS3_APPLICATION}"/ + # Copy versioned MockNccl translation layer into the ns3 app tree (flat layout). + # SIMAI_NCCL_VERSION selects which mock version to build (default: v2.30). + MOCK_VERSION="${SIMAI_NCCL_VERSION:-v2.30}" + MOCK_SRC="${SCRIPT_DIR}/../../../SimCCL/src/mock/${MOCK_VERSION}" + if [ ! -d "$MOCK_SRC" ]; then + echo "[ERROR] SimCCL mock version '${MOCK_VERSION}' not found at: ${MOCK_SRC}" >&2 + echo "Supported versions: v2.20, v2.30" >&2 + exit 1 + fi + rm -rf "${NS3_APPLICATION}"/SimCCL + mkdir -p "${NS3_APPLICATION}"/SimCCL/mock + cp -r "$MOCK_SRC"/* "${NS3_APPLICATION}"/SimCCL/mock/ cd "${NS3_DIR}/simulation" CC='gcc' CXX='g++' ./ns3 configure -d debug --enable-mtp diff --git a/astra-sim-alibabacloud/build/simai_analytical/CMakeLists.txt b/astra-sim-alibabacloud/build/simai_analytical/CMakeLists.txt index 0e76ac090..f63767042 100644 --- a/astra-sim-alibabacloud/build/simai_analytical/CMakeLists.txt +++ b/astra-sim-alibabacloud/build/simai_analytical/CMakeLists.txt @@ -16,6 +16,21 @@ endif() # Setup project project (AstraSim_Analytical) +# SimCCL mock version selection (matches ns3 build default) +if(DEFINED ENV{SIMAI_NCCL_VERSION}) + set(MOCK_VERSION $ENV{SIMAI_NCCL_VERSION}) +else() + set(MOCK_VERSION "v2.30") +endif() + +# Global include paths for SimCCL mock resolution +include_directories("${PROJECT_SOURCE_DIR}/../../") +include_directories("${PROJECT_SOURCE_DIR}/../../../") +# Create symlink: build_dir/SimCCL/mock/ -> actual mock version dir +file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/SimCCL") +file(CREATE_LINK "${PROJECT_SOURCE_DIR}/../../../SimCCL/src/mock/${MOCK_VERSION}" "${CMAKE_BINARY_DIR}/SimCCL/mock" SYMBOLIC) +include_directories("${CMAKE_BINARY_DIR}") + # Compile AstraSim library add_subdirectory("${PROJECT_SOURCE_DIR}/../../" AstraSim) diff --git a/astra-sim-alibabacloud/build/simai_phy/CMakeLists.txt b/astra-sim-alibabacloud/build/simai_phy/CMakeLists.txt index bb64a2eb9..89d07ea3a 100644 --- a/astra-sim-alibabacloud/build/simai_phy/CMakeLists.txt +++ b/astra-sim-alibabacloud/build/simai_phy/CMakeLists.txt @@ -9,6 +9,22 @@ set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ${MPI_LINK_FLAGS} -std=c++ project (SimAi_Phy) add_definitions(-DPHY_MTP) + +# SimCCL mock version selection (matches ns3/analytical build default) +if(DEFINED ENV{SIMAI_NCCL_VERSION}) + set(MOCK_VERSION $ENV{SIMAI_NCCL_VERSION}) +else() + set(MOCK_VERSION "v2.30") +endif() + +# Global include paths for SimCCL mock resolution +include_directories("${PROJECT_SOURCE_DIR}/../../") +include_directories("${PROJECT_SOURCE_DIR}/../../../") +# Create symlink: build_dir/SimCCL/mock/ -> actual mock version dir +file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/SimCCL") +file(CREATE_LINK "${PROJECT_SOURCE_DIR}/../../../SimCCL/src/mock/${MOCK_VERSION}" "${CMAKE_BINARY_DIR}/SimCCL/mock" SYMBOLIC) +include_directories("${CMAKE_BINARY_DIR}") + add_subdirectory("${PROJECT_SOURCE_DIR}/../../" AstraSim) add_subdirectory ("${PROJECT_SOURCE_DIR}/../../astra-sim/network_frontend/phynet" simai_phynet) diff --git a/docs/CN/Tutorial.md b/docs/CN/Tutorial.md new file mode 100644 index 000000000..de9043cf7 --- /dev/null +++ b/docs/CN/Tutorial.md @@ -0,0 +1,412 @@ +> 本文为 [Tutorial.md](../Tutorial.md) 的中文翻译，以英文版本为准。 + +# 简介 + +SimAI 是一个综合性的大规模 AI 训练仿真工具包，提供三种主要仿真场景： + +1. **SimAI-Analytical** - 一种分析型仿真工具，抽象底层网络通信细节。采用简化方法，使用 busbw（总线带宽）来估算集合/点对点通信的通信时间，实现快速场景验证。主要应用场景包括（但不限于）： + + * *性能分析*：比较不同模型的完成时间（如研究 Expert 数量对 MoE 模型训练性能的影响） + + * *框架级并行参数优化*：平衡和优化 TP/EP/PP 参数，分析端到端时延效果 + + * *Scale-up 探索*：研究特定场景优化中不同 scale-up 域的并行参数性能 + + * *Scale-out 带宽选择*：研究不同 GPU 性能下的性价比带宽配置 + +> 💡 *当前开源二进制通过 `-nv`/`-nic`/`-n_p_s` 自动计算 busbw（默认方式；见下方 Analytical 仿真部分）。手动 `busbw.yaml` 路径（`-busbw`）在当前开源二进制中尚未接线——该用法请参考早期 SimAI 版本。欢迎联系我们了解更多详情。✨* + +2. **SimAI-Simulation(NS-3)** - 一种高保真全栈仿真工具，理论上可与任何纯网络模拟器集成。提供 LLM 训练过程中通信行为的细粒度还原。当前支持 NS-3 作为网络后端（我们鼓励集成新的网络仿真工具）。主要研究方向包括： + + * *集合通信算法研究*：为非交换机架构和其他新兴网络拓扑设计和优化集合通信流量模式 + + * *网络协议研究*：评估和优化不同架构下的网络协议、拥塞控制算法、路由机制等底层网络技术 + + * *新型网络架构设计*：探索创新性网络架构 + +> 💡 我们强烈鼓励研究人员基于 SimAI-Simulation 进行创新扩展和突破性研究，发表于顶级会议。加入我们的社区或通过邮件联系我们——我们致力于为有前景的研究方向提供技术支持！✨ + +3. **SimAI-Physical(TODO)** + +各组件的更多功能请参考 [SimCCL](https://github.com/aliyun/SimCCL) 和 [ns-3-alibabacloud](https://github.com/aliyun/ns-3-alibabacloud)。 + +# 🛠️ 环境搭建 + +通常情况下，运行 SimAI 需要使用 [AICB](https://github.com/aliyun/aicb?tab=readme-ov-file#generate-workloads-for-simulation-simai) 工具生成 Workload 文件。为创建精确的 Workload，您可能需要使用 AIOB 功能确定各计算 kernel 的时间，这需要 GPU 环境。因此，我们建议直接在最新的 **NGC 镜像**中运行 SimAI 全栈工具包。 + +> 💡 **重要提示**：SimAI-Simulation 编译需要移除 ninja（NGC 镜像中预装）。使用以下命令移除： +> ```bash +> apt remove ninja-build && pip uninstall ninja +> ``` + +编译指南： + +```bash +# Clone the repository +$ git clone https://github.com/aliyun/SimAI.git +$ cd ./SimAI/ + +# Clone submodules +$ git submodule update --init --recursive +# Make sure to use the newest commit +$ git submodule update --remote + +# Compile SimAI-Analytical +$ ./scripts/build.sh -c analytical + +# Compile SimAI-Simulation (ns3) +$ ./scripts/build.sh -c ns3 +``` + +# 🌐 SimAI-Analytical 使用方法 +## 📝 Workload 生成 + +使用 [AICB](https://github.com/aliyun/aicb) 中的 [SimAI-WorkloadGenerator](https://github.com/aliyun/aicb?tab=readme-ov-file#generate-workloads-for-simulation-simai) 功能生成仿真用工作负载。这将产生类似 [workload_analytical.txt](../../example/workload_analytical.txt) 的 `.txt` 文件，包含： + +- `model_parallel_NPU_group`：表示 Tensor Parallelism 大小 +- `ep`：表示 Expert 模型并行大小 +- `pp`：表示流水线模型并行大小 +- `vpp`：Virtual Pipeline Parallelism（默认：`--num-layers-per-virtual-pipeline-stage=1`，实现最小 PP bubble） + +> 💡 *更多详情请参考 [AICB Workload Tutorial](https://github.com/aliyun/aicb/blob/master/training/tutorial.md#workload)* + +## 🔧 Busbw 设置 + +SimAI-Analytical 通过直接指定 busbw 来估算集合通信时间，从而抽象底层网络细节。要为不同场景自定义通信 busbw，可使用如下格式的 [busbw.yaml](../../example/busbw.yaml) 文件： + +```yaml +test +TP: + allreduce,: 300 # TP 中 AllReduce busbw 300GB/s + allgather,: 280 + reducescatter,: 280 + alltoall,: 230 +DP: + allreduce,: null + allgather,: 380 # DP 中 AllGather busbw 380GB/s + reducescatter,: 380 + alltoall,: null +EP: + allreduce,: null + allgather,: 45 # DP_EP 中 AllGather busbw 45GB/s + reducescatter,: 45 # DP_EP 中 ReduceScatter busbw 45GB/s + alltoall,: 80 # EP 中 AlltoAll busbw 80GB/s +``` +> 🔍 *对自动 busbw 计算（考虑集群规模、架构、并行参数、小消息调整和延迟）感兴趣？欢迎联系讨论！* ✨ + +## 🖥️ Analytical 仿真 + +使用以下命令运行 analytical 仿真（自动 busbw —— 默认方式，当前二进制支持）： + +```bash +$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- +``` + +> 注意：下面的 `-busbw example/busbw.yaml`（用户自定义 busbw）用法在当前开源 analytical 二进制中**不受支持**——`-busbw` 参数不会被解析，命令会打印用法并退出。此处仅作保留参考；手动 `busbw.yaml` 用法请参考早期 SimAI 版本。 + +```bash +$ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml +``` + +### 必需参数 + +| 参数 | 完整形式 | 描述 | +|:---------:|:----------|:------------| +| `-w` | `--workload` | 指定 Workload 文件路径 | +| `-g` | `--gpus` | 指定仿真 GPU 规模 | +| `-g_p_s` | `--gpus-per-server` | 指定 Scale-up 大小 | +| `-r` | `--result` | 指定输出文件路径和前缀（默认：`./results/`）
建议包含仿真参数，如
`A100-llama405b-tp8-pp16-dp128-ga16-ep1-NVL8` | +| `-busbw` | `--bus-bandwidth` | 指定 busbw 文件路径
（建议直接修改 `example/busbw.yaml`） | + +### 可选参数 + +| 参数 | 完整形式 | 描述 | +|:---------:|:----------|:------------| +| `-v` | `--visual` | 是否生成可视化文件 | + +### 通信组重叠比例 + +以下参数指定通信组的重叠比例（默认：0，表示无重叠）： + +| 参数 | 完整形式 | 描述 | 范围 | +|:---------:|:----------|:------------|:------| +| `-dp_o` | `--dp-overlap-ratio` | DP 重叠比例 | [0.0-1.0] | +| `-ep_o` | `--ep-overlap-ratio` | EP 重叠比例 | [0.0-1.0] | +| `-tp_o` | `--tp-overlap-ratio` | TP 重叠比例 | [0.0-1.0] | +| `-pp_o` | `--pp-overlap-ratio` | PP 重叠比例 | [0.0-1.0] | + +> 📝 *由于重叠策略多样且重叠比例依赖场景，我们优先使用简单高效的方式直接指定重叠条件。* + + +## 结果分析 + +### 原始数据 + +正常运行 SimAI-Analytical 将生成如下图所示的 CSV 输出。 + +第二行包含摘要信息，包括暴露时间以及每个通信组的计算时间的绝对值和百分比，以及一次迭代的端到端时间。下面是每个具体层的操作详情。 + +simai_raw + + +### 可视化 + +如果运行 SimAI-Analytical 时指定 `-v`，将生成以下内容： + +simai_visual + +# SimAI-Simulation 使用方法 +## 📝 Workload 生成 + +使用与 SimAI-Analytical 相同的 workload，由 [AICB](https://github.com/aliyun/aicb) 中的 [SimAI-WorkloadGenerator](https://github.com/aliyun/aicb?tab=readme-ov-file#generate-workloads-for-simulation-simai) 功能生成。 + +## 🔧 TOPO 设置 +运行 SimAI-Simulator 前，需要生成 ns-3-alibabacloud 可识别的拓扑文件。 +### 拓扑模板 +为增强便利性，我们提供了5种常见架构的模板，包括 Spectrum-X、单平面 AlibabaHPN、双平面 Alibaba HPN 和 DCN+。可设置参数 `-topo`。以下五张图给出了概览。 +如需了解更多关于 dual-ToR 和 dual-plane 的信息，请阅读 [HPN 7.0](https://ennanzhai.github.io/pub/sigcomm24-hpn.pdf) 的文章。 + + + + + + + + + + + + + +
Spectrum-X

Spectrum-X（一个 Pod）

DCN+

DCN+SingleToR（一个 Pod）

DCN+

DCN+DualToR（一个 Pod）

AlibabaHPN_SinglePlane

单平面 AlibabaHPN（一个 Pod）

AlibabaHPN_DualPlane

双平面 AlibabaHPN（一个 Pod）

+ + +以下命令生成图中所示的 8 GPU Spectrum-X 拓扑： +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo Spectrum-X -g 8 -psn 1 +``` +Spectrum + +以下表格给出了各层级参数的描述，并展示了每个模板的默认参数。用户可更改 `-topo` 名称和对应的 `-g` 来生成相应结构的拓扑。注意，如果未输入 gpu 数量，则会生成每个模板一个 Pod 的拓扑。 +（当前不支持超过一个 Pod 的 GPU 规模。） + +| 参数层级 | 参数 | 描述 | +|-----------------|------------|-------------| +| 整体结构 | `-topo` | 拓扑模板 | +| | `-g` | GPU 数量 | +| | `--dp` | 启用双平面，默认单平面 | +| | `--ro` | 启用 rail-optimized | +| | `--dt` | 启用双网卡和双 ToR | +| | `-er` | 错误率 | +| 主机内部 | `-gps` | 每服务器 GPU 数 | +| | `-gt` | GPU 类型 | +| | `-nsps` | 每服务器 NV switch 数 | +| | `-nvbw` | NVLink 带宽 | +| | `-nl` | NVLink 延迟 | +| | `-l` | NIC 延迟 | +| 段内 | `-bw` | NIC 到 ASW 带宽 | +| | `-asw` | ASW 交换机数量 | +| | `-nps` | 每交换机 NIC 数（每 ASW 连接的 GPU 数） | +| Pod 内 | `-psn` | PSW 交换机数量 | +| | `-apbw` | ASW 到 PSW 带宽 | +| | `-app` | 每 PSW 的 ASW 数 | + + +| 参数层级 | 参数 | Spectrum-X | AlibabaHPN 单平面 | AlibabaHPN 双平面 | DCN+ 双平面 | DCN+ 单平面 | +|-----------------|------------|-------------|-------------------------|-----------------------|---------------|---------------| +| 整体结构 | `-topo` | `Spectrum-X`| `AlibabaHPN` | `AlibabaHPN` | `DCN+` | `DCN+` | +| | `-g` | 4096 | 15360 | 15360 | 512 | 512 | +| | `--dp` | false | false | true | false | false | +| | `--ro` | true | true | true | false | false | +| | `--dt` | false | true | true | true | false | +| | `-er` | 0 | 0 | 0 | 0 | 0 | +| 主机内部 | `-gps` | 8 | 8 | 8 | 8 | 8 | +| | `-gt` | H100 | H100 | H100 | H100 | A100 | +| | `-nsps` | 1 | 1 | 1 | 1 | 1 | +| | `-nvbw` | 2880Gbps | 2880Gbps | 2880Gbps | 2880Gbps | 2880Gbps | +| | `-nl` | 0.000025ms | 0.000025ms | 0.000025ms | 0.000025ms | 0.000025ms | +| | `-l` | 0.0005ms | 0.0005ms | 0.0005ms | 0.0005ms | 0.0005ms | +| 段内 | `-bw` | 400Gbps | 200Gbps | 200Gbps | 200Gbps | 400Gbps | +| | `-asw` | 64 | 120 | 120 | 2 | 1 | +| | `-nps` | 64 | 128 | 128 | 128 | 64 | +| Pod 内 | `-psn` | 64 | 120 | 120 | 8 | 4 | +| | `-apbw` | 400Gbps | 400Gbps | 400Gbps | 400Gbps | 400Gbps | +| | `-app` | 64 | 240 | 120 | 8 | 4 | + +可根据不同的 `-topo` 名称和参数生成各模板的拓扑： +对于 64 GPU、16 asn、16 psn 的双平面 AlibabaHPN，使用以下命令： +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo AlibabaHPN --dp -g 64 -asn 16 -psn 16 +``` +对于 128 GPU、2 asn、8 psn 的双 ToR DCN，使用以下命令： +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo DCN+ --dt -g 128 -asn 2 -psn 8 +``` +请注意 `--ro` `--dt` `--dp` 对 `Spectrum-X` 无效（固定为 Rail-Optimized Single-ToR Single Plane），`--ro` `--dt` 对 `AlibabaHPN` 无效（Rail-Optimized 双平面，可为单平面或双平面），`--ro` `--dp` 对 `DCN+` 无效（非 Rail-Optimized 单平面，可为 Single ToR 或 Dual ToR）。 + +用户可自定义拓扑。例如，如果要构建一个 32 GPU、200Gbps 带宽、A100、8 psn 的 rail-optimized single ToR 拓扑，输入以下命令： +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -g 32 -bw 200Gbps -gt A100 -psn 8 --ro +``` + +## 🖥️ SimAI-NS3 仿真 + +```bash +$ AS_SEND_LAT=3 AS_NVLS_ENABLE=1 ./bin/SimAI_simulator -t 16 -w ./example/microAllReduce.txt -n ./Spectrum-X_8g_8gps_400Gbps_H100 -c astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +| 环境变量名 | 描述 | 默认值 | +|---------------------------|----------------------------------|-------------------------------------------| +| `AS_LOG_LEVEL` | 日志级别 | `DEBUG`, `INFO`, `WARNING`, `ERROR`, `UNKNOWN`；默认为 `INFO` | +| `AS_PXN_ENABLE` | 启用 PXN | `0/1`；默认为 `false` | +| `AS_NVLS_ENABLE` | 启用 NVLS | `0/1`；默认为 `false` | +| `AS_SEND_LAT` | 设置数据包发送延迟 | 默认为 `6`，单位为 `us` | +| `AS_NVLSTREE_ENABLE` | 启用 NVLSTREE | 默认为 `false` | + +| 参数 | 描述 | 默认值 | +|----------------------------|------------------------------------------|--------------------------------------------------------------------| +| `-t --thread` | 多线程加速的线程数 | 默认为 `1`；如启用多线程，控制线程数在 `8` 到 `16` 之间。| +| `-w --workload` | workload 路径 | `./microAllReduce.txt` | +| `-n --network-topo` | 网络拓扑路径 | None | + +## RING VS NVLS +### workload +```bash +HYBRID_TRANSFORMER_FWD_IN_BCKWD model_parallel_NPU_group: 8 ep: 1 pp: 1 vpp: 8 ga: 1 all_gpus: 32 checkpoints: 0 checkpoint_initiates: 0 +6 +embedding_layer -1 556000 ALLREDUCE 16777216 1 NONE 0 1 NONE 0 1 +embedding_layer -1 556000 ALLREDUCE 33554432 1 NONE 0 1 NONE 0 1 +embedding_layer -1 556000 ALLREDUCE 67108864 1 NONE 0 1 NONE 0 1 +embedding_layer -1 556000 ALLREDUCE 134217728 1 NONE 0 1 NONE 0 1 +embedding_layer -1 556000 ALLREDUCE 268435456 1 NONE 0 1 NONE 0 1 +embedding_layer -1 556000 ALLREDUCE 536870912 1 NONE 0 1 NONE 0 1 + +``` +### NVLS 拓扑文件和运行命令 +```bash +cd SimAI +./scripts/build.sh -c ns3 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py --ro -g 32 -gt H100 -bw 400Gbps -nvbw 1360Gbps +AS_SEND_LAT=12 AS_NVLS_ENABLE=1 ./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt -n ./Rail_Opti_SingleToR_32g_8gps_400Gbps_H100 -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` +### RING 拓扑文件和运行命令 +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py --ro -g 32 -gt H100 -bw 400Gbps -nvbw 1440Gbps +AS_SEND_LAT=2 AS_PXN_ENABLE=1 ./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt -n ./Rail_Opti_SingleToR_32g_8gps_400Gbps_H100 -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` +### 结果 +| 消息大小 | NVLS | RING | +|----------|--------|--------| +| 16M | 148.88 | 141.84 | +| 32M | 178.04 | 153.68 | +| 64M | 197.38 | 160.60 | +| 128M | 208.70 | 163.85 | +| 256M | 214.87 | 165.72 | +| 512M | 218.09 | 166.68 | + + +## Spectrum-X 架构 VS DCN+ 架构 +### workload +```bash +HYBRID_TRANSFORMER_FWD_IN_BCKWD model_parallel_NPU_group: 8 ep: 1 pp: 1 vpp: 8 ga: 1 all_gpus: 256 checkpoints: 0 checkpoint_initiates: 0 +1 +embedding_layer -1 556000 NONE 0 1 NONE 0 1 ALLREDUCE 536870912 1 +``` +### 网络拓扑文件 +```bash +# DCN+ 拓扑文件（Single ToR，非 rail-optimized） +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo DCN+ -g 256 -psn 64 -bw 400Gbps +# Spectrum 拓扑文件（Single ToR，Rail-optimized） +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py -topo Spectrum-X -g 256 +``` +### 运行命令 +```bash +# DCN+ 运行命令 +AS_SEND_LAT=2 ./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt -n ./DCN+SingleToR_256g_8gps_400Gbps_H100 -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# HPN7.0 运行命令 +AS_SEND_LAT=2 ./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt -n ./Spectrum-X_256g_8gps_400Gbps_H100 -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` +| 消息大小 | Spectrum-X | DCN-SingleToR | +|----------|-------------|--------------| +| 16M | 33.095585 | 23.332048 | +| 32M | 38.572166 | 25.762846 | +| 64M | 42.049648 | 23.677116 | +| 128M | 44.036110 | 35.209461 | +| 256M | 45.101425 | 36.205692 | +| 512M | 45.653648 | 36.242008 | + +# SimAI-Physical 使用方法 +当前仿真器兼容 ns3 离散事件模拟器作为网络后端，同时支持物理网络后端用于物理包注入。 + +## 编译 +SimAI-Phy 当前使用 roceV2 协议进行流量生成。编译过程需要依赖与 RDMA 物理设备相关的 libverbs，以及 MPI 程序。编译前，请确认您的环境能成功运行基础的 RDMA perftest 流量生成工具，并支持相关 MPI 程序。 +```bash +# Clone the repository +$ git clone https://github.com/aliyun/SimAI.git +$ cd ./SimAI/ + +# Clone submodules +$ git submodule update --init --recursive +# Make sure to use the newest commit +$ git submodule update --remote + +# Compile SimAI-Analytical +$ ./scripts/build.sh -c analytical + +# Compile SimAI-Simulation (ns3) +$ ./scripts/build.sh -c ns3 + +# Compile SimAI-phynet (phynet) +$ sudo yum install openmpi openmpi-devel +$ export MPI_INCLUDE_PATH=/usr/include/openmpi-x86_64/ +$ export MPI_BIN_PATH=/usr/lib64/openmpi/bin/mpic++ +$ ./scripts/build.sh -c phy +``` +## Workload 生成 +SimAI-Phy 物理流量生成所需的 workload 与 Sim-Simulation 相同，通过 AICB 生成。 + +### 示例 workload +```bash +HYBRID_TRANSFORMER_FWD_IN_BCKWD model_parallel_NPU_group: 2 ep: 1 pp: 1 vpp: 8 ga: 1 all_gpus: 2 checkpoints: 0 checkpoint_initiates: 0 +10 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +mlp_norm -1 1055000 ALLGATHER 1073741824 1055000 NONE 0 1055000 NONE 0 100 +``` +## 准备主机列表 +这里主要任务是准备启动 MPI 程序所需的 iplist，与 nccl-test 不同。这里的 IP 数量应与实际参与物理流量生成的网卡数量一致，而非参与物理流量生成的节点数量。 +```bash +33.255.199.130 +33.255.199.129 +``` +## 运行 +### MPI 运行 +```bash +/usr/lib64/openmpi/bin/mpirun -np 2 -host 33.255.199.130,33.255.199.129 --allow-run-as-root -x AS_LOG_LEVEL=0 ./bin/SimAI_phynet ./hostlist -g 2 -w ./example/microAllReduce.txt +``` +以下输出表明程序已运行完成。 + +Sim-phynet_finished + +## MPI 程序参数设置 + +| 参数 | 描述 | 默认值 | +|------------------|-------------------------------------------------------|---------------| +| -np | 进程数 | NULL | +| -host | IP 列表 | NULL | +| --allow-run-as-root | 允许 MPI 程序以 root 权限运行 | FALSE | +## SimAI-phynet 参数设置 + +| 参数 | 描述 | 默认值 | +|------------------|------------------------------|----------------------------------------------------------| +| hostlist | 主机 IP 列表 | NULL | +| -w --workload | workload 路径 | ./microAllReduce.txt | +| -i --gid_index | 网络拓扑路径 | 0 | +| -g --gpus | GPU 数量 | 8（应与主机 IP 列表中的 IP 数量一致） | diff --git a/docs/CN/configuration/build-options.md b/docs/CN/configuration/build-options.md new file mode 100644 index 000000000..16e6fa7c4 --- /dev/null +++ b/docs/CN/configuration/build-options.md @@ -0,0 +1,77 @@ +# SimAI 构建选项 + +> [English Version](../../configuration/build-options.md) + +## 构建模式 + +| 模式 | 命令 | 输出二进制 | 说明 | +|---|---|---|---| +| analytical | `./scripts/build.sh -c analytical` | `bin/SimAI_analytical` | 快速 busbw 仿真（无网络建模） | +| ns3 | `./scripts/build.sh -c ns3` | `bin/SimAI_simulator` | 完整 ns3 网络仿真（高保真度） | +| phy | `./scripts/build.sh -c phy` | `bin/SimAI_phynet` | 物理 RDMA 流量生成（需要 IB 硬件） | + +## 清理构建 + +```bash +./scripts/build.sh -l ns3 # 清理 ns3 构建产物 +./scripts/build.sh -l analytical # 清理 analytical 构建 +./scripts/build.sh -l phy # 清理 physical 构建 +``` + +## SimAI_simulator 参数（ns3 模式） + +| 参数 | 长格式 | 说明 | 默认值 | +|---|---|---|---| +| `-t` | `--thread` | 多线程加速的线程数 | 1 | +| `-w` | `--workload` | Workload 文件路径 | 必填 | +| `-n` | `--network-topo` | 网络拓扑文件路径 | 必填 | +| `-c` | `--config` | 仿真配置文件路径 | 必填 | + +典型用法： +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +> **线程建议**：多线程模式推荐使用 8-16 个线程。更高的线程数可减少大规模拓扑的仿真时间。 + +## SimAI_analytical 参数 + +| 参数 | 长格式 | 说明 | 默认值 | +|---|---|---|---| +| `-w` | `--workload` | Workload 文件路径 | 必填 | +| `-g` | `--gpus` | GPU 总数 | 必填 | +| `-g_p_s` | `--gpus-per-server` | 每台服务器 GPU 数（纵向扩展大小） | 必填 | +| `-r` | `--result` | 输出文件路径/前缀 | `./results/` | +| `-busbw` | `--bus-bandwidth` | busbw.yaml 文件路径（用户自定义 busbw）。**当前开源二进制未接线** — 见下方说明 | 不适用 | +| `-v` | `--visual` | 生成可视化文件 | 关闭 | +| `-nv` | - | NVLink 带宽（GB/s），用于自动 busbw（默认方法） | 推荐 | +| `-nic` | - | NIC 带宽（GB/s），用于自动 busbw（默认方法） | 推荐 | +| `-n_p_s` | - | 每台服务器 NIC 数，用于自动 busbw（默认方法） | 推荐 | + +典型用法（自动 busbw —— 默认方式，当前二进制支持）： +```bash +./bin/SimAI_analytical \ + -w ./example/workload_analytical.txt \ + -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 \ + -r example- +``` + +> 注意：下面的 `-busbw example/busbw.yaml`（用户自定义 busbw）用法在当前开源 analytical 二进制中**不受支持**——`-busbw` 参数不会被解析，命令会打印用法并退出。此处仅作保留参考；手动 `busbw.yaml` 用法请参考早期 SimAI 版本。 +```bash +./bin/SimAI_analytical \ + -w example/workload_analytical.txt \ + -g 9216 -g_p_s 8 \ + -r test- \ + -busbw example/busbw.yaml +``` + +## 拓扑生成器参数 + +详见 [快速入门 - 生成拓扑](../getting_started/quickstart.md#2-生成拓扑)。 + +--- + +> 最后编辑：2026-06-25 diff --git a/docs/CN/configuration/env-variables.md b/docs/CN/configuration/env-variables.md new file mode 100644 index 000000000..a56620250 --- /dev/null +++ b/docs/CN/configuration/env-variables.md @@ -0,0 +1,65 @@ +# SimAI 环境变量参考 + +> [English Version](../../configuration/env-variables.md) + +## 仿真环境变量 + +以下变量控制 `SimAI_simulator`（ns3 模式）的行为。 + +| 变量 | 默认值 | 单位 | 作用域 | 说明 | +|---|---|---|---|---| +| `AS_SEND_LAT` | 未设置（使用表查询） | **纳秒 (ns)** | entry.h（ns3 前端） | 覆盖每个 flow 的发送延迟。最高优先级——覆盖所有 send_lat_table 查询。 | +| `AS_NVLS_ENABLE` | `0` | - | MockNcclGroup.cc | 启用 NVLS 算法用于 AllReduce（H20/H100/H800） | +| `AS_NVLSTREE_ENABLE` | `0` | - | MockNcclGroup.cc | 启用 NVLS Tree 算法 | +| `AS_PXN_ENABLE` | `0` | - | MockNcclGroup.cc | 启用 PXN 跨节点代理 | +| `AS_LOG_LEVEL` | `INFO` | - | 系统全局 | 日志级别：`DEBUG`、`INFO`、`WARNING`、`ERROR` | + +### AS_SEND_LAT 详细说明 + +**单位：纳秒 (ns)**。send_lat_table 默认值范围为 6000 到 22000 ns，取决于算法、协议和链路类型。 + +| 示例值 | 含义 | +|---|---| +| 未设置 | 使用 send_lat_table[algo][proto]（Ring+LL+NVLINK=7200ns，PAT+Simple+NET=22000ns 等） | +| `AS_SEND_LAT=6` | 6 ns — 实际上等于禁用发送延迟（用于快速测试） | +| `AS_SEND_LAT=7200` | 7200 ns = 7.2 μs — 匹配 Ring+LL+NVLINK 表值 | +| `AS_SEND_LAT=22000` | 22000 ns = 22 μs — 匹配 PAT+Simple+NET 表值 | + +> **警告**：部分旧文档（Tutorial.md）将 AS_SEND_LAT 描述为“单位 μs，默认 6”。这是不准确的。实际单位为**纳秒**，已通过代码分析和实验验证确认。 + +#### 快速命令 + +```bash +# 前置 +cd SimAI/ && ./scripts/build.sh -c ns3 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +# 基线（使用 send_lat_table） +./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# 覆盖发送延迟（7200ns = Ring+LL+NVLINK 表值） +AS_SEND_LAT=7200 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +完整 A/B 实验指南请见 [快速入门](../getting_started/quickstart.md#5-验证-as_send_lat-效果ab-实验)。 + +## SimCCL 集成变量 + +SimCCL（集合通信建模层）有其自身的环境变量。详见： +- [SimCCL 环境变量](../../../SimCCL/docs/CN/configuration/env-variables.md) — Standalone 变量 +- [SimCCL 集成指南](../../../SimCCL/docs/CN/integration/integration-with-simai.md) — 跨模块变量 + +## 与 send_lat_table 的关系 + +`entry.h` 中的 `send_lat_table` 为每个（算法、协议、链路类型）组合提供发送延迟值。当未设置 `AS_SEND_LAT` 时，每个 flow 使用与其算法和协议匹配的表值。设置后，所有 flow 统一使用覆盖值。 + +表机制的详细分析请参见 [send-lat-analysis.md](send-lat-analysis.md)。 + +--- + +> 最后编辑：2026-06-25 diff --git a/docs/CN/configuration/send-lat-analysis.md b/docs/CN/configuration/send-lat-analysis.md new file mode 100644 index 000000000..01bd9b27c --- /dev/null +++ b/docs/CN/configuration/send-lat-analysis.md @@ -0,0 +1,178 @@ +# send_lat 深入分析：send_latency_table vs AS_SEND_LAT + +> [English Version](../../configuration/send-lat-analysis.md) + +## 概述 + +SimAI 使用 `send_lat`（发送延迟）值来延迟 ns3 仿真中的数据包发送。这模拟了发起 NCCL 集合通信操作的软件开销。有两种机制控制此值。 + +--- + +## 1. send_latency_table（表查询） + +### 位置 +`astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h:134-152` + +### 工作原理 + +两个 7x3 表，按 `[algorithm][protocol]` 索引： +- `send_lat_table_nvlink[7][3]` — 用于同节点（NVLINK）通信 +- `send_lat_table_net[7][3]` — 用于跨节点（NET/IB）通信 + +算法索引：0=Tree, 1=Ring, 2=CollNetDirect, 3=CollNetChain, 4=NVLS, 5=NVLS_TREE, 6=PAT +协议索引：0=LL, 1=LL128, 2=Simple + +### 表值（单位：纳秒） + +#### NVLINK 表（节点内） +| 算法 | LL | LL128 | Simple | +|---|---|---|---| +| Tree | 7400 | 15250 | 12400 | +| Ring | 7200 | 15900 | 11800 | +| CollNetDirect | 0 | 0 | 3700 | +| CollNetChain | 0 | 0 | 2800 | +| NVLS | 0 | 0 | 25000 | +| NVLS_TREE | 0 | 0 | 25000 | +| PAT | 0 | 0 | 12000 | + +#### NET 表（节点间） +| 算法 | LL | LL128 | Simple | +|---|---|---|---| +| Tree | 11800 | 22500 | 22400 | +| Ring | 9300 | 18000 | 22400 | +| CollNetDirect | 0 | 0 | 31000 | +| CollNetChain | 0 | 0 | 30000 | +| NVLS | 0 | 0 | 18000 | +| NVLS_TREE | 0 | 0 | 20900 | +| PAT | 0 | 0 | 22000 | + +值 `0` 表示不支持的组合 → 回退到默认值 6000 ns。 + +### 链路类型检测（entry.h:159-162） + +```cpp +int gpus_per_node = request->flowTag.gpus_per_node; +bool same_node = (src / gpus_per_node) == (dst / gpus_per_node); +const int (*table)[3] = same_node ? send_lat_table_nvlink : send_lat_table_net; +``` + +### 调用栈 + +``` +NcclFlowModel::insert_packets() + → front_end_sim_send() + → entry.h: SendFlow() + → 读取 flowTag.algorithm, protocol, gpus_per_node + → 判断 same_node → 选择 nvlink 或 net 表 + → table[algo][proto] → send_lat (ns) + → AS_SEND_LAT 覆盖检查（最高优先级） + → send_lat *= 1000 → 将 ns 转换为 ps（ns3 Time 单位） + → appCon.Start(Time(send_lat)) +``` + +--- + +## 2. AS_SEND_LAT（环境变量覆盖） + +### 位置 +`entry.h:168-177` + +### 工作原理 + +```cpp +const char* send_lat_env = std::getenv("AS_SEND_LAT"); +if (send_lat_env) { + send_lat = std::stoi(send_lat_env); // 覆盖表查询 +} +send_lat *= 1000; // ns → ps +``` + +### 作用域 + +**AS_SEND_LAT 不在 SimCCL 中**。它作用于 ns3 网络前端层。 + +| 属性 | send_latency_table | AS_SEND_LAT | +|---|---|---| +| 位置 | entry.h（ns3 前端） | entry.h（ns3 前端） | +| 粒度 | 按 (algorithm, protocol, link_type) | 全局单一值 | +| 优先级 | 较低 | **最高**（覆盖表） | +| 用例 | 生产环境仿真 | A/B 实验、快速校准 | +| 单位 | 纳秒 | 纳秒 | + +### send_latency_table 能否完全替代 AS_SEND_LAT？ + +**不能。** 它们服务于不同目的： + +- `send_latency_table`：细粒度的按 (algo, proto, link) 延迟。用于精确的生产仿真。 +- `AS_SEND_LAT`：快速全局覆盖，用于 A/B 测试。设置 `AS_SEND_LAT=6` 使所有通信使用 6000 ps，无论算法/协议。 + +**建议**：生产环境使用表，AS_SEND_LAT 仅用于调试/对比实验。 + +--- + +## 3. 与 NCCL 2.30 延迟模型的比较 + +### NCCL 源码 +`nccl-2.30/src/graph/tuning.cc:150-174`（基础/硬件延迟表），`L380-435`（延迟计算） + +### 关键差异 + +SimAI 的 `send_lat_table` 是 NCCL 延迟模型的 **分桶近似**： + +| 方面 | SimAI entry.h | NCCL tuning.cc | +|---|---|---| +| 模型类型 | 固定 3D 表查询 | 每算法动态计算 | +| Ring 延迟 | 固定值（如 LL-NET 9300 ns） | `baseLat + nsteps * (intraLat + netOverhead)` | +| PAT 延迟 | 固定 22000 ns（Simple-NET） | `log2(nNodes) * (interLat/3.5) + nRanks * 2.8` | +| Tree 延迟 | 固定值 | `baseLat + 2*log2(nNodes)*interLat` | +| 依赖拓扑 | 仅链路类型（同/跨节点） | nNodes, nRanks, nsteps, netOverhead, cpuArch | +| 硬件特定 | 否（所有 GPU 使用同一表） | 按 CPU 厂商调整 `netOverhead` | + +**含义**：SimAI 的表提供合理但近似的 send_lat。对于 PAT 具体而言，表使用固定的 22us（NET-Simple），而 NCCL 计算 `log2(N)*(interLat/3.5)+N*2.8`，取决于实际节点数和节点间延迟。 + +--- + +## 4. 如何实验 + +### 前置条件 + +```bash +cd SimAI/ +./scripts/build.sh -c ns3 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +### 使用 send_lat_table 运行（默认） + +```bash +./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# 查看结果 +head -3 ncclFlowModel_EndToEnd.csv +``` + +### 使用 AS_SEND_LAT 覆盖运行 + +```bash +AS_SEND_LAT=7200 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +### 对比结果 + +| 条件 | Total Time | 备注 | +|---|---|---| +| 基线（send_lat_table） | 389,404 | 按(算法,协议,链路)查表 | +| AS_SEND_LAT=6 (6ns) | 1,772 | 发送延迟可忽略 | +| AS_SEND_LAT=6000 (6us) | 169,604 | 统一 6us | +| AS_SEND_LAT=7200 (7.2us) | 203,204 | 匹配 Ring+LL+NVLINK | + +详细分析请见 [快速入门 - A/B 实验](../getting_started/quickstart.md#5-验证-as_send_lat-效果ab-实验)。 + +--- + +> 最后编辑：2026-06-25 diff --git a/docs/CN/getting_started/installation.md b/docs/CN/getting_started/installation.md new file mode 100644 index 000000000..f6ad6192f --- /dev/null +++ b/docs/CN/getting_started/installation.md @@ -0,0 +1,115 @@ +# SimAI 安装指南 + +> [English Version](../../getting_started/installation.md) + +## 前置条件 + +| 要求 | 版本 | 备注 | +|---|---|---| +| 操作系统 | Linux (Ubuntu 20.04+) | 在 Ubuntu 20.04/22.04 上测试通过 | +| GCC/G++ | 9.4.0+ | 需要 C++17 支持 | +| Python3 | 3.8+ | 用于拓扑生成脚本 | +| CMake | 3.14+ | ns3 构建使用 | +| GPU/CUDA | **不需要** | 仿真只需 CPU | +| ninja | **不能安装** | ns3 构建与 ninja 冲突 | + +> **注意**：如使用 NGC 容器镜像（推荐用于 AICB workload 生成），需先移除 ninja： +> ```bash +> apt remove ninja-build && pip uninstall ninja +> ``` + +## 克隆和初始化 + +```bash +git clone https://github.com/aliyun/SimAI.git +cd SimAI/ + +# 初始化所有子模块（SimCCL、ns-3-alibabacloud、aicb 等） +git submodule update --init --recursive +git submodule update --remote +``` + +## 编译 SimAI-Analytical + +快速 busbw 仿真（无网络建模）： + +```bash +./scripts/build.sh -c analytical +``` + +输出：`bin/SimAI_analytical` + +## 编译 SimAI-Simulation (ns3) + +完整 ns3 网络仿真： + +```bash +./scripts/build.sh -c ns3 +``` + +输出：`bin/SimAI_simulator` + +> **注意**：此命令会删除并重建 `astra-sim-alibabacloud/extern/network_backend/ns3-interface/`。首次构建需 5-15 分钟。 + +## 编译 SimAI-Physical + +物理 RDMA 流量生成（需要 InfiniBand 硬件 + MPI）： + +```bash +# 安装 MPI 依赖 +sudo yum install openmpi openmpi-devel # 或 apt install libopenmpi-dev + +# 设置 MPI 路径 +export MPI_INCLUDE_PATH=/usr/include/openmpi-x86_64/ +export MPI_BIN_PATH=/usr/lib64/openmpi/bin/mpic++ + +# 构建 +./scripts/build.sh -c phy +``` + +输出：`bin/SimAI_phynet` + +## Docker 环境（推荐） + +```bash +# 进入容器 +docker exec -it bash + +# 容器内：编译 ns3 模式 +cd /path/to/SimAI +./scripts/build.sh -c ns3 + +# 验证 +ls -la bin/SimAI_simulator +``` + +## 清理构建 + +```bash +./scripts/build.sh -l ns3 # 清理 ns3 构建产物 +./scripts/build.sh -l analytical # 清理 analytical 构建 +./scripts/build.sh -l phy # 清理 physical 构建 +``` + +## 验证安装 + +```bash +# 检查二进制是否存在 +ls bin/SimAI_simulator bin/SimAI_analytical + +# 快速冒烟测试（8 GPU，单节点） +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# 预期：仿真完成，生成 ncclFlowModel_EndToEnd.csv +ls ncclFlowModel_EndToEnd.csv +``` + +--- + +> 最后编辑：2026-06-25 diff --git a/docs/CN/getting_started/quickstart.md b/docs/CN/getting_started/quickstart.md new file mode 100644 index 000000000..b5ae9048a --- /dev/null +++ b/docs/CN/getting_started/quickstart.md @@ -0,0 +1,203 @@ +# SimAI 快速入门指南 + +> [English Version](../../getting_started/quickstart.md) + +## 端到端运行指南 + +### 1. 编译 SimAI-Simulation (ns3) + +```bash +cd SimAI/ +./scripts/build.sh -c ns3 +``` + +产物：`bin/SimAI_simulator`（指向 ns3 构建输出的符号链接）。 + +### 2. 生成拓扑 + +使用拓扑生成脚本： + +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +#### 关键参数 + +| 参数 | 说明 | 单位/取值 | 示例 | +|---|---|---|---| +| `-topo` | 拓扑模板名称 | `Spectrum-X`, `AlibabaHPN`, `DCN+` | `-topo Spectrum-X` | +| `-g` | GPU 总数 | 整数 | `-g 8` | +| `-gt` | GPU 类型 | `A100`, `H100`, `H800`, `H20` | `-gt H20` | +| `-gps` | 每台服务器 GPU 数 | 整数（默认：8） | `-gps 8` | +| `-bw` | NIC 带宽（横向扩展） | 如 `100Gbps`, `200Gbps`, `400Gbps` | `-bw 200Gbps` | +| `-nvbw` | NVLink 带宽（纵向扩展） | 如 `2400Gbps`, `2880Gbps` | `-nvbw 2400Gbps` | +| `--ro` | Rail-Optimized 拓扑 | 标志位（无需值） | `--ro` | +| `-psn` | PSW 交换机数量 | 整数 | `-psn 64` | +| `--dp` | 双平面 | 标志位 | `--dp` | +| `--dt` | 双 ToR | 标志位 | `--dt` | +| `-nl` | NVLink 延迟 | 如 `0.000025ms` | `-nl 0.000025ms` | +| `-l` | NIC 延迟 | 如 `0.0005ms` | `-l 0.0005ms` | + +#### 输出文件名规则 + +输出文件命名格式：`{模板}_{g}g_{gps}gps_{bw}_{gt}` + +示例： +- `--ro -g 8 -gt H20 -bw 200Gbps` → `Rail_Opti_SingleToR_8g_8gps_200Gbps_H20` +- `-topo Spectrum-X -g 128 -gt A100 -bw 100Gbps` → `Spectrum-X_128g_8gps_100Gbps_A100` + +### 3. 运行仿真 + +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +#### 仿真器参数 + +| 参数 | 说明 | 默认值 | +|---|---|---| +| `-t` | 线程数（多线程加速） | 1 | +| `-w` | Workload 文件路径 | 必填 | +| `-n` | 网络拓扑文件路径 | 必填 | +| `-c` | 配置文件路径 | 必填 | + +#### 环境变量（命令前设置） + +| 变量 | 单位 | 默认值 | 说明 | +|---|---|---|---| +| `AS_SEND_LAT` | **纳秒 (ns)** | 未设置（使用 send_lat_table） | 覆盖每个 flow 的发送延迟。覆盖所有表查询。 | +| `AS_NVLS_ENABLE` | - | `0` | 启用 NVLS 算法用于 AllReduce | +| `AS_PXN_ENABLE` | - | `0` | 启用 PXN 跨节点代理 | +| `AS_LOG_LEVEL` | - | `INFO` | 日志级别：DEBUG/INFO/WARNING/ERROR | + +> **重要提示**：`AS_SEND_LAT` 的单位是**纳秒**，不是微秒。send_lat_table 默认值范围为 6000-22000 ns（6-22 μs）。设置 `AS_SEND_LAT=6` 意味着 6 ns（实际上等于禁用发送延迟）。 + +### 4. 预期输出 + +仿真在当前工作目录生成以下文件： + +| 文件 | 说明 | +|---|---| +| `ncclFlowModel_EndToEnd.csv` | 端到端迭代时序汇总 | +| `ncclFlowModel_detailed_N.csv` | 逐层详细时序（N = 节点数） | +| `ncclFlowModel_detailed_flows.csv` | SimCCL 点对点流分解 | +| `ncclFlowModel_*_dimension_utilization_*.csv` | 通信组利用率 | + +验证输出： +```bash +ls ncclFlowModel_*.csv +head -5 ncclFlowModel_EndToEnd.csv +``` + +### 5. 验证 AS_SEND_LAT 效果（A/B 实验） + +**设计原则**：单变量实验。相同拓扑、相同 workload，仅改变 `AS_SEND_LAT`。 + +```bash +# 基线：使用 send_lat_table（无覆盖） +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# 记录：EndToEnd.csv 中的 total time + +# 实验组：使用 AS_SEND_LAT=6000 覆盖（6 μs，接近默认 Ring+LL+NVLINK=7200ns） +AS_SEND_LAT=6000 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# 记录：EndToEnd.csv 中的 total time +``` + +#### 实测结果（8 GPU, H20, microAllReduce.txt） + +| 条件 | Total Time | 备注 | +|---|---|---| +| 基线（send_lat_table 查表） | 389,404 | 表值：Ring+LL+NVLINK=7200ns 等 | +| AS_SEND_LAT=6（6 ns） | 1,772 | 发送延迟可忽略不计 | +| AS_SEND_LAT=6000（6 μs） | 169,604 | 统一 6μs，低于表平均值 | +| AS_SEND_LAT=7200（7.2 μs） | 203,204 | 匹配 Ring+LL+NVLINK 表值 | + +**分析**： +- 基线使用按（算法、协议、链路类型）分类的表值，平均值高于任何单一值 +- `AS_SEND_LAT` 将所有 flow 覆盖为同一个值，失去了按类型差异化的能力 +- 端到端时间受拓扑、拥塞、协议、算法和数据量等多因素影响——不仅仅是 send_lat +- 使用此实验验证发送延迟变化的相对影响 + +--- + +## 冒烟测试（8 GPU，单节点） + +完整单节点测试可在 1 分钟内完成： + +```bash +# 1. 生成拓扑 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +# 2. 运行 +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# 3. 验证 +cat ncclFlowModel_EndToEnd.csv | head -3 +``` + +此测试验证节点内通信（NVLINK）。所有 flow 在单个节点内完成。 + +--- + +## 跨节点测试（16 GPU，2 节点） + +验证跨节点（NET/IB）行为： + +### 1. 创建 16 GPU Workload + +修改 `example/microAllReduce.txt`，设置 `all_gpus: 16`： + +```bash +# 复制并修改 +cp example/microAllReduce.txt example/microAllReduce_16g.txt +# 编辑：将 "all_gpus: 8" 改为 "all_gpus: 16" +sed -i 's/all_gpus: 8/all_gpus: 16/' example/microAllReduce_16g.txt +``` + +### 2. 生成 16 GPU 拓扑（2 节点 x 8 GPU） + +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 16 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +输出文件：`Rail_Opti_SingleToR_16g_8gps_200Gbps_H20` + +### 3. 运行跨节点仿真 + +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce_16g.txt \ + -n ./Rail_Opti_SingleToR_16g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +此命令同时测试 NVLINK（节点内）和 NET（节点间）两种通信路径。send_lat_table 对不同链路类型使用不同的延迟值。 + +--- + +## 相关文档 + +- [安装指南](installation.md) — 所有模式的构建说明 +- [环境变量](../configuration/env-variables.md) — 完整变量参考 +- [构建选项](../configuration/build-options.md) — 构建模式和参数 +- [send_lat 分析](../configuration/send-lat-analysis.md) — send_lat 机制深入分析 +- [SimCCL 集成](../../../SimCCL/docs/CN/integration/integration-with-simai.md) — SimCCL + SimAI 集成 + +--- + +> 最后编辑：2026-06-25 diff --git a/docs/README.md b/docs/README.md new file mode 100644 index 000000000..d51d98c07 --- /dev/null +++ b/docs/README.md @@ -0,0 +1,27 @@ +# SimAI Documentation + +## Navigation / 文档导航 + +### English + +| Category | Document | Description | +|---|---|---| +| **Tutorial** | [Tutorial.md](Tutorial.md) | Complete SimAI usage guide (Analytical + Simulation + Physical) | +| **Getting Started** | [installation.md](getting_started/installation.md) | Build instructions for all modes | +| | [quickstart.md](getting_started/quickstart.md) | Quick start guide with end-to-end examples | +| **Configuration** | [env-variables.md](configuration/env-variables.md) | Environment variable reference | +| | [build-options.md](configuration/build-options.md) | Build configuration options | +| | [send-lat-analysis.md](configuration/send-lat-analysis.md) | Send latency mechanism deep dive | + +### 中文 (Chinese) + +| 分类 | 文档 | 说明 | +|---|---|---| +| **教程** | [Tutorial.md](CN/Tutorial.md) | SimAI 完整使用指南（Analytical + Simulation + Physical） | +| **入门** | [installation.md](CN/getting_started/installation.md) | 编译指南 | +| | [quickstart.md](CN/getting_started/quickstart.md) | 快速开始（端到端示例） | +| **配置** | [env-variables.md](CN/configuration/env-variables.md) | 环境变量参考 | +| | [build-options.md](CN/configuration/build-options.md) | 构建配置选项 | +| | [send-lat-analysis.md](CN/configuration/send-lat-analysis.md) | Send latency 机制分析 | + +> Note: 中文文档以英文版本为准，如有出入请参考英文原文。 diff --git a/docs/Tutorial.md b/docs/Tutorial.md index 04d637e33..292de4876 100755 --- a/docs/Tutorial.md +++ b/docs/Tutorial.md @@ -12,7 +12,7 @@ SimAI is a comprehensive large-scale AI training simulation toolkit that provide * *Scale-out Bandwidth Selection*: Research cost-effective bandwidth configurations for various GPU performances -> 💡 *Currently supports manual busbw.yaml configuration. Automatic busbw inference based on parallel scenarios will be open-sourced soon. Stay tuned and feel free to contact us for more details. ✨* +> 💡 *The current open-source binary computes busbw automatically from `-nv`/`-nic`/`-n_p_s` (the default; see the Analytical Simulation section below). The manual `busbw.yaml` path (`-busbw`) is not wired in the current open-source binary — refer to earlier SimAI versions for that workflow. Feel free to contact us for more details. ✨* 2. **SimAI-Simulation(NS-3)** - A high-fidelity, full-stack simulation tool that can theoretically integrate with any pure network simulator. It provides fine-grained reproduction of communication behaviors during LLM training. Currently supports NS-3 as the network backend (we encourage integration of new network simulation tools). Key research areas include: @@ -94,7 +94,13 @@ EP: ## 🖥️ Analytical Simulation -To run the analytical simulation, use the following command: +To run the analytical simulation, use the following command (automatic busbw — the default, supported by the current binary): + +```bash +$ ./bin/SimAI_analytical -w ./example/workload_analytical.txt -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 -r example- +``` + +> Note: The `-busbw example/busbw.yaml` form below (user-defined busbw) is **not supported** by the current open-source analytical binary — the `-busbw` flag is not parsed, so the command prints usage and exits. It is kept for reference; see earlier SimAI versions for the manual `busbw.yaml` workflow. ```bash $ ./bin/SimAI_analytical -w example/workload_analytical.txt -g 9216 -g_p_s 8 -r test- -busbw example/busbw.yaml diff --git a/docs/configuration/build-options.md b/docs/configuration/build-options.md new file mode 100644 index 000000000..00ba7d7fe --- /dev/null +++ b/docs/configuration/build-options.md @@ -0,0 +1,77 @@ +# SimAI Build Options + +> [中文版](../CN/configuration/build-options.md) + +## Build Modes + +| Mode | Command | Output Binary | Description | +|---|---|---|---| +| analytical | `./scripts/build.sh -c analytical` | `bin/SimAI_analytical` | Fast busbw-based simulation (no network modeling) | +| ns3 | `./scripts/build.sh -c ns3` | `bin/SimAI_simulator` | Full ns3 network simulation (high fidelity) | +| phy | `./scripts/build.sh -c phy` | `bin/SimAI_phynet` | Physical RDMA traffic generation (requires IB hardware) | + +## Clean Build + +```bash +./scripts/build.sh -l ns3 # Clean ns3 build artifacts +./scripts/build.sh -l analytical # Clean analytical build +./scripts/build.sh -l phy # Clean physical build +``` + +## SimAI_simulator Parameters (ns3 mode) + +| Parameter | Long Form | Description | Default | +|---|---|---|---| +| `-t` | `--thread` | Number of threads for multi-threading acceleration | 1 | +| `-w` | `--workload` | Path to workload file | Required | +| `-n` | `--network-topo` | Path to network topology file | Required | +| `-c` | `--config` | Path to simulation config file | Required | + +Typical usage: +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +> **Threading**: Recommended 8-16 threads for multi-threaded mode. Higher thread count reduces simulation time for large topologies. + +## SimAI_analytical Parameters + +| Parameter | Long Form | Description | Default | +|---|---|---|---| +| `-w` | `--workload` | Path to workload file | Required | +| `-g` | `--gpus` | Total GPU count | Required | +| `-g_p_s` | `--gpus-per-server` | GPUs per server (scale-up size) | Required | +| `-r` | `--result` | Output file path/prefix | `./results/` | +| `-busbw` | `--bus-bandwidth` | Path to busbw.yaml (user-defined busbw). **Not wired in the current open-source binary** — see note below | n/a | +| `-v` | `--visual` | Generate visualization files | Off | +| `-nv` | - | NVLink bandwidth (GB/s) for auto busbw (default method) | Recommended | +| `-nic` | - | NIC bandwidth (GB/s) for auto busbw (default method) | Recommended | +| `-n_p_s` | - | NICs per server for auto busbw (default method) | Recommended | + +Typical usage (automatic busbw — the default, supported by the current binary): +```bash +./bin/SimAI_analytical \ + -w ./example/workload_analytical.txt \ + -g 9216 -nv 360 -nic 48.5 -n_p_s 8 -g_p_s 8 \ + -r example- +``` + +> Note: The `-busbw example/busbw.yaml` form below (user-defined busbw) is **not supported** by the current open-source analytical binary — the `-busbw` flag is not parsed, so the command prints usage and exits. It is kept for reference; see earlier SimAI versions for the manual `busbw.yaml` workflow. +```bash +./bin/SimAI_analytical \ + -w example/workload_analytical.txt \ + -g 9216 -g_p_s 8 \ + -r test- \ + -busbw example/busbw.yaml +``` + +## Topology Generator Parameters + +See [Quick Start - Generate Topology](../getting_started/quickstart.md#2-generate-topology) for full parameter reference. + +--- + +> Last edited: 2026-06-25 diff --git a/docs/configuration/env-variables.md b/docs/configuration/env-variables.md new file mode 100644 index 000000000..079f499cb --- /dev/null +++ b/docs/configuration/env-variables.md @@ -0,0 +1,65 @@ +# SimAI Environment Variables Reference + +> [中文版](../CN/configuration/env-variables.md) + +## Simulation Environment Variables + +These variables control the behavior of `SimAI_simulator` (ns3 mode). + +| Variable | Default | Unit | Scope | Description | +|---|---|---|---|---| +| `AS_SEND_LAT` | Not set (use table) | **nanoseconds (ns)** | entry.h (ns3 frontend) | Override per-flow send latency. Highest priority — overrides all send_lat_table lookups. | +| `AS_NVLS_ENABLE` | `0` | - | MockNcclGroup.cc | Enable NVLS algorithm for AllReduce (H20/H100/H800) | +| `AS_NVLSTREE_ENABLE` | `0` | - | MockNcclGroup.cc | Enable NVLS Tree algorithm | +| `AS_PXN_ENABLE` | `0` | - | MockNcclGroup.cc | Enable PXN cross-node proxy | +| `AS_LOG_LEVEL` | `INFO` | - | System-wide | Log verbosity: `DEBUG`, `INFO`, `WARNING`, `ERROR` | + +### AS_SEND_LAT Details + +**Unit: nanoseconds (ns)**. The send_lat_table default values range from 6000 to 22000 ns depending on algorithm, protocol, and link type. + +| Example Value | Meaning | +|---|---| +| Not set | Use send_lat_table[algo][proto] (7200ns for Ring+LL+NVLINK, 22000ns for PAT+Simple+NET, etc.) | +| `AS_SEND_LAT=6` | 6 ns — effectively disables send latency (for fast testing) | +| `AS_SEND_LAT=7200` | 7200 ns = 7.2 us — matches Ring+LL+NVLINK table value | +| `AS_SEND_LAT=22000` | 22000 ns = 22 us — matches PAT+Simple+NET table value | + +> **Warning**: Some older documentation (Tutorial.md) describes AS_SEND_LAT as "unit is us, default 6". This is inaccurate. The actual unit is **nanoseconds** as confirmed by code analysis and experiment. + +#### Quick Commands + +```bash +# Prerequisites +cd SimAI/ && ./scripts/build.sh -c ns3 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +# Baseline (use send_lat_table) +./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# Override send latency (7200ns = Ring+LL+NVLINK table value) +AS_SEND_LAT=7200 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +For full A/B experiment guide, see [Quick Start](../getting_started/quickstart.md#5-verify-as_send_lat-effect-ab-experiment). + +## SimCCL Integration Variables + +SimCCL (the collective communication modeling layer) has its own environment variables. See: +- [SimCCL Environment Variables](../../SimCCL/docs/configuration/env-variables.md) — Standalone variables +- [SimCCL Integration Guide](../../SimCCL/docs/integration/integration-with-simai.md) — Cross-module variables + +## Relationship with send_lat_table + +The `send_lat_table` in `entry.h` provides per-(algorithm, protocol, link_type) send latency values. When `AS_SEND_LAT` is not set, each flow uses the table value matching its algorithm and protocol. When set, ALL flows use the override value. + +For detailed analysis of the table mechanism, see [send-lat-analysis.md](send-lat-analysis.md). + +--- + +> Last edited: 2026-06-25 diff --git a/docs/configuration/send-lat-analysis.md b/docs/configuration/send-lat-analysis.md new file mode 100644 index 000000000..a887146bc --- /dev/null +++ b/docs/configuration/send-lat-analysis.md @@ -0,0 +1,178 @@ +# send_lat Deep Dive: send_latency_table vs AS_SEND_LAT + +> [中文版](../CN/configuration/send-lat-analysis.md) + +## Overview + +SimAI uses a `send_lat` (send latency) value to delay packet transmission in ns3 simulation. This models the software overhead of initiating a NCCL collective communication operation. Two mechanisms control this value. + +--- + +## 1. send_latency_table (Table Lookup) + +### Location +`astra-sim-alibabacloud/astra-sim/network_frontend/ns3/entry.h:134-152` + +### How It Works + +Two 7x3 tables indexed by `[algorithm][protocol]`: +- `send_lat_table_nvlink[7][3]` — for same-node (NVLINK) communication +- `send_lat_table_net[7][3]` — for cross-node (NET/IB) communication + +Algorithm index: 0=Tree, 1=Ring, 2=CollNetDirect, 3=CollNetChain, 4=NVLS, 5=NVLS_TREE, 6=PAT +Protocol index: 0=LL, 1=LL128, 2=Simple + +### Table Values (unit: nanoseconds) + +#### NVLINK Table (intra-node) +| Algorithm | LL | LL128 | Simple | +|---|---|---|---| +| Tree | 7400 | 15250 | 12400 | +| Ring | 7200 | 15900 | 11800 | +| CollNetDirect | 0 | 0 | 3700 | +| CollNetChain | 0 | 0 | 2800 | +| NVLS | 0 | 0 | 25000 | +| NVLS_TREE | 0 | 0 | 25000 | +| PAT | 0 | 0 | 12000 | + +#### NET Table (inter-node) +| Algorithm | LL | LL128 | Simple | +|---|---|---|---| +| Tree | 11800 | 22500 | 22400 | +| Ring | 9300 | 18000 | 22400 | +| CollNetDirect | 0 | 0 | 31000 | +| CollNetChain | 0 | 0 | 30000 | +| NVLS | 0 | 0 | 18000 | +| NVLS_TREE | 0 | 0 | 20900 | +| PAT | 0 | 0 | 22000 | + +Value `0` means unsupported combination -> falls back to default 6000 ns. + +### Link Type Detection (entry.h:159-162) + +```cpp +int gpus_per_node = request->flowTag.gpus_per_node; +bool same_node = (src / gpus_per_node) == (dst / gpus_per_node); +const int (*table)[3] = same_node ? send_lat_table_nvlink : send_lat_table_net; +``` + +### Call Stack + +``` +NcclFlowModel::insert_packets() + → front_end_sim_send() + → entry.h: SendFlow() + → Read flowTag.algorithm, protocol, gpus_per_node + → Determine same_node → select nvlink or net table + → table[algo][proto] → send_lat (ns) + → AS_SEND_LAT override check (highest priority) + → send_lat *= 1000 → convert ns to ps (ns3 Time unit) + → appCon.Start(Time(send_lat)) +``` + +--- + +## 2. AS_SEND_LAT (Environment Variable Override) + +### Location +`entry.h:168-177` + +### How It Works + +```cpp +const char* send_lat_env = std::getenv("AS_SEND_LAT"); +if (send_lat_env) { + send_lat = std::stoi(send_lat_env); // overrides table lookup +} +send_lat *= 1000; // ns → ps +``` + +### Scope + +**AS_SEND_LAT is NOT in SimCCL**. It operates in the ns3 network frontend layer. + +| Property | send_latency_table | AS_SEND_LAT | +|---|---|---| +| Location | entry.h (ns3 frontend) | entry.h (ns3 frontend) | +| Granularity | Per-(algorithm, protocol, link_type) | Global single value | +| Priority | Lower | **Highest** (overrides table) | +| Use case | Production simulation | A/B experiments, quick calibration | +| Unit | Nanoseconds | Nanoseconds | + +### Can send_latency_table Fully Replace AS_SEND_LAT? + +**No.** They serve different purposes: + +- `send_latency_table`: Fine-grained per-(algo, proto, link) latency. For accurate production simulations. +- `AS_SEND_LAT`: Quick global override for A/B testing. Setting `AS_SEND_LAT=6` makes ALL communications use 6000 ps regardless of algorithm/protocol. + +**Recommendation**: Use table for production, AS_SEND_LAT for debugging/comparison experiments only. + +--- + +## 3. Comparison with NCCL 2.30 Latency Model + +### NCCL Source +`nccl-2.30/src/graph/tuning.cc:150-174` (base/hw latency tables), `L380-435` (latency calculation) + +### Key Differences + +SimAI's `send_lat_table` is a **bucketed approximation** of NCCL's latency model: + +| Aspect | SimAI entry.h | NCCL tuning.cc | +|---|---|---| +| Model type | Fixed 3D table lookup | Dynamic calculation per-algorithm | +| Ring latency | Fixed value (e.g., 9300 ns for LL-NET) | `baseLat + nsteps * (intraLat + netOverhead)` | +| PAT latency | Fixed 22000 ns (Simple-NET) | `log2(nNodes) * (interLat/3.5) + nRanks * 2.8` | +| Tree latency | Fixed value | `baseLat + 2*log2(nNodes)*interLat` | +| Depends on topology | Only link type (same/cross node) | nNodes, nRanks, nsteps, netOverhead, cpuArch | +| Hardware-specific | Not (same table for all GPUs) | Per-CPU-vendor `netOverhead` adjustment | + +**Implication**: SimAI's table gives reasonable but approximate send_lat. For PAT specifically, the table uses a fixed 22us for NET-Simple, while NCCL computes `log2(N)*(interLat/3.5)+N*2.8` which depends on the actual node count and inter-node latency. + +--- + +## 4. How to Experiment + +### Prerequisites + +```bash +cd SimAI/ +./scripts/build.sh -c ns3 +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +### Running with send_lat_table (default) + +```bash +./bin/SimAI_simulator -t 8 -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# Check results +head -3 ncclFlowModel_EndToEnd.csv +``` + +### Running with AS_SEND_LAT override + +```bash +AS_SEND_LAT=7200 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +### Comparing Results + +| Condition | Total Time | Notes | +|---|---|---| +| Baseline (send_lat_table) | 389,404 | Per-(algo,proto,link) table values | +| AS_SEND_LAT=6 (6ns) | 1,772 | Negligible send latency | +| AS_SEND_LAT=6000 (6us) | 169,604 | Uniform 6us | +| AS_SEND_LAT=7200 (7.2us) | 203,204 | Matches Ring+LL+NVLINK | + +For detailed analysis of what these numbers mean, see [Quick Start - A/B Experiment](../getting_started/quickstart.md#5-verify-as_send_lat-effect-ab-experiment). + +--- + +> Last edited: 2026-06-25 diff --git a/docs/getting_started/installation.md b/docs/getting_started/installation.md new file mode 100644 index 000000000..d238cf7b3 --- /dev/null +++ b/docs/getting_started/installation.md @@ -0,0 +1,132 @@ +# SimAI Installation Guide + +> [中文版](../CN/getting_started/installation.md) + +## Prerequisites + +| Requirement | Version | Notes | +|---|---|---| +| OS | Linux (Ubuntu 20.04+) | Tested on Ubuntu 20.04/22.04 | +| GCC/G++ | 9.4.0+ | C++17 support required | +| Python3 | 3.8+ | For topology generation scripts | +| CMake | 3.14+ | Used by ns3 build | +| GPU/CUDA | **Not required** | Simulation is CPU-only | +| ninja | **Must NOT be installed** | ns3 build conflicts with ninja | + +> **Note**: If using NGC container images (recommended for AICB workload generation), remove ninja first: +> ```bash +> apt remove ninja-build && pip uninstall ninja +> ``` + +## Clone and Initialize + +```bash +git clone https://github.com/aliyun/SimAI.git +cd SimAI/ + +# Initialize all submodules (SimCCL, ns-3-alibabacloud, aicb, etc.) +git submodule update --init --recursive +git submodule update --remote +``` + +## Compile SimAI-Analytical + +Fast busbw-based simulation (no network modeling): + +```bash +./scripts/build.sh -c analytical +``` + +Output: `bin/SimAI_analytical` + +## Compile SimAI-Simulation (ns3) + +Full-stack ns3 network simulation: + +```bash +./scripts/build.sh -c ns3 +``` + +Output: `bin/SimAI_simulator` + +> **Note**: Default mock version is v2.30 (protocol-aware, PAT support). To use legacy v2.20: `SIMAI_NCCL_VERSION=v2.20 ./scripts/build.sh -c ns3` + +> **Note**: This command removes and rebuilds `astra-sim-alibabacloud/extern/network_backend/ns3-interface/`. First build takes 5-15 minutes depending on hardware. + +## Compile SimAI-Physical + +Physical RDMA traffic generation (requires InfiniBand hardware + MPI): + +```bash +# Install MPI dependencies +sudo yum install openmpi openmpi-devel # or apt install libopenmpi-dev + +# Set MPI paths +export MPI_INCLUDE_PATH=/usr/include/openmpi-x86_64/ +export MPI_BIN_PATH=/usr/lib64/openmpi/bin/mpic++ + +# Build +./scripts/build.sh -c phy +``` + +Output: `bin/SimAI_phynet` + +## Docker Environment (Recommended) + +```bash +# Enter container +docker exec -it bash + +# Inside container: compile ns3 mode +cd /path/to/SimAI +./scripts/build.sh -c ns3 + +# Verify +ls -la bin/SimAI_simulator +``` + +## Clean Build + +To remove build artifacts and start fresh: + +```bash +./scripts/build.sh -l ns3 # Clean ns3 build +./scripts/build.sh -l analytical # Clean analytical build +./scripts/build.sh -l phy # Clean physical build +``` + +## Verify Installation + +```bash +# Check binaries exist +ls bin/SimAI_simulator bin/SimAI_analytical + +# Quick smoke test (8 GPU, single node) +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# Expected: simulation completes, generates ncclFlowModel_EndToEnd.csv +ls ncclFlowModel_EndToEnd.csv +``` + +--- + +## Compile SimCCL Standalone (Optional) + +For independent collective operation analysis without ns3: + +```bash +cd SimCCL/standalone +bash build.sh v2.30 +``` + +Output: `build/simccl-standalone` (~300KB, no ns3 dependency, CPU-only) + +--- + +> Last edited: 2026-07-07 diff --git a/docs/getting_started/quickstart.md b/docs/getting_started/quickstart.md new file mode 100644 index 000000000..7efffc730 --- /dev/null +++ b/docs/getting_started/quickstart.md @@ -0,0 +1,204 @@ +# SimAI Quick Start Guide + +> [中文版](../CN/getting_started/quickstart.md) + +## End-to-End Running Guide + +### 1. Compile SimAI-Simulation (ns3) + +```bash +cd SimAI/ +./scripts/build.sh -c ns3 +``` + +Produces: `bin/SimAI_simulator` (symlink to the ns3 build output). + +### 2. Generate Topology + +Use the topology generator script: + +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +#### Key Parameters + +| Parameter | Description | Unit/Values | Example | +|---|---|---|---| +| `-topo` | Template name | `Spectrum-X`, `AlibabaHPN`, `DCN+` | `-topo Spectrum-X` | +| `-g` | Total GPU count | Integer | `-g 8` | +| `-gt` | GPU type | `A100`, `H100`, `H800`, `H20` | `-gt H20` | +| `-gps` | GPUs per server | Integer (default: 8) | `-gps 8` | +| `-bw` | NIC bandwidth (scale-out) | e.g., `100Gbps`, `200Gbps`, `400Gbps` | `-bw 200Gbps` | +| `-nvbw` | NVLink bandwidth (scale-up) | e.g., `2400Gbps`, `2880Gbps` | `-nvbw 2400Gbps` | +| `--ro` | Rail-optimized topology | Flag (no value) | `--ro` | +| `-psn` | PSW switch number | Integer | `-psn 64` | +| `--dp` | Dual-plane | Flag | `--dp` | +| `--dt` | Dual-ToR | Flag | `--dt` | +| `-nl` | NVLink latency | e.g., `0.000025ms` | `-nl 0.000025ms` | +| `-l` | NIC latency | e.g., `0.0005ms` | `-l 0.0005ms` | + +#### Output Filename Convention + +Output file is named: `{template}_{g}g_{gps}gps_{bw}_{gt}` + +Examples: +- `--ro -g 8 -gt H20 -bw 200Gbps` → `Rail_Opti_SingleToR_8g_8gps_200Gbps_H20` +- `-topo Spectrum-X -g 128 -gt A100 -bw 100Gbps` → `Spectrum-X_128g_8gps_100Gbps_A100` + +### 3. Run Simulation + +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +#### Simulator Parameters + +| Parameter | Description | Default | +|---|---|---| +| `-t` | Number of threads (multi-threading acceleration) | 1 | +| `-w` | Path to workload file | Required | +| `-n` | Path to network topology file | Required | +| `-c` | Path to config file | Required | + +#### Environment Variables (set before command) + +| Variable | Unit | Default | Description | +|---|---|---|---| +| `AS_SEND_LAT` | **nanoseconds (ns)** | Not set (use send_lat_table) | Override per-flow send latency. Overrides all table lookups. | +| `AS_NVLS_ENABLE` | - | `0` | Enable NVLS algorithm for AllReduce | +| `AS_PXN_ENABLE` | - | `0` | Enable PXN cross-node proxy | +| `AS_LOG_LEVEL` | - | `INFO` | Log verbosity: DEBUG/INFO/WARNING/ERROR | + +> **Important**: `AS_SEND_LAT` unit is **nanoseconds**, not microseconds. The send_lat_table default values range from 6000-22000 ns (6-22 us). Setting `AS_SEND_LAT=6` means 6 ns (effectively disabling send latency). + +### 4. Expected Output + +Simulation generates these files in the current working directory: + +| File | Description | +|---|---| +| `ncclFlowModel_EndToEnd.csv` | End-to-end iteration timing summary | +| `ncclFlowModel_detailed_N.csv` | Per-layer detailed timing (N = node count) | +| `ncclFlowModel_detailed_flows.csv` | SimCCL point-to-point flow decomposition | +| `ncclFlowModel_*_dimension_utilization_*.csv` | Communication group utilization | + +Verify output: +```bash +ls ncclFlowModel_*.csv +head -5 ncclFlowModel_EndToEnd.csv +``` + +### 5. Verify AS_SEND_LAT Effect (A/B Experiment) + +**Design principle**: Single-variable experiment. Same topo, same workload, only change `AS_SEND_LAT`. + +```bash +# Baseline: use send_lat_table (no override) +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# Record: total time from EndToEnd.csv + +# Experiment: override with AS_SEND_LAT=6000 (6 us, close to default Ring+LL+NVLINK=7200ns) +AS_SEND_LAT=6000 ./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +# Record: total time from EndToEnd.csv +``` + +#### Measured Results (8 GPU, H20, microAllReduce.txt) + +| Condition | Total Time | Notes | +|---|---|---| +| Baseline (send_lat_table) | 389,404 | Table values: Ring+LL+NVLINK=7200ns, etc. | +| AS_SEND_LAT=6 (6 ns) | 1,772 | Negligible send latency | +| AS_SEND_LAT=6000 (6 us) | 169,604 | Uniform 6us, lower than table average | +| AS_SEND_LAT=7200 (7.2 us) | 203,204 | Matches Ring+LL+NVLINK table value | + +**Analysis**: +- The baseline uses per-(algorithm, protocol, link_type) table values, averaging higher than any single value +- `AS_SEND_LAT` overrides ALL flows to one value, losing per-type differentiation +- End-to-end time is affected by topology, congestion, protocol, algorithm, and payload size — not just send_lat +- Use this experiment to verify the relative impact of send latency changes + +--- + +## Smoke Test (8 GPU, Single Node) + +Complete single-node test in under 1 minute: + +```bash +# 1. Generate topology +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 8 -gt H20 -bw 200Gbps -nvbw 2400Gbps + +# 2. Run +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce.txt \ + -n ./Rail_Opti_SingleToR_8g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf + +# 3. Verify +cat ncclFlowModel_EndToEnd.csv | head -3 +``` + +This tests intra-node communication (NVLINK). All flows stay within one node. + +--- + +## Cross-Node Test (16 GPU, 2 Nodes) + +To verify cross-node (NET/IB) behavior: + +### 1. Create 16 GPU Workload + +Modify `example/microAllReduce.txt` to set `all_gpus: 16`: + +```bash +# Copy and modify +cp example/microAllReduce.txt example/microAllReduce_16g.txt +# Edit: change "all_gpus: 8" to "all_gpus: 16" +sed -i 's/all_gpus: 8/all_gpus: 16/' example/microAllReduce_16g.txt +``` + +### 2. Generate 16 GPU Topology (2 nodes x 8 GPUs) + +```bash +python3 ./astra-sim-alibabacloud/inputs/topo/gen_Topo_Template.py \ + --ro -g 16 -gt H20 -bw 200Gbps -nvbw 2400Gbps +``` + +Output: `Rail_Opti_SingleToR_16g_8gps_200Gbps_H20` + +### 3. Run Cross-Node Simulation + +```bash +./bin/SimAI_simulator -t 8 \ + -w ./example/microAllReduce_16g.txt \ + -n ./Rail_Opti_SingleToR_16g_8gps_200Gbps_H20 \ + -c ./astra-sim-alibabacloud/inputs/config/SimAI.conf +``` + +This exercises both NVLINK (intra-node) and NET (inter-node) communication paths. The send_lat_table applies different values for each link type. + +--- + +## Related Documentation + +- [Installation Guide](installation.md) — Build instructions for all modes +- [Environment Variables](../configuration/env-variables.md) — Complete variable reference +- [Build Options](../configuration/build-options.md) — Build modes and parameters +- [send_lat Analysis](../configuration/send-lat-analysis.md) — Deep dive into send_lat mechanism +- [SimCCL Integration](../../SimCCL/docs/integration/integration-with-simai.md) — SimCCL + SimAI integration +- [SimCCL Standalone Guide](../../SimCCL/docs/getting_started/quickstart.md) — Independent flow analysis + +--- + +> Last edited: 2026-07-07