[NPU]Ascend NPU Performance Profiling Guide and Ascend NPU Operator Development Guide (#25384)

.codespellrc

@@ -1,3 +1,3 @@
 [codespell]
-ignore-words-list = ans, als, hel, boostrap, childs, te, vas, hsa, ment, cann, thi, makro, wil, rouge, PRIS, ather, MIS, medias, allready, inout, nd, fo, visibles, nothink, renderD, ond, tbe
+ignore-words-list = ans, als, hel, boostrap, childs, te, vas, hsa, ment, cann, thi, makro, wil, rouge, PRIS, ather, MIS, medias, allready, inout, nd, fo, visibles, nothink, renderD, ond, tbe, CopyIn
 skip = *.json, *.jsonl, *.patch, *.txt, *.lock

docs_new/docs.json

@@ -905,6 +905,8 @@
                   "docs/hardware-platforms/ascend-npus/ascend_npu_qwen3_5_examples",
                   "docs/hardware-platforms/ascend-npus/ascend_npu_glm5_examples",
                   "docs/hardware-platforms/ascend-npus/ascend_npu_environment_variables",
+                  "docs/hardware-platforms/ascend-npus/ascend_npu_profiling",
+                  "docs/hardware-platforms/ascend-npus/ascend_npu_operator_development",
                   "docs/hardware-platforms/ascend-npus/ascend_npu_faq"
                 ]
               },

docs_new/docs/hardware-platforms/ascend-npus/ascend_npu_operator_development.mdx

@@ -0,0 +1,513 @@
+---
+title: "Ascend NPU Operator Development Guide"
+description: "How to develop custom operators (Ascend C / Triton) for Ascend NPU and integrate them into the SGLang inference engine."
+---
+
+## Overview
+
+[SGL-Kernel-NPU](https://github.com/sgl-project/sgl-kernel-npu) is the official
+operator library provided by the SGLang framework for Ascend NPU. It includes
+two types of operator implementations:
+
+1. **Ascend C operators**: High-performance C++ kernels written in Ascend C,
+   compiled into `libsgl_kernel_npu.so`, and registered through PyTorch's custom
+   operator mechanism (`TORCH_LIBRARY_FRAGMENT`). Called in SGLang via
+   `torch.ops.npu.<op_name>()`.
+2. **Triton operators**: Python kernels written in Triton, adapted for Ascend
+   NPU. Called directly via `from sgl_kernel_npu.xxx import ...`.
+
+When SGLang detects an NPU device, it automatically loads `sgl_kernel_npu` and
+uses its operators in place of GPU counterparts, providing optimized inference
+on Ascend hardware.
+
+## Directory Structure
+
+```text
+sgl-kernel-npu/
+├── csrc/                          # Ascend C operator C++ sources
+│   ├── CMakeLists.txt             # Build configuration
+│   ├── pytorch_extensions.cpp     # PyTorch op registration (core integration file)
+│   └── <op_name>/                 # One directory per operator
+│       ├── op_host/               # Host-side code (validation, tiling, launch)
+│       │   ├── <op_name>.cpp
+│       │   └── tiling/            # Optional: tiling data
+│       └── op_kernel/             # Device-side code (Ascend C kernel on AICore)
+│           └── <op_name>_kernel.cpp
+├── include/
+│   └── sgl_kenel_npu_ops.h        # C++ interface declarations
+├── python/
+│   └── sgl_kernel_npu/
+│       └── sgl_kernel_npu/
+│           ├── __init__.py        # Loads libsgl_kernel_npu.so
+│           ├── attention/         # Triton attention kernels
+│           ├── norm/              # Triton normalization kernels
+│           ├── activation/        # Triton activation kernels
+│           ├── fla/               # Triton linear attention kernels
+│           ├── mamba/             # Triton Mamba kernels
+│           ├── moe/               # Triton MoE kernels
+│           └── sample/            # Triton speculative decoding kernels
+├── tests/
+│   └── python/sgl_kernel_npu/     # One test file per operator
+├── build.sh                       # Build script
+└── CMakeLists.txt                 # Root CMake configuration
+```
+
+## Developing Ascend C Operators
+
+A complete Ascend C operator consists of two parts:
+
+- **Device part**: Kernel code running on the NPU AICore, responsible for actual
+  computation. Written using the Ascend C API.
+- **Host part**: Code running on the CPU, responsible for parameter validation,
+  data pre-processing, tiling, and kernel launch.
+
+We recommend starting with the
+[helloworld](https://github.com/sgl-project/sgl-kernel-npu/tree/main/csrc/helloworld)
+example, a simple operator that performs element-wise addition on two tensors.
+
+### Step 1: Create the operator directory and files
+
+Create a new operator directory under `csrc/`, following the `op_host/` +
+`op_kernel/` structure:
+
+```text
+csrc/<op_name>/
+├── op_host/
+│   └── <op_name>.cpp
+└── op_kernel/
+    └── <op_name>_kernel.cpp
+```
+
+### Step 2: Write the Device-side Kernel (op_kernel)
+
+Device-side code runs on AICore and follows the Ascend C programming model. The
+core structure is a class with `Init()` and `Process()` methods, plus an
+`extern "C"` entry function.
+
+Using helloworld as an example:
+
+```cpp
+// csrc/helloworld/op_kernel/kernel_helloworld.cpp
+#include "kernel_operator.h"
+
+constexpr int32_t BUFFER_NUM = 2;
+
+class KernalHelloworld {
+public:
+    __aicore__ inline KernalHelloworld() {}
+
+    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
+    {
+        // Compute workload for current block
+        this->blockLength = totalLength / AscendC::GetBlockNum();
+        this->tileNum = 8;
+        this->tileLength = this->blockLength / this->tileNum / BUFFER_NUM;
+
+        // Set global memory buffers
+        xGm.SetGlobalBuffer((__gm__ half *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
+        yGm.SetGlobalBuffer((__gm__ half *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
+        zGm.SetGlobalBuffer((__gm__ half *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
+
+        // Initialize pipeline queues
+        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half));
+        pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(half));
+        pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(half));
+    }
+
+    __aicore__ inline void Process()
+    {
+        int32_t loopCount = this->tileNum * BUFFER_NUM;
+        for (int32_t i = 0; i < loopCount; i++) {
+            CopyIn(i);    // Move data from Global Memory to Local Memory
+            Compute(i);   // Compute on Local Memory
+            CopyOut(i);   // Move results back to Global Memory
+        }
+    }
+
+private:
+    __aicore__ inline void CopyIn(int32_t progress) { /* data copy-in... */ }
+    __aicore__ inline void Compute(int32_t progress) { /* core computation... */ }
+    __aicore__ inline void CopyOut(int32_t progress) { /* data copy-out... */ }
+
+private:
+    AscendC::TPipe pipe;
+    AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
+    AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
+    AscendC::GlobalTensor<half> xGm, yGm, zGm;
+    uint32_t blockLength, tileNum, tileLength;
+};
+
+// Entry function: the compile tool auto-generates aclrtlaunch_<op_name>.h from this name
+extern "C" __global__ __aicore__ void helloworld(
+    GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
+{
+    KernalHelloworld op;
+    op.Init(x, y, z, totalLength);
+    op.Process();
+}
+```
+
+**Key points:**
+- Class methods must be marked with `__aicore__`, indicating they run on AICore.
+- Use `AscendC::TPipe` + `AscendC::TQue` to build a pipeline that overlaps data
+  movement and computation.
+- The entry function must be declared `extern "C" __global__ __aicore__`. The
+  compile tool generates a host-callable launch header
+  `aclrtlaunch_<func_name>.h` from the function name.
+- Simple operators (e.g., helloworld, cache_assign, lora) do not need extra
+  workspace memory. Complex operators (e.g., mla_preprocess, alloc_extend,
+  build_tree) require temporary workspace memory and are compiled separately in
+  `CMakeLists.txt`.
+
+For more in-depth Ascend C programming knowledge, refer to the
+[Ascend C Kernel Development Guide](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850alpha001/opdevg/Ascendcopdevg/atlas_ascendc_10_0001.html).
+
+### Step 3: Write the Host-side Code (op_host)
+
+Host-side code is responsible for passing PyTorch Tensors to the kernel and
+launching it. The key macro is `EXEC_KERNEL_CMD` (located in
+`csrc/utils/torch_helper.h`).
+
+```cpp
+// csrc/helloworld/op_host/helloworld.cpp
+#include "defines.h"                   // Provides HOST_API macro
+#include "torch_helper.h"              // Provides EXEC_KERNEL_CMD macro
+#include "aclrtlaunch_helloworld.h"    // Auto-generated by compile tool
+
+namespace sglang {
+namespace npu_kernel {
+
+HOST_API at::Tensor helloworld(const at::Tensor &x, const at::Tensor &y)
+{
+    // Create output tensor
+    at::Tensor z = at::empty_like(x);
+
+    // Define block count
+    uint32_t blockDim = 8;
+
+    // Compute total element count
+    uint32_t totalLength = 1;
+    for (uint32_t size : x.sizes()) {
+        totalLength *= size;
+    }
+
+    // Launch kernel via EXEC_KERNEL_CMD macro
+    EXEC_KERNEL_CMD(helloworld, blockDim, x, y, z, totalLength);
+    return z;
+}
+
+} // namespace npu_kernel
+} // namespace sglang
+```
+
+**Key points:**
+- The namespace must be `sglang::npu_kernel`.
+- Function signatures follow the pattern
+  `at::Tensor <op_name>(const at::Tensor &input, ...)`.
+- For operators with multiple outputs, use
+  `std::tuple<at::Tensor, at::Tensor, ...>` or non-const reference parameters.
+
+### Step 4: Declare the C++ Interface (include/sgl_kenel_npu_ops.h)
+
+Add the operator function declaration in `include/sgl_kenel_npu_ops.h`:
+
+```cpp
+// include/sgl_kenel_npu_ops.h
+namespace sglang {
+namespace npu_kernel {
+
+at::Tensor helloworld(const at::Tensor &x, const at::Tensor &y);
+
+} // namespace npu_kernel
+} // namespace sglang
+```
+
+### Step 5: Register the PyTorch Custom Operator (pytorch_extensions.cpp)
+
+Register the operator in `csrc/pytorch_extensions.cpp` in two steps: define the
+schema and bind the implementation.
+
+```cpp
+// csrc/pytorch_extensions.cpp
+namespace {
+
+// 1. Define operator schema (used by torch.compile, etc.)
+TORCH_LIBRARY_FRAGMENT(npu, m)
+{
+    m.def("helloworld(Tensor x, Tensor y) -> Tensor");
+    // ... other operator schemas ...
+}
+
+// 2. Bind implementation for the PrivateUse1 device (i.e., NPU)
+TORCH_LIBRARY_IMPL(npu, PrivateUse1, m)
+{
+    m.impl("helloworld", TORCH_FN(sglang::npu_kernel::helloworld));
+    // ... other operator implementations ...
+}
+
+} // namespace
+```
+
+**Schema conventions:**
+- The namespace is fixed to `npu`. In SGLang, operators are called via
+  `torch.ops.npu.<op_name>()`.
+- Output tensor parameters use the `Tensor(a!)` mutating annotation.
+- Optional parameters use the `Tensor?` annotation, with `c10::optional<T>`
+  handling in the impl.
+- For detailed schema syntax, see the
+  [PyTorch Schema Reference](https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/native/README.md#func).
+
+**Implementation binding rules:**
+- The device name is fixed to `PrivateUse1` (PyTorch NPU backend identifier).
+- Use the `TORCH_FN` macro to bind to the implementation function.
+- For complex operators with optional parameters, use lambda expressions to
+  unpack the arguments.
+
+### Step 6: Update the Build Configuration (csrc/CMakeLists.txt)
+
+Add the new operator's source files to `csrc/CMakeLists.txt`:
+
+**For operators not requiring workspace** (simple operators), add kernel source
+to `no_workspace_kernel`:
+
+```cmake
+ascendc_library(no_workspace_kernel STATIC
+    # ... existing kernel files ...
+    ${PROJECT_OP_SRC_BASE}/<op_name>/op_kernel/<op_name>_kernel.cpp
+)
+```
+
+**For operators requiring workspace** (complex operators), add kernel source to
+`workspace_kernel` with the `-DHAVE_WORKSPACE -DHAVE_TILING` compile flags:
+
+```cmake
+ascendc_library(workspace_kernel STATIC
+    # ... existing kernel files ...
+    ${PROJECT_OP_SRC_BASE}/<op_name>/op_kernel/<op_name>_kernel.cpp
+)
+```
+
+**Add host source files to `OP_SRCS`:**
+
+```cmake
+FILE(GLOB OP_SRCS
+    # ... existing host files ...
+    ${PROJECT_OP_SRC_BASE}/<op_name>/op_host/<op_name>.cpp
+)
+```
+
+### Step 7: Build
+
+Build following the steps in the
+[python/sgl_kernel_npu/README.md](https://github.com/sgl-project/sgl-kernel-npu/blob/main/python/sgl_kernel_npu/README.md):
+
+```bash
+cd sgl-kernel-npu
+
+# Build all modules
+bash build.sh
+
+# Install the sgl_kernel_npu wheel
+pip install output/sgl_kernel_npu*.whl
+```
+
+The compiled `libsgl_kernel_npu.so` is copied into
+`python/sgl_kernel_npu/sgl_kernel_npu/lib/` and loaded by the Python package.
+
+## Developing Triton Operators
+
+Triton operators are located under `python/sgl_kernel_npu/sgl_kernel_npu/`,
+organized by function category:
+
+```text
+python/sgl_kernel_npu/sgl_kernel_npu/
+├── attention/     # Attention (decode_attention, sinks_attention)
+├── norm/          # Normalization (rmsnorm, fused_qk_norm, l1_norm)
+├── activation/    # Activation (swiglu_oai, swiglu_quant)
+├── fla/           # Linear attention (chunk, cumsum, wy_fast)
+├── mamba/         # Mamba-related (causal_conv1d, state_update)
+├── moe/           # MoE-related (mul_add, zero_experts)
+└── sample/        # Speculative decoding (verify_tree_greedy)
+```
+
+**Development steps:**
+
+1. Create a new `.py` file in the appropriate category subdirectory.
+2. Write the kernel using the Triton language, using existing operators in the
+   same directory as templates.
+3. Export functions in the corresponding `__init__.py` if needed.
+4. Write tests under `tests/python/sgl_kernel_npu/`.
+
+**Note:** Many Triton operators are adapted from SGLang's GPU Triton kernels
+(e.g., comments in `fla/utils.py` note the original source). Pay special
+attention to differences between NPU and GPU when adapting.
+
+## Integrating Operators into SGLang
+
+### Ascend C Operator Integration
+
+After completing the [Steps 1-7](#developing-ascend-c-operators) above (writing
+the kernel, registering the torch op, building), and installing the
+`sgl-kernel-npu` wheel, call the operator in SGLang as follows:
+
+```python
+import sgl_kernel_npu  # Loading the library auto-triggers libsgl_kernel_npu.so loading
+
+# Call the operator
+result = torch.ops.npu.helloworld(x, y)
+```
+
+Real-world usage in SGLang (from `sglang/srt/speculative/eagle_utils.py`):
+
+```python
+torch.ops.npu.build_tree_kernel_efficient(
+    parent_list, selected_index, verified_seq_len, tree_mask,
+    positions, retrive_index, retrive_next_token,
+    retrive_next_sibling, topk, depth, draft_token_num, tree_mask_mode
+)
+```
+
+### Triton Operator Integration
+
+Import and call directly via Python:
+
+```python
+from sgl_kernel_npu.attention.decode_attention import decode_attention_fwd
+from sgl_kernel_npu.norm.rmsnorm_bias import rmsnorm_bias
+from sgl_kernel_npu.mamba.causal_conv1d import causal_conv1d_fwd
+
+# Direct function call
+output = decode_attention_fwd(q, k, v, ...)
+```
+
+Real-world usage in SGLang (from `sglang/srt/models/llama.py`):
+
+```python
+from sgl_kernel_npu.norm.split_qkv_rmsnorm_rope import split_qkv_rmsnorm_rope
+```
+
+### sgl-kernel-npu Wheel Update Process
+
+Since SGLang and sgl-kernel-npu are separate Python packages, dependency updates
+require a multi-PR workflow:
+
+1. **Submit sgl-kernel-npu PR**: Add/modify operators in the sgl-kernel-npu
+   repository, ensuring all tests pass.
+2. **Bump sgl-kernel-npu version**: Update the version number in sgl-kernel-npu.
+   Merging triggers an automatic PyPI release.
+3. **Reference the new version in SGLang**:
+   - Update the `sgl-kernel-npu` version requirement in SGLang's
+     `python/pyproject.toml`.
+   - Use the new operator in SGLang code.
+
+If not urgent, you can wait for a regular release (typically within one week).
+
+## Writing Unit Tests
+
+Each operator needs a corresponding unit test under
+`tests/python/sgl_kernel_npu/`, using Python's `unittest` framework.
+
+Test file naming convention: `test_<op_name>.py`
+
+```python
+# tests/python/sgl_kernel_npu/test_helloworld.py
+import unittest
+import torch
+import sgl_kernel_npu
+
+class TestHelloworld(unittest.TestCase):
+    def test_helloworld_basic(self):
+        x = torch.randn(1024, dtype=torch.bfloat16, device="npu")
+        y = torch.randn(1024, dtype=torch.bfloat16, device="npu")
+
+        z = torch.ops.npu.helloworld(x, y)
+        expected = x + y
+
+        torch.testing.assert_close(z, expected)
+
+if __name__ == "__main__":
+    unittest.main()
+```
+
+Run tests:
+
+```bash
+python tests/python/sgl_kernel_npu/test_helloworld.py
+```
+
+**Testing checklist:**
+- Cover typical input shapes (power-of-2 sizes and non-standard sizes).
+- Cover different data types (bf16 / fp16, etc.).
+- For operators with in-place behavior, verify correctness of output tensors.
+- Compare against PyTorch native computation to verify accuracy.
+
+## Code Style
+
+### Pre-commit Checks
+
+sgl-kernel-npu uses pre-commit for consistent code style:
+
+```bash
+pip3 install pre-commit
+cd sgl-kernel-npu
+pre-commit install
+pre-commit run --all-files
+```
+
+**Note:** If `pre-commit run --all-files` fails the first time, run it again to
+ensure all lint errors are auto-fixed. All code must pass checks before
+submitting a PR.
+
+### C++ Code Style
+
+- Use the C++17 standard.
+- Place all operator implementations under the `sglang::npu_kernel` namespace.
+- Follow existing code style; format using `.clang-format`.
+- Use `TORCH_CHECK` for error checking (not the standard GE `OP_ADD` macros).
+- Do not include unnecessary GE registration code (e.g., `OP_ADD()` macros).
+
+### Python Code Style
+
+- Follow PEP 8.
+- Use snake_case for file and function names.
+- When adapting from SGLang GPU code, note the original source in the file
+  header.
+
+### General Principles
+
+- Avoid code duplication: extract shared functions for any repeated code blocks
+  over 5 lines.
+- Minimize device synchronization: reduce CPU-NPU sync operations like
+  `tensor.item()` or `tensor.cpu()`.
+- Keep functions pure: avoid in-place argument modification.
+- Keep files concise: split files exceeding 2,000 lines.
+
+## Submitting a PR
+
+1. **Fork the repo**: Fork
+   [sgl-kernel-npu](https://github.com/sgl-project/sgl-kernel-npu) on GitHub,
+   then clone locally.
+2. **Create a branch**: Create a new branch from `main`, e.g.,
+   `feature/add-my-op`.
+3. **Develop and test**: Develop the operator and write tests following the
+   steps above. Ensure all tests pass.
+4. **Run pre-commit**: Ensure code formatting compliance.
+5. **Commit and push**:
+   ```bash
+   git add .
+   git commit -m "feat: add <op_name> operator"
+   git push origin feature/add-my-op
+   ```
+6. **Create a PR**: Open a Pull Request on GitHub from your branch to
+   `sgl-project/sgl-kernel-npu:main`.
+7. **Wait for CI and review**: CI checks include linting, compilation, and
+   operator tests. After passing, wait for maintainer review and merge.
+
+## References
+
+- [SGL-Kernel-NPU Official Repository](https://github.com/sgl-project/sgl-kernel-npu)
+- [SGL-Kernel-NPU Contribution Guide](https://github.com/sgl-project/sgl-kernel-npu/blob/main/docs/developer_guide/contribution_guide.md)
+- [Ascend C Kernel Development Guide](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850alpha001/opdevg/Ascendcopdevg/atlas_ascendc_10_0001.html)
+- [PyTorch Custom Ops Schema Reference](https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/native/README.md#func)
+- [helloworld Example Operator](https://github.com/sgl-project/sgl-kernel-npu/tree/main/csrc/helloworld)
+- [SGLang Contribution Guide](/docs/hardware-platforms/ascend-npus/ascend_contribution_guide)

docs_new/docs/hardware-platforms/ascend-npus/ascend_npu_profiling.mdx

@@ -0,0 +1,531 @@
+---
+title: "Ascend NPU Performance Profiling Guide"
+metatags:
+  description: "Ascend NPU performance profiling guide: use SGLang's built-in PyTorch Profiler for operator-level performance analysis of Ascend NPU inference services."
+---
+
+During inference serving, it is sometimes necessary to monitor the internal
+execution flow of the serving framework to identify performance issues. By
+collecting start/end timestamps of key flows, identifying critical functions or
+iterations, recording key events, and gathering relevant information, you can
+quickly locate performance bottlenecks.
+
+This guide walks you through the complete workflow of collecting performance
+data in an SGLang Ascend NPU inference service — from preparation, collection,
+and analysis to visualization — helping you get started with performance
+profiling quickly.
+
+For more profiling scenarios (e.g., Nsight Systems, PD disaggregation, etc.),
+see [SGLang Benchmark and Profiling](/docs/developer_guide/benchmark_and_profiling).
+
+## Ascend PyTorch Profiler
+
+SGLang has built-in PyTorch Profiler support. Through the Ascend `torch_npu`
+backend, you can directly collect NPU operator-level performance data. No
+additional packages are required — profiling start/stop is controlled via API
+requests.
+
+### 1. Environment Setup
+
+Launch an SGLang online service and set the `SGLANG_TORCH_PROFILER_DIR`
+environment variable to control where performance files are saved. Once the
+service starts, profiling is ready on standby.
+
+```shell Command
+# Set the performance data output directory
+export SGLANG_TORCH_PROFILER_DIR=./sglang_profile
+
+# Start SGLang server (use local model path or HuggingFace model id)
+sglang serve \
+  --model-path /path/to/your/model \
+  --attention-backend ascend \
+  --host 0.0.0.0 --port 30000 \
+  --tp-size 1 \
+  --max-running-requests 128
+```
+
+<Note>
+On Ascend NPU, SGLang uses `torch_npu._apply_patches()` to automatically
+redirect PyTorch Profiler's CUDA activity to NPU, so
+`activities: ["CPU", "GPU"]` actually captures NPU operator events.
+</Note>
+
+**Profiling-related environment variables:**
+
+<table>
+  <thead>
+    <tr>
+      <th>Variable</th>
+      <th>Description</th>
+      <th>Default</th>
+    </tr>
+  </thead>
+  <tbody>
+    <tr>
+      <td><code>SGLANG_TORCH_PROFILER_DIR</code></td>
+      <td>Trace file output directory</td>
+      <td><code>/tmp</code></td>
+    </tr>
+    <tr>
+      <td><code>SGLANG_PROFILE_WITH_STACK</code></td>
+      <td>Record Python call stack (True / False)</td>
+      <td><code>True</code></td>
+    </tr>
+    <tr>
+      <td><code>SGLANG_PROFILE_RECORD_SHAPES</code></td>
+      <td>Record operator input shapes (True / False)</td>
+      <td><code>True</code></td>
+    </tr>
+  </tbody>
+</table>
+
+### 2. Collection Methods
+
+SGLang provides four collection methods. The core differences are **whether you
+need to manually send `/start_profile` and `/stop_profile`**. All four methods
+produce identical results — choose the most convenient one.
+
+**Method comparison:**
+
+<table>
+  <thead>
+    <tr>
+      <th>Method</th>
+      <th>Manual start_profile</th>
+      <th>Manual stop_profile</th>
+      <th>Notes</th>
+    </tr>
+  </thead>
+  <tbody>
+    <tr>
+      <td>A: API manual start/stop</td>
+      <td>Yes</td>
+      <td>Yes</td>
+      <td>Maximum flexibility for precise control</td>
+    </tr>
+    <tr>
+      <td>B: API auto-stop</td>
+      <td>Yes</td>
+      <td>No</td>
+      <td>Set <code>num_steps</code>, auto-stops and generates output</td>
+    </tr>
+    <tr>
+      <td>C: bench_serving --profile</td>
+      <td>No</td>
+      <td>No</td>
+      <td>Benchmark + profiling in one command</td>
+    </tr>
+    <tr>
+      <td>D: sglang.profiler CLI</td>
+      <td>No</td>
+      <td>No</td>
+      <td>Standalone profiling CLI tool</td>
+    </tr>
+  </tbody>
+</table>
+
+#### Method A: API Manual Start/Stop
+
+Send `/start_profile` to start → send workload requests → send `/stop_profile`
+to stop. After stopping, the server automatically parses the data — **no need to
+manually call `analyse()`**.
+
+```bash Command
+# Step 1: Start profiling (no num_steps, requires manual stop)
+curl -X POST http://127.0.0.1:30000/start_profile \
+  -H "Content-Type: application/json" \
+  -d '{
+    "output_dir": "./sglang_profile",
+    "start_step": 1,
+    "activities": ["CPU", "GPU"]
+  }'
+
+# Step 2: Send workload requests (using curl as example)
+curl http://127.0.0.1:30000/generate \
+  -H "Content-Type: application/json" \
+  -d '{"text": "Hello", "sampling_params": {"max_new_tokens": 10}}'
+
+# Step 3: Stop profiling
+curl -X POST http://127.0.0.1:30000/stop_profile
+```
+
+<Note>
+`/stop_profile` returns `"Stop profiling. This will take some time."` — the
+server needs time to flush trace data to disk and parse it. Wait for the
+response to complete.
+This method takes a significant amount of time to parse
+profiling data;consider using **Method B** instead to avoid lengthy waits.
+</Note>
+
+#### Method B: API Auto-Stop
+
+Specify `num_steps` in the `/start_profile` request. Profiling stops
+automatically after N steps and generates output — **no need to manually send
+`/stop_profile`**.
+
+```bash Command
+# num_steps=10, wait 3 warmup steps, auto-stop after 10 steps
+curl -X POST http://127.0.0.1:30000/start_profile \
+  -H "Content-Type: application/json" \
+  -d '{
+    "output_dir": "./sglang_profile",
+    "start_step": 3,
+    "num_steps": 10,
+    "activities": ["CPU", "GPU"]
+  }'
+
+# Just send workload — no /stop_profile needed
+curl http://127.0.0.1:30000/generate \
+  -H "Content-Type: application/json" \
+  -d '{"text": "Hello", "sampling_params": {"max_new_tokens": 32}}'
+```
+
+#### Method C: bench_serving --profile
+
+Use SGLang's built-in `bench_serving` with the `--profile` flag.
+**Automatically handles `/start_profile` and `/stop_profile`** — no manual API
+calls needed.
+
+```bash Command
+# With --profile-steps: auto-stops after N steps and generates output
+python -m sglang.bench_serving \
+  --backend sglang \
+  --base-url http://127.0.0.1:30000 \
+  --model /path/to/your/model \
+  --tokenizer /path/to/your/model \
+  --dataset-name random \
+  --random-input-len 1024 \
+  --random-output-len 100 \
+  --num-prompts 10 \
+  --profile \
+  --profile-steps 10 \
+  --profile-output-dir ./sglang_profile
+
+# Without --profile-steps: /stop_profile sent automatically after benchmark
+python -m sglang.bench_serving \
+  --backend sglang \
+  --base-url http://127.0.0.1:30000 \
+  --model /path/to/your/model \
+  --tokenizer /path/to/your/model \
+  --dataset-name random \
+  --random-input-len 1024 \
+  --random-output-len 100 \
+  --num-prompts 10 \
+  --profile \
+  --profile-output-dir ./sglang_profile
+```
+
+<Note>
+`--profile-steps N` sends `"num_steps": N` to the server's `/start_profile`, so
+the server auto-stops and parses data after N steps — bench_serving skips
+sending `/stop_profile`.
+</Note>
+
+<Note>
+`bench_serving --profile` creates a timestamp subdirectory inside
+`--profile-output-dir` (e.g. `<output_dir>/<timestamp>/`). The output path is
+shown in the server log as `Profiling done. Traces are saved to: <path>`.
+</Note>
+
+**`bench_serving --profile` parameters:**
+
+<table>
+  <thead>
+    <tr><th>Parameter</th><th>Description</th></tr>
+  </thead>
+  <tbody>
+    <tr><td><code>--profile</code></td><td>Enable auto profiling start/stop</td></tr>
+    <tr><td><code>--profile-steps N</code></td><td>Auto-stop after N steps (skips /stop_profile)</td></tr>
+    <tr><td><code>--profile-output-dir</code></td><td>Trace output directory</td></tr>
+  </tbody>
+</table>
+
+#### Method D: sglang.profiler CLI
+
+Use the `sglang.profiler` CLI module, which automatically sends
+`/start_profile` and waits for completion. **Start `sglang.profiler` first,
+then send inference requests** (otherwise there are no steps to capture and the
+profiler will wait indefinitely).
+
+```bash Command
+# Terminal 1: Start sglang.profiler first (sends /start_profile, then waits for completion)
+python3 -m sglang.profiler \
+  --url http://127.0.0.1:30000 \
+  --output-dir ./my_profiles \
+  --num-steps 3 \
+  --cpu --gpu &
+```
+```bash Command
+# Terminal 2: Immediately send inference requests to provide steps for profiling
+curl http://127.0.0.1:30000/generate \
+  -H "Content-Type: application/json" \
+  -d '{"text": "Hello", "sampling_params": {"max_new_tokens": 32}}'
+```
+
+A simpler and more reliable approach is to use `bench_serving --profile`, which
+handles both steps automatically:
+
+```bash Command
+python3 -m sglang.bench_serving \
+  --backend sglang \
+  --base-url http://127.0.0.1:30000 \
+  --model /path/to/your/model \
+  --tokenizer /path/to/your/model \
+  --dataset-name random \
+  --random-input-len 128 \
+  --random-output-len 32 \
+  --num-prompts 10 \
+  --profile \
+  --profile-steps 3 \
+  --profile-output-dir ./my_profiles
+```
+
+<Note>
+`sglang.profiler` is essentially a CLI wrapper around the `/start_profile` API.
+Advanced options like `--profile-by-stage` are also supported. On Ascend NPU,
+trace flushing is asynchronous and may take a while — the CLI may occasionally
+block waiting for flush. If it times out, use Method B (API auto-stop) or
+Method C (bench_serving --profile) instead.
+</Note>
+
+**`sglang.profiler` CLI parameters:**
+
+<table>
+  <thead>
+    <tr><th>Parameter</th><th>Description</th></tr>
+  </thead>
+  <tbody>
+    <tr><td><code>--url</code></td><td>SGLang server address</td></tr>
+    <tr>
+      <td><code>--output-dir</code></td>
+      <td>Output directory (defaults to <code>SGLANG_TORCH_PROFILER_DIR</code>)</td>
+    </tr>
+    <tr><td><code>--num-steps</code></td><td>Number of steps to profile</td></tr>
+    <tr>
+      <td><code>--profile-by-stage</code></td>
+      <td>Profile prefill / decode stages separately</td>
+    </tr>
+    <tr><td><code>--profile-prefix</code></td><td>Trace filename prefix</td></tr>
+    <tr>
+      <td><code>--cpu</code> / <code>--gpu</code> / <code>--mem</code> / <code>--rpd</code></td>
+      <td>Activity types to collect</td>
+    </tr>
+  </tbody>
+</table>
+
+### 3. Full Parameter Reference
+
+All methods ultimately send a `/start_profile` request to the server. The full
+set of supported parameters:
+
+<table>
+  <thead>
+    <tr>
+      <th>Parameter</th>
+      <th>Description</th>
+      <th>Default</th>
+    </tr>
+  </thead>
+  <tbody>
+    <tr>
+      <td><code>output_dir</code></td>
+      <td>
+        Output directory. Falls back to
+        <code>SGLANG_TORCH_PROFILER_DIR</code> or <code>/tmp</code>
+      </td>
+      <td><code>/tmp</code></td>
+    </tr>
+    <tr>
+      <td><code>num_steps</code></td>
+      <td>
+        Number of steps. If set, profiling auto-stops — no /stop_profile needed
+      </td>
+      <td>None</td>
+    </tr>
+    <tr>
+      <td><code>start_step</code></td>
+      <td>
+        Step index to start profiling (inclusive), for skipping warmup
+      </td>
+      <td>0</td>
+    </tr>
+    <tr>
+      <td><code>activities</code></td>
+      <td>
+        Activity types: CPU, GPU, MEM, RPD. On Ascend NPU, primarily CPU and GPU
+      </td>
+      <td><code>["CPU", "GPU"]</code></td>
+    </tr>
+    <tr>
+      <td><code>profile_by_stage</code></td>
+      <td>Profile prefill and decode stages separately</td>
+      <td><code>false</code></td>
+    </tr>
+    <tr>
+      <td><code>with_stack</code></td>
+      <td>
+        Record Python call stack. Also controllable via
+        <code>SGLANG_PROFILE_WITH_STACK</code>
+      </td>
+      <td><code>true</code></td>
+    </tr>
+    <tr>
+      <td><code>record_shapes</code></td>
+      <td>
+        Record operator input shapes. Also controllable via
+        <code>SGLANG_PROFILE_RECORD_SHAPES</code>
+      </td>
+      <td><code>true</code></td>
+    </tr>
+    <tr>
+      <td><code>profile_prefix</code></td>
+      <td>Prefix for trace filenames</td>
+      <td>None</td>
+    </tr>
+    <tr>
+      <td><code>profile_stages</code></td>
+      <td>
+        Stages to profile, e.g. <code>["prefill", "decode"]</code>.
+        Requires <code>profile_by_stage</code>
+      </td>
+      <td>None</td>
+    </tr>
+  </tbody>
+</table>
+
+### 4. Finding Output Files
+
+**The server log explicitly indicates where traces are saved.** You can find
+them via:
+
+- **When profiling starts**: server log outputs
+  `Profiling starts. Traces will be saved to: <path> (with profile id: <id>)`
+```
+[2026-05-19 13:23:15] Profiling starts. Traces will be saved to: /tmp/1779196995.6948605 (with profile id: 1779196995.6979997)
+[2026-05-19 13:23:15] [WARNING] [350443] profiler.py: Invalid parameter export_type: None, reset it to text.
+[2026-05-19 13:23:15] [WARNING] [350443] profiler.py: Invalid parameter export_type: None, reset it to text.
+[2026-05-19 13:23:15] INFO:     127.0.0.1:40714 - "POST /start_profile HTTP/1.1" 200 OK
+```
+- **When profiling stops**: server log outputs
+  `Profiling done. Traces are saved to: <path>`
+```
+[2026-05-19 13:23:17] Stop profiling...
+[2026-05-19 13:23:17] [WARNING] [350443] profiler.py: Incorrect schedule: Stop profiler while current state is RECORD which may result in incomplete parsed data.
+[rank0]:[W519 13:23:17.084812760 compiler_depend.ts:3136] Warning: The indexFromRank 0is not equal indexFromCurDevice 4 , which might be normal if the number of devices on your collective communication server is inconsistent.Otherwise, you need to check if the current device is correct when calling the interface.If it's incorrect, it might have introduced an error. (function operator())
+[2026-05-19 13:23:17] [INFO] [352725] profiler.py: Start parsing profiling data: /tmp/1779196995.6948605/localhost.localdomain_350443_20260519132315700_ascend_pt
+[2026-05-19 13:23:22] [INFO] [352734] profiler.py: CANN profiling data parsed in a total time of 0:00:04.022310
+[2026-05-19 13:23:32] [INFO] [352725] profiler.py: All profiling data parsed in a total time of 0:00:14.305669
+[2026-05-19 13:23:32] Profiling done. Traces are saved to: /tmp/1779196995.6948605
+```
+- **CLI output**: `sglang.profiler` outputs `Dump profiling traces to <path>`
+```
+Dump profiling traces to /tmp/1779243331.3219
+Waiting for 10 steps and the trace to be flushed.... (profile_by_stage=False)
+```
+
+The directory structure is
+`<output_dir>/<hostname>_<pid>_<timestamp>_ascend_pt/`. When using Method C
+(`bench_serving --profile`), a timestamp subdirectory is added:
+`<output_dir>/<timestamp>/`. Always check the server log for the exact path:
+`Profiling done. Traces are saved to: <path>`.
+
+### 5. Viewing Results
+
+After profiling stops (either `/stop_profile` returns or `num_steps`
+auto-triggers), the server **automatically parses the raw data**. The
+`ASCEND_PROFILER_OUTPUT` directory directly contains the following visualization
+files — **no need to manually call `analyse()`**:
+
+<table>
+  <thead>
+    <tr><th>File</th><th>Description</th></tr>
+  </thead>
+  <tbody>
+    <tr>
+      <td><code>trace_view.json</code></td>
+      <td>
+        Chrome Tracing format. Open in
+        <a href="https://www.hiascend.com/document/detail/zh/mindstudio/81RC1/GUI_baseddevelopmenttool/msascendinsightug/Insight_userguide_0002.html">MindStudio Insight</a>
+      </td>
+    </tr>
+    <tr><td><code>analysis.db</code></td><td>Database-format performance data</td></tr>
+    <tr>
+      <td><code>ascend_pytorch_profiler_0.db</code></td>
+      <td>Database-format performance data</td>
+    </tr>
+    <tr><td><code>kernel_details.csv</code></td><td>Kernel-level data</td></tr>
+    <tr><td><code>operator_details.csv</code></td><td>Operator-level data</td></tr>
+    <tr><td><code>step_trace_time.csv</code></td><td>Step trace timing data</td></tr>
+  </tbody>
+</table>
+
+<Note>
+`trace_view.json` can also be opened using Chrome's built-in
+`chrome://tracing` or [Perfetto UI](https://ui.perfetto.dev/).
+</Note>
+
+<Note>
+If you need to merge distributed trace files in a multi-node deployment, set
+`"merge_profiles": true` in the `/start_profile` request. Note: on Ascend NPU,
+the merger has limited support for the `*_ascend_pt` format — check
+`trace_view.json` on each node individually. See
+[Benchmark and Profiling](/docs/developer_guide/benchmark_and_profiling#profiler-trace-merger-for-distributed-traces)
+for details.
+</Note>
+
+### 6. Re-parsing Raw Data (Optional)
+
+If you need to **re-parse existing data** with different parameters, or if
+profiling was interrupted and `ASCEND_PROFILER_OUTPUT` was not auto-generated,
+use `torch_npu`'s `analyse()` tool:
+
+```python
+from torch_npu.profiler.profiler import analyse
+analyse("./sglang_profile/<hostname>_*_ascend_pt/")
+```
+
+<Note>
+Normally **no need** to manually run `analyse()` — the server already parses
+data automatically. Only use this for re-parsing or handling interrupted data.
+</Note>
+
+## Best Practices
+
+### Common Notes
+
+- **Finding output**: Check the server log for
+  `Profiling starts. Traces will be saved to: <path>` and
+  `Profiling done. Traces are saved to: <path>`, or `sglang.profiler` output for
+  `Dump profiling traces to <path>`.
+- **Control trace file size**: Reduce the number of requests and output length
+  using `--num-prompts` and `--random-output-len` to avoid trace files too large
+  for browsers.
+- **Warmup iterations**: Set `start_step` to skip the first few warmup steps and
+  capture performance data under steady state.
+- **Profile step count**: Large values for `num_steps` or `--profile-steps` can
+  lead to lengthy profiling data parsing times. Reduce these values
+  appropriately when you only need a quick overview.
+- **CUDA Graph impact**: To see the full Python call stack → operator mapping in
+  traces, add `--disable-cuda-graph` when starting the server. Note that this
+  reduces decode performance — only use during profiling. To analyze CUDA Graph
+  capture specifically, use `--enable-profile-cuda-graph` — traces are saved to
+  `SGLANG_TORCH_PROFILER_DIR/graph_capture_profile/`.
+- **Multi-node deployment**: In multi-node environments, performance data is
+  distributed across nodes. On Ascend NPU, the `merge_profiles` feature has
+  limited support — check `*_ascend_pt/ASCEND_PROFILER_OUTPUT/trace_view.json`
+  on each node individually. In PD disaggregation mode, prefill and decode
+  workers must be profiled separately — see
+  [Profile In PD Disaggregation Mode](/docs/developer_guide/benchmark_and_profiling#profile-in-pd-disaggregation-mode).
+
+## See Also
+
+- [SGLang Benchmark and Profiling](/docs/developer_guide/benchmark_and_profiling)
+  — General SGLang profiling guide
+- [Ascend NPU Quickstart](/docs/hardware-platforms/ascend-npus/ascend_npu_quick_start)
+  — Ascend NPU environment setup
+- [Ascend NPU Optimization](/docs/hardware-platforms/ascend-npus/ascend_npu_optimization)
+  — Ascend NPU optimization parameters
+- [Ascend NPU Performance Testing](/docs/hardware-platforms/ascend-npus/ascend_npu_performance_testing)
+  — Ascend NPU performance benchmarking
+- [Ascend NPU Environment Variables](/docs/hardware-platforms/ascend-npus/ascend_npu_environment_variables)
+  — Environment variable reference