From 1bd5c0bd5e44bfb5ea426925131c1013d63c26c7 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Wed, 29 Jul 2026 11:29:05 +0800 Subject: [PATCH 1/7] feat: add Ascend950 direct SDMA backend --- AGENTS.md | 2 +- README.md | 20 +- docs/BUILD_VERIFICATION.md | 31 +- docs/SDMA_TRANSPORT.md | 247 +++--- docs/diagrams/architecture-overview.drawio | 4 +- .../diagrams/architecture-overview.drawio.svg | 2 +- ...026-07-28-a5-direct-sdma-implementation.md | 373 +++++++++ .../specs/2026-07-28-a5-direct-sdma-design.md | 222 ++++++ scripts/README.md | 15 + scripts/ci/control/build_blue.sh | 14 +- scripts/ci/provision/cann.sh | 4 +- scripts/ci/provision/common.sh | 47 +- scripts/ci/provision/verify.sh | 4 +- src/comm/CMakeLists.txt | 4 + src/comm/comm_wrap.cpp | 5 +- src/comm/sdma/tilexr_sdma_a5_backend.cpp | 710 ++++++++++++++++++ src/comm/sdma/tilexr_sdma_a5_backend.h | 138 ++++ src/comm/sdma/tilexr_sdma_transport.cpp | 74 +- src/comm/sdma/tilexr_sdma_transport.h | 32 +- src/comm/tilexr_comm.cpp | 37 +- src/comm/tilexr_comm.h | 3 +- src/include/tilexr_sdma.h | 36 +- src/include/tilexr_sdma_a5.h | 206 +++++ src/include/tilexr_sdma_a5_types.h | 200 +++++ tests/ci/test_control_contract.py | 23 +- tests/sdma/CMakeLists.txt | 49 +- tests/sdma/build.sh | 5 +- tests/sdma/demo/run_tilexr_sdma_demo.sh | 16 +- tests/sdma/demo/tilexr_sdma_demo.cpp | 137 ++-- tests/sdma/demo/tilexr_sdma_demo_kernel.cpp | 77 +- tests/sdma/run_tests.sh | 1 + .../unit/test_tilexr_sdma_a5_validation.cpp | 106 +++ tests/sdma/unit/test_tilexr_sdma_metadata.cpp | 63 ++ .../unit/test_tilexr_sdma_source_guard.cpp | 61 ++ .../test_tilexr_sdma_transport_disabled.cpp | 32 +- 35 files changed, 2752 insertions(+), 248 deletions(-) create mode 100644 docs/plans/2026-07-28-a5-direct-sdma-implementation.md create mode 100644 docs/specs/2026-07-28-a5-direct-sdma-design.md create mode 100644 src/comm/sdma/tilexr_sdma_a5_backend.cpp create mode 100644 src/comm/sdma/tilexr_sdma_a5_backend.h create mode 100644 src/include/tilexr_sdma_a5.h create mode 100644 src/include/tilexr_sdma_a5_types.h create mode 100644 tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp diff --git a/AGENTS.md b/AGENTS.md index 100266d3..699500fc 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -6,7 +6,7 @@ Ask the user to install missing `ascendc-development` or `superpowers-neo` skill TileXR is a C++14 communication runtime for Huawei Ascend NPUs. It provides tile-level synchronization, optional collectives and EP libraries, IPC/MTE communication, registered-memory UDMA, and opt-in local SDMA. -- Current build target: CANN 9.1.0 and NPU driver 25.5.0 or later. +- Current build target: CANN 9.1.0 and NPU driver 25.1.rc1 or later. - Core runtime targets Ascend 910B and 910A5. UDMA data-plane validation requires A5 / Ascend950 / 950 hardware. ## Key Paths diff --git a/README.md b/README.md index 24b4a743..3bdfa02d 100644 --- a/README.md +++ b/README.md @@ -27,13 +27,13 @@ Instead of stalling every rank at coarse barriers, TileXR splits a phase into ti - **Standalone EP dispatch/combine MVP**: `libtilexr-ep.so`, `libtilexr_ep_dispatch_kernel.so`, and `libtilexr_ep_combine_kernel.so` provide TileXR-native MoE EP dispatch/combine routes under `src/ep`, independent from HCCL window helpers, `ops-transformer`, and shmem. Same-node paths use IPC peer-memory windows; cross-node dispatch/combine use TileXR-registered UDMA workspaces. - **Tile-level synchronization**: device-side flag regions and magic values support reusable fine-grained synchronization rounds. - **Registered-memory UDMA path**: host code registers ordinary `aclrtMalloc` device memory with `TileXRUDMARegister`; device kernels use `tilexr_udma.h` wrappers for put/get/signal. -- **On-card SDMA transport**: an opt-in (`TILEXR_ENABLE_SDMA=1`) local GM-to-GM copy path. Host code queries it with `TileXRSDMAAvailable` / `TileXRGetSDMAWorkspaceDev`; device kernels use `tilexr_sdma.h` (`SDMACopyNbi`, `SDMAWait`). Separate from UDMA: SDMA is local to one device, UDMA targets registered remote memory. +- **On-card SDMA transport**: an opt-in (`TILEXR_ENABLE_SDMA=1`) local GM-to-GM copy path. A2/A3 use the PTO backend; A5 / Ascend950 uses TileXR's direct STARS SQ backend initialized through CANN's built-in `ShmemSdmaStarsQuery`. Host code queries it with `TileXRSDMAAvailable` / `TileXRGetSDMAWorkspaceDev`; device kernels use `tilexr_sdma.h` (`SDMACopyNbi`, `SDMAWait`). - **Operator simulator**: `op-simulator/` supports functional/performance simulation for selected AICore kernels without physical hardware. ## System Requirements - **User**: root access or membership in the Ascend driver user group is typically required for CANN runfile installation and NPU device operations -- **NPU driver**: 25.5.0 or later, check with `npu-smi info` +- **NPU driver**: 25.1.rc1 or later, check with `npu-smi info` - **CANN**: current build scripts and CMake are aligned to CANN 9.1.0 - **Core supported chips**: Ascend 910B, 910A5 - **UDMA runtime validation target**: A5 / Ascend950 / 950 only @@ -147,7 +147,7 @@ TileXR/ |-- src/ | |-- comm/ # Core communication runtime | | |-- udma/ # TileXR-owned HCCP/RA UDMA transport -| | `-- sdma/ # On-card PTO SDMA local copy transport +| | `-- sdma/ # On-card PTO and A5 direct SDMA backends | |-- collectives/ # Optional TileXR collectives library | |-- ep/ # Standalone TileXR EP dispatch MVP | `-- include/ # Public C/C++ and device headers @@ -246,12 +246,13 @@ If UDMA is unavailable, communicator initialization continues without setting `E SDMA is a first-class local on-card GM-to-GM copy path, separate from UDMA. It is disabled by default and enabled with `TILEXR_ENABLE_SDMA=1`. -- `TileXRComm::InitSDMA()` owns a `TileXRSDMATransport` beside the UDMA transport. When enabled, it creates a PTO `pto::comm::sdma::SdmaWorkspaceManager`, stores its device workspace address in `CommArgs::sdmaWorkspacePtr`, and sets `ExtraFlag::SDMA`. +- `TileXRComm::InitSDMA()` owns a `TileXRSDMATransport` beside the UDMA transport. Runtime SoC selection keeps PTO `SdmaWorkspaceManager` on A2/A3 and selects TileXR's 48-channel direct backend on A5 / Ascend950. - Host queries: `TileXRSDMAAvailable(comm, &available)` and `TileXRGetSDMAWorkspaceDev(comm, &workspace)`. The workspace pointer is owned by `TileXRComm` and must not be freed. - Device API: `src/include/tilexr_sdma.h` provides `TileXR::SDMACopyNbi` and `TileXR::SDMAWait`, accepting raw same-device GM pointers. It does not register memory or validate buffer ownership. -- PTO SDMA header differences across CANN 9.0.0 / 9.1.0 are isolated in `src/include/tilexr_sdma_compat.h`. +- A5 initialization calls CANN's built-in `ShmemSdmaStarsQuery`, maps RTSQ doorbells on Host, and lets AIV write the two-SQE data/completion sequence. TileXR does not ship a custom AICPU kernel or OPP package for SDMA. +- PTO-specific compatibility remains isolated in `src/include/tilexr_sdma_compat.h`; A5 uses the installed `tilexr_sdma_a5.h` and `tilexr_sdma_a5_types.h` headers. -Enabled initialization is best-effort: if PTO SDMA headers or runtime resources are unavailable, communicator initialization continues without setting `ExtraFlag::SDMA`, and `SDMACopyNbi` returns event handle `0` while `SDMAWait` reports completion. See [docs/SDMA_TRANSPORT.md](docs/SDMA_TRANSPORT.md) for the full transport guide. +Enabled initialization is best-effort: if the selected backend or runtime resources are unavailable, communicator initialization continues without setting `ExtraFlag::SDMA`, and `SDMACopyNbi` returns event handle `0` while `SDMAWait` reports completion. See [docs/SDMA_TRANSPORT.md](docs/SDMA_TRANSPORT.md) for the full transport guide. ## Dependencies @@ -296,7 +297,8 @@ See: Build and run the SDMA unit tests against a selected CANN install, then run the data-plane demo on a device: ```bash -bash tests/sdma/build.sh /path/to/cann +bash tests/sdma/build.sh /path/to/cann # Ascend910B default +bash tests/sdma/build.sh /path/to/cann Ascend950 # A5 / Ascend950 bash tests/sdma/run_tests.sh /path/to/cann bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 0 64 4096 1048576 ``` @@ -304,10 +306,10 @@ bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 0 64 4096 1048576 Expected demo success line: ```text -PASS TileXR SDMA copied bytes correctly +PASS TileXR SDMA copied bytes on block(s), channels .., iterations ``` -The unit tests are hardware-free; the demo requires a usable driver HAL/device runtime and resolves `libascend_hal.so` from `/usr/local/Ascend/driver/lib64/driver`. See [docs/SDMA_TRANSPORT.md](docs/SDMA_TRANSPORT.md) for enablement, the host/device API, CANN 9.0.0 / 9.1.0 acceptance steps, and current validation status. +The unit tests are hardware-free; the demo requires a usable driver HAL/device runtime and resolves `libascend_hal.so` from `/usr/local/Ascend/driver/lib64/driver`. A5 kernels are built with `-O2`, need no custom OPP, and do not link PTO's `libnnopbase.so`; the shared Host library retains PTO support for A2/A3. See [docs/SDMA_TRANSPORT.md](docs/SDMA_TRANSPORT.md) for enablement, APIs, backend details, and the hardware acceptance matrix. ## Collectives Validation diff --git a/docs/BUILD_VERIFICATION.md b/docs/BUILD_VERIFICATION.md index d5d79cd8..c5abd522 100644 --- a/docs/BUILD_VERIFICATION.md +++ b/docs/BUILD_VERIFICATION.md @@ -1,6 +1,6 @@ # TileXR Build Verification -**Updated:** 2026-05-29 +**Updated:** 2026-07-28 This checklist reflects the current TileXR codebase. The core runtime builds `libtile-comm.so` without a compile-time or link-time shmem dependency. @@ -16,7 +16,7 @@ Expected: - CANN 9.1.0 environment is visible through `ASCEND_HOME_PATH`. - `scripts/common_env.sh` detects architecture and SOC information. -- NPU driver version is 25.5.0 or later. +- NPU driver version is 25.1.rc1 or later. ## Build Core Runtime @@ -37,12 +37,33 @@ Check dynamic dependencies: ```bash ldd install/lib/libtile-comm.so | grep -E "ascendcl|runtime|ascend_hal|profapi" ldd install/lib/libtile-comm.so | grep -i shmem || true +readelf -d install/lib/libtile-comm.so | grep -E "RPATH|RUNPATH" || true ``` Expected: - CANN runtime libraries are resolved. -- The shmem grep prints nothing for the current TileXR UDMA implementation. +- The shmem grep prints nothing; A5 SDMA dynamically invokes the CANN built-in + query and does not add a shmem or custom OPP link dependency. +- Any RPATH/RUNPATH output does not contain a CANN `devlib` directory. + +## Build And Run SDMA Checks + +```bash +cd /path/to/TileXR +bash tests/sdma/build.sh "$ASCEND_HOME_PATH" Ascend950 +bash tests/sdma/run_tests.sh "$ASCEND_HOME_PATH" +``` + +Use the default `Ascend910B` target to compile-check the preserved PTO path: + +```bash +bash tests/sdma/build.sh "$ASCEND_HOME_PATH" +``` + +On A5 / Ascend950 hardware, run the direct data-plane matrix documented in +[SDMA_TRANSPORT.md](SDMA_TRANSPORT.md). The A5 kernel must retain the repository's +`-O2` compile option, and no custom OPP environment setting is required. ## Build UDMA Tests @@ -119,7 +140,7 @@ The final grep should print nothing for a clean run. | Symptom | Likely Cause | Action | | --- | --- | --- | | Missing CANN headers | `common_env.sh` not sourced or CANN path mismatch | Source the environment and confirm CANN 9.1.0 layout | -| Cannot find `ascend_hal` | `devlib` path missing | Use the current top-level CMake configuration | +| Cannot find `ascend_hal` | Driver HAL path missing | Add `/usr/local/Ascend/driver/lib64/driver` to `LD_LIBRARY_PATH`; do not use CANN `devlib` | | Demo target skipped | `bisheng` unavailable | Install/compiler configure `bisheng`, or run host-only tests | | UDMA disabled in demo | Unsupported hardware or HCCP/RA runtime unavailable | Use A5 / Ascend950 / 950 and check CANN driver/runtime libraries | | shmem appears in `ldd libtile-comm.so` | Unexpected dependency regression | Inspect `src/comm/CMakeLists.txt` and source includes | @@ -139,6 +160,8 @@ ldd shmem check: UDMA host tests: UDMA all-gather demo: UDMA put-signal demo: +SDMA unit/build checks: +SDMA A5 data-plane matrix: Log directory: Errors or warnings: diff --git a/docs/SDMA_TRANSPORT.md b/docs/SDMA_TRANSPORT.md index d0e6343d..6bb86e10 100644 --- a/docs/SDMA_TRANSPORT.md +++ b/docs/SDMA_TRANSPORT.md @@ -1,21 +1,30 @@ # TileXR SDMA Transport -TileXR SDMA transport provides a first-class local on-card GM-to-GM copy path. -It is separate from UDMA: SDMA is local to one device, while UDMA targets -registered remote memory on supported A5/Ascend950 systems. +TileXR SDMA is an opt-in, same-device GM-to-GM copy transport. It is separate +from UDMA: SDMA moves local device memory, while UDMA accesses registered remote +memory on supported A5 / Ascend950 systems. -## Enablement +TileXR supports two runtime-selected SDMA backends: -SDMA is disabled by default. Enable it explicitly: +- Ascend 910A/910B-class A2/A3 devices use the existing PTO SDMA backend. +- Ascend950/A5 devices use TileXR's direct STARS SQ submission backend. + +Both backends preserve the same Host and AIV APIs. SDMA remains best-effort: a +backend initialization failure leaves the communicator usable and reports SDMA +as unavailable. + +## Requirements And Enablement + +The project baseline is CANN 9.1.0 with NPU driver `25.1.rc1` or later. Enable +SDMA before communicator creation: ```bash export TILEXR_ENABLE_SDMA=1 ``` -When disabled, `TileXRComm` initialization does not create STARS streams and -`TileXRSDMAAvailable` reports `false`. Enabled initialization is best-effort: -if PTO SDMA headers or runtime resources are unavailable, normal communicator -initialization continues without setting `ExtraFlag::SDMA`. +When disabled, `TileXRComm` does not create SDMA resources, +`TileXRSDMAAvailable` returns `false`, and `CommArgs::extraFlag` does not contain +`ExtraFlag::SDMA`. ## Host API @@ -26,146 +35,166 @@ TileXRSDMAAvailable(comm, &available); TileXRGetSDMAWorkspaceDev(comm, &workspace); ``` -The workspace pointer is owned by `TileXRComm`; callers must not free it. A -successful query with `available=false` and `workspace=nullptr` means SDMA was +The workspace is owned by `TileXRComm`; callers must not free it. A successful +query with `available == false` and `workspace == nullptr` means SDMA was disabled or unavailable for that communicator. +Callers must synchronize all streams that can execute TileXR device APIs before +`TileXRCommDestroy`. Destroy does not cancel an in-flight AIV kernel or SDMA +event; it releases the communicator-owned workspace, STARS streams, and RTSQ +mappings. + ## Device API ```cpp #include "tilexr_sdma.h" -uint64_t event = TileXR::SDMACopyNbi(args, dst, src, bytes, 0); -bool ok = TileXR::SDMAWait(args, event, 0); +uint64_t event = TileXR::SDMACopyNbi(args, dst, src, bytes, channel); +bool ok = TileXR::SDMAWait(args, event, channel); ``` -The API accepts raw same-device GM pointers. It does not register memory and -does not validate whether pointers belong to TileXR communication buffers. +The API accepts same-device GM pointers. It does not register memory or validate +buffer ownership. Event 0 is the unavailable or invalid no-op result, and +`SDMAWait(args, 0, channel)` succeeds immediately. -`SDMACopyNbi` returns event handle `0` when SDMA is disabled or arguments are -invalid. `SDMAWait(args, 0, channelGroup)` returns `true`, matching the no-op -completion path. The default channel group is -`TILEXR_SDMA_AUTO_CHANNEL_GROUP`, which resolves to the current AI Core block -index; the demo passes channel group `0` explicitly. +`TILEXR_SDMA_AUTO_CHANNEL_GROUP` resolves to the current AIV block index. A5 +provides 48 channels numbered 0 through 47. Each channel permits one outstanding +event; different channels can progress concurrently. -## Implementation Notes +## A2/A3 PTO Backend -`TileXRComm::InitSDMA()` owns a `TileXRSDMATransport` beside the existing UDMA -transport. When `TILEXR_ENABLE_SDMA=1`, the transport creates a PTO -`pto::comm::sdma::SdmaWorkspaceManager`, obtains its device workspace address, -stores that address in `CommArgs::sdmaWorkspacePtr`, and sets -`ExtraFlag::SDMA`. +On A2/A3, `TileXRSDMATransport` owns a PTO +`pto::comm::sdma::SdmaWorkspaceManager` and publishes its device workspace in +`CommArgs::sdmaWorkspacePtr`. Device wrappers create a PTO `SdmaSession`, post +`__sdma_put_async`, and wait through the PTO event helper. Compatibility details +are isolated in `tilexr_sdma_compat.h`. -Device wrappers build a PTO `SdmaSession` from `args->sdmaWorkspacePtr`, post -`__sdma_put_async`, and wait with PTO's SDMA event wait helper. TileXR uses a -256-byte scratch tile and current defaults of one queue, -1 MiB block bytes, and communication block offset `0`. +## A5 Direct Backend -## CANN Compatibility +On Ascend950/A5, TileXR owns the Host initializer and AIV submission path. It +does not ship a TileXR AICPU binary or operator package. -The implementation is validated with the PTO SDMA interfaces in CANN 9.0.0 and -CANN 9.1.0. TileXR isolates PTO SDMA header differences in -`tilexr_sdma_compat.h`. +During communicator initialization, the Host backend: -Runtime must load `libascend_hal.so` from the driver path, typically: +1. Dynamically resolves CANN runtime and opapi entry points from `libruntime.so` + and `libopapi.so`. +2. Creates 48 device-only STARS streams. +3. Calls CANN's built-in `ShmemSdmaStarsQuery` system operator to obtain SQ/CQ + metadata. Driver `25.1.rc1` may return the known status `507018` after filling + all fields except the unsupported SQ register-base field; TileXR accepts only + that exact validated partial result and isolates each follow-up query in a + disposable context. +4. Cross-checks SQ state with the Host HAL and maps each RTSQ doorbell through + `halResAddrMap`. +5. Uploads a versioned TileXR-owned 48-channel workspace and publishes it through + `CommArgs::sdmaWorkspacePtr`. -```text -/usr/local/Ascend/driver/lib64/driver/libascend_hal.so -``` +The AIV backend atomically claims a channel, writes a data-copy SQE followed by a +64-byte completion-copy SQE, cleans the touched cache lines, orders the stores, +and updates the mapped RTSQ tail. `SDMAWait` validates the event generation, +polls the completion record with a fixed upper bound, and releases the channel +only after completion. A timed-out channel remains busy so later submissions +cannot reuse an uncertain queue. -Do not put `${ASCEND_HOME_PATH}/${ARCH}-linux/devlib` into runtime RPATH. The -devlib HAL can cause `aclInit` failures such as `500000` with `init soc version -failed`. +No TileXR target builds or installs a custom OPP, and this feature needs no +custom OPP environment setting. `reference/` remains comparison-only and is not +an include, source, or link dependency. -## Build And Test +## Build -Build tests against a selected CANN install: +Build the core and SDMA tests against CANN 9.1.0: ```bash -bash tests/sdma/build.sh /path/to/cann +bash tests/sdma/build.sh /path/to/cann # Ascend910B demo target +bash tests/sdma/build.sh /path/to/cann Ascend950 # A5 demo target bash tests/sdma/run_tests.sh /path/to/cann ``` -Run data-plane demo: +The supported demo targets are `Ascend910B` and `Ascend950`. The A5 device kernel +must be compiled with optimization enabled; the repository build uses `-O2`. +Repeated-channel submissions are not reliable with the Bisheng default +low-optimization code generation used in the investigated CANN 9.1.0 setup. + +The A5 device kernel does not include PTO headers and does not link +`libnnopbase.so`. The shared Host library retains its PTO dependency so the same +installation can continue to initialize the A2/A3 backend. + +The demo bounds Host stream synchronization at 60 seconds. The A5 busy-channel +check is reported only on A5; the PTO build reports that check as skipped. + +## Hardware Validation + +The runner's second argument is the physical device selected through +`ASCEND_RT_VISIBLE_DEVICES`. On the named A5 server, pass `5`, not logical device +`0`: ```bash -bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 0 64 4096 1048576 +TILEXR_SDMA_DEMO_CHANNEL=0 \ +TILEXR_SDMA_DEMO_ITERATIONS=3 \ +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 64 4096 1048576 + +TILEXR_SDMA_DEMO_CHANNEL=47 \ +TILEXR_SDMA_DEMO_ITERATIONS=3 \ +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 64 4096 1048576 + +TILEXR_SDMA_DEMO_CHANNEL=0 \ +TILEXR_SDMA_DEMO_BLOCKS=4 \ +TILEXR_SDMA_DEMO_ITERATIONS=3 \ +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 4096 + +TILEXR_SDMA_DEMO_CHANNEL=1 \ +TILEXR_SDMA_DEMO_ITERATIONS=3 \ +TILEXR_SDMA_DEMO_REPEATS=3 \ +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 4096 ``` -Expected demo success line: +Every run must report a nonzero event, a successful wait, the expected +generation, and a full byte-for-byte destination comparison. A skipped or +unavailable backend is not an A5 data-plane pass. -```text -PASS TileXR SDMA copied bytes correctly -``` +## Validation Record -## Acceptance +Production A5 acceptance was exercised on: -For release validation, run the unit tests and demo against both CANN versions: +- hardware: `Ascend950PR_9589` / Ascend950PR; +- CANN: 9.1.0 at `/home/pkg/b061/cann-9.1.T560`; +- driver: `25.1.rc1.b188`; +- physical device: 5; +- initialization: all 48 distinct STARS channels live simultaneously; +- channel 0 and channel 47: 64 B, 4 KiB, and 1 MiB, three iterations each; +- channels 0 through 3: four concurrent AIV blocks at 4 KiB, three iterations; +- channel 1: three init/use/destroy repetitions at 4 KiB, three iterations each; +- busy-channel behavior: a second outstanding submission returned event 0; +- validation: every successful run compared the full source and destination. -```bash -bash tests/sdma/build.sh "${TILEXR_CANN_90_HOME}" -bash tests/sdma/run_tests.sh "${TILEXR_CANN_90_HOME}" -bash tests/sdma/demo/run_tilexr_sdma_demo.sh "${TILEXR_CANN_90_HOME}" 0 64 4096 1048576 +The A2/A3 PTO path was previously exercised on Ascend 910B3 with CANN 9.0.0 and +9.1.0 and driver `25.5.0`. That driver value records the tested machine; it is +not TileXR's minimum supported driver. -bash tests/sdma/build.sh "${TILEXR_CANN_91_HOME}" -bash tests/sdma/run_tests.sh "${TILEXR_CANN_91_HOME}" -bash tests/sdma/demo/run_tilexr_sdma_demo.sh "${TILEXR_CANN_91_HOME}" 0 64 4096 1048576 -``` +## Runtime Dependency Checks -## Current Validation Status - -Local validation in this branch on the development host: - -- CANN 9.1.0 build passed with `TILEXR_HAVE_PTO_SDMA: ON`. -- CANN 9.1.0 SDMA unit tests passed. -- Local demo binary and kernel built. -- Local demo runtime is blocked because this environment has no usable driver - HAL/device runtime. With devlib HAL stripped, the runner reports - `libascend_hal.so not found`; if devlib HAL is allowed, the symptom can become - `aclInit ret=500000`. - -Hardware validation on `blue`: - -- Verification directory: - `/home//tilexr_sdma_verify_20260530/TileXR`. -- Hardware: 8 x Ascend 910B3, driver `25.5.0`, device 0 used for demo. -- CANN 9.0.0 path: `/home/gsn3/Ascend/cann-9.0.0`. -- CANN 9.1.0 path: - `/home//tilexr_sdma_verify_20260530/TileXR/env/cann/cann-9.1.0`. -- Both CANN versions configured with `TILEXR_HAVE_PTO_SDMA: ON` and linked - `libnnopbase.so` from `${ARCH}-linux/lib64` while resolving - `libascend_hal.so` from `/usr/local/Ascend/driver/lib64/driver`. -- `bash tests/sdma/run_tests.sh ` passed for both CANN 9.0.0 and 9.1.0. -- `bash tests/sdma/demo/run_tilexr_sdma_demo.sh 0 64 4096 1048576` - passed for both versions. Logs show STARS stream creation, STARS query - completion, `ExtraFlag::SDMA`, non-null SDMA workspace, and: +Runtime must resolve the real driver HAL, typically: ```text -PASS TileXR SDMA copied 64 bytes correctly -PASS TileXR SDMA copied 4096 bytes correctly -PASS TileXR SDMA copied 1048576 bytes correctly +/usr/local/Ascend/driver/lib64/driver/libascend_hal.so ``` -Remaining validation: - -- Broader parameter/performance matrix below; current demo covers one device, - one stream, channel group `0`, and sizes 64 B, 4 KiB, and 1 MiB. - -## Deferred Stress And Performance Scope +Do not put `${ASCEND_HOME_PATH}/${ARCH}-linux/devlib` in runtime RPATH or +RUNPATH. Validate the installed artifacts with: -Deferred validation scope includes multi-block channel-group assignment, -multi-stream concurrency, long-loop stability, parameter matrix tests for queue -settings, and MTE/SDMA bandwidth and latency comparisons. +```bash +ldd install/lib/libtile-comm.so | grep libascend_hal +readelf -d install/lib/libtile-comm.so | grep -E 'RPATH|RUNPATH' || true +readelf -d tests/sdma/install/bin/tilexr_sdma_demo | grep -E 'RPATH|RUNPATH' || true +``` -## TODO: C And Performance Parameter Testing +The HAL path must not contain `devlib`. No artifact may depend on a TileXR custom +OPP; the A5 device kernel must not have a `libnnopbase.so` dependency. -Future validation should cover: +## Failure Semantics -- C-facing API and parameter shape, if SDMA is exposed beyond the current C++ - device wrapper and host query API. -- `queue_num`, `block_bytes`, and `channelGroup` matrix coverage. -- Multiple transfer sizes, stream counts, loop counts, and concurrency levels. -- Runtime perf counters and DFX logs sufficient to distinguish SDMA issue, - wait, and copy completion costs. -- Bandwidth and latency comparisons against MTE `DataCopy`. -- Two-version CANN acceptance for CANN 9.0.0 and CANN 9.1.0. +Any missing symbol, malformed query result, context-health failure, HAL +cross-check mismatch, RTSQ mapping failure, allocation failure, or upload failure +disables SDMA for that communicator. Cleanup releases only TileXR-owned streams, +mappings, and workspace memory. It does not reset a device or destroy an +application-owned context or stream. diff --git a/docs/diagrams/architecture-overview.drawio b/docs/diagrams/architecture-overview.drawio index 4e73d891..01d9fd12 100644 --- a/docs/diagrams/architecture-overview.drawio +++ b/docs/diagrams/architecture-overview.drawio @@ -59,7 +59,7 @@ - + @@ -70,7 +70,7 @@ - + diff --git a/docs/diagrams/architecture-overview.drawio.svg b/docs/diagrams/architecture-overview.drawio.svg index 881fe960..c3df3e1b 100644 --- a/docs/diagrams/architecture-overview.drawio.svg +++ b/docs/diagrams/architecture-overview.drawio.svg @@ -1,4 +1,4 @@ -
TileXR — Layered Architecture
Applications, Examples & Integrations
vLLM-Ascend Integration
(integrations/vllm_ascend)
Tests / Demos / Simulator
TileXR Libraries
libtile-comm.so
Core Runtime: ranks, shared buffers,
peer memory, CommArgs, DFX
libtilexr-collectives.so
(optional)
AllGather / AllToAll
libtilexr-ep.so
MoE EP Dispatch MVP
Public API & Device Headers (src/include)
tilexr_api.h
(lifecycle / CommArgs)
tilexr_collectives.h
tilexr_ep.h
tilexr_sync.h
(device flag rounds)
tilexr_udma.h
tilexr_sdma.h (device)
Transports / Data Plane
IPC / MTE
peer-memory windows
(peerMems[], IPC_DATA_OFFSET)
UDMA (HCCP / RA)
registered remote memory
put / get / signal / quiet
SDMA (PTO)
on-card GM-to-GM copy
opt-in: TILEXR_ENABLE_SDMA=1
Platform Runtime
CANN 9.1.0 — ACL / Runtime APIs
Driver HAL (libascend_hal.so) — driver ≥ 25.5.0
Hardware
Ascend NPU — 910B / 910A5 / Ascend950 (A5, UDMA data-plane target)
Text is not SVG - cannot display \ No newline at end of file +
TileXR — Layered Architecture
Applications, Examples & Integrations
vLLM-Ascend Integration
(integrations/vllm_ascend)
Tests / Demos / Simulator
TileXR Libraries
libtile-comm.so
Core Runtime: ranks, shared buffers,
peer memory, CommArgs, DFX
libtilexr-collectives.so
(optional)
AllGather / AllToAll
libtilexr-ep.so
MoE EP Dispatch MVP
Public API & Device Headers (src/include)
tilexr_api.h
(lifecycle / CommArgs)
tilexr_collectives.h
tilexr_ep.h
tilexr_sync.h
(device flag rounds)
tilexr_udma.h
tilexr_sdma.h (device)
Transports / Data Plane
IPC / MTE
peer-memory windows
(peerMems[], IPC_DATA_OFFSET)
UDMA (HCCP / RA)
registered remote memory
put / get / signal / quiet
SDMA (PTO / A5 direct)
on-card GM-to-GM copy
opt-in: TILEXR_ENABLE_SDMA=1
Platform Runtime
CANN 9.1.0 — ACL / Runtime APIs
Driver HAL (libascend_hal.so) — driver ≥ 25.1.rc1
Hardware
Ascend NPU — 910B / 910A5 / Ascend950 (A5, UDMA data-plane target)
Text is not SVG - cannot display \ No newline at end of file diff --git a/docs/plans/2026-07-28-a5-direct-sdma-implementation.md b/docs/plans/2026-07-28-a5-direct-sdma-implementation.md new file mode 100644 index 00000000..e1931c47 --- /dev/null +++ b/docs/plans/2026-07-28-a5-direct-sdma-implementation.md @@ -0,0 +1,373 @@ +# A5 Direct SDMA Implementation Plan + +Date: 2026-07-28 +Status: Approved for execution +Design: `docs/specs/2026-07-28-a5-direct-sdma-design.md` + +## Goal + +Implement the approved Ascend950/A5 direct-SDMA backend inside TileXR while +preserving the existing A2/A3 PTO backend and public API. The production path +must use CANN's built-in system AICPU `ShmemSdmaStarsQuery`, map Host RTSQ +doorbells, publish a TileXR-owned 48-channel workspace, and let AIV submit the +two-SQE data/completion sequence directly. + +The project-wide supported baseline becomes CANN 9.1.0 with driver +`25.1.rc1` or later. The validated hardware baseline is +`Ascend950PR_9589`, driver `25.1.rc1.b188`, device 5. + +## Scope + +- Add a fixed-width, versioned A5 workspace and SQE ABI shared by Host and AIV. +- Add a production A5 Host backend to `libtile-comm.so`, including dynamic + built-in-op loading, 48 STARS streams, query compatibility, RTSQ mappings, + workspace publication, and complete failure cleanup. +- Dispatch `SDMACopyNbi` and `SDMAWait` to A5 direct submission for dav-3510 + kernels while keeping PTO behavior on A2/A3. +- Extend SDMA unit, integration, build, demo, dependency, and hardware checks. +- Change requirements and CI provisioning from driver `25.5.0+` to + `25.1.rc1+`, including Huawei RC/build-suffix comparison. +- Replace the temporary preflight and probe with the production path, and + retire obsolete A5 PoC documents. + +## Non-Goals + +- Do not build, install, embed, or load a TileXR custom AICPU kernel or OPP. +- Do not set or document `ASCEND_CUSTOM_OPP_PATH`. +- Do not include or link any source under `reference/`; it remains evidence + only. +- Do not change `CommArgs`, public Host API signatures, the opt-in environment + variable, or event-zero semantics. +- Do not replace PTO SDMA on A2/A3, make SDMA mandatory for communicator init, + add more than one outstanding event per channel, or claim Host/simulator + checks prove the A5 data plane. +- Do not add the CANN `devlib` directory to runtime RPATH/RUNPATH. + +## Authoritative References + +- Approved behavior and acceptance: the design linked above. +- Existing Host ownership and fallback: `src/comm/sdma/tilexr_sdma_transport.*` + and `src/comm/tilexr_comm.*`. +- Existing device API and PTO adapter: `src/include/tilexr_sdma.h` and + `src/include/tilexr_sdma_compat.h`. +- Probe evidence to extract before deletion: + `tests/sdma/a5_aicpu_probe/{main.asc,sdma_probe_kernel.asc,sdma_probe_types.h}`. +- Build/install wiring: `src/comm/CMakeLists.txt` and `tests/sdma/CMakeLists.txt`. +- Driver enforcement: `scripts/ci/provision/{common,cann,verify}.sh` and + `tests/ci/test_control_contract.py`. +- Build and hardware expectations: `docs/BUILD_VERIFICATION.md` and + `docs/SDMA_TRANSPORT.md`. + +`reference/` may be consulted to verify layouts and runtime behavior, but it is +not an implementation dependency and must not appear in active target include +paths, sources, or link inputs. + +## Dependency Order + +Tasks 1 and 2 establish contracts needed by Tasks 3 and 4. Task 5 depends on +the production Host and device paths. Task 6 is independent of Tasks 2-5 at the +file level but must land before final validation. Task 7 follows functional +implementation and baseline migration. Task 8 is the final acceptance gate; +PoC removal is not considered complete until its evidence is covered there. + +## Task 1: Define The A5 ABI And Test Seams + +**Objective and role:** Define the only Host/AIV binary contract for the A5 +backend and the pure validation/event helpers needed for deterministic unit +tests. This prevents the probe's diagnostic structs from becoming an implicit +production ABI. + +**Background and prerequisites:** Follow the approved 48-channel, +64-byte-aligned workspace, two-SQE completion protocol, one-outstanding-event +rule, and event-zero contract. Reuse the validated 64-byte A5 SQE field layout +from the probe, with explicit static assertions for every relied-upon size and +offset. + +**Modification scope:** Add an installed shared header such as +`src/include/tilexr_sdma_a5_types.h`; update `src/include/tilexr_sdma_types.h`, +`src/comm/CMakeLists.txt`, and focused SDMA metadata/header tests. Keep Host-only +runtime handles and ACL objects out of the installed ABI. + +**Constraints and non-goals:** Use fixed-width integer fields, C++14-compatible +constructs, device-visible addresses, a nonzero magic/version, 48 channels, and +64-byte cache-line alignment. Do not expose a new public Host function or reuse +PTO workspace layout. Do not use implementation-defined event bitfields; use +mask/shift helpers with invalid/stale detection. + +**Acceptance and verification:** Build Host and dav-3510 header compile tests; +unit-test structure sizes/offsets, queue wraparound, maximum length checks, +channel range, event encode/decode, generation wrap avoiding zero, malformed +events, and stale-generation rejection. + +**Artifacts and interfaces:** A stable A5 workspace header, channel record, +completion payload/record, SQE definition, backend identifier, and pure helpers +consumed by the Host initializer, AIV adapter, and tests. + +## Task 2: Add Dynamic CANN And Driver Runtime Adapters + +**Objective and role:** Isolate all optional A5 runtime calls behind an owned, +injectable adapter so missing symbols or system OPP support become a clean SDMA +capability failure and Host logic can be tested without hardware. + +**Background and prerequisites:** The built-in opapi lifecycle requires tensor +creation/destruction, `aclnnShmemSdmaStarsQueryGetWorkspaceSize`, and +`aclnnShmemSdmaStarsQuery`. Stream/SQ/CQ/logical-CQ/physical-die discovery, +HAL query, and `halResAddrMap`/`halResAddrUnmap` are also required. The current +core target already links ACL/runtime/real driver HAL; the built-in opapi entry +points must be resolved dynamically. + +**Modification scope:** Add private files under `src/comm/sdma/` for the A5 +loader/runtime function table and RAII resource wrappers; wire them into +`src/comm/CMakeLists.txt`; add fake-function-table unit tests under +`tests/sdma/unit/`. + +**Constraints and non-goals:** Load only CANN/runtime shared libraries expected +from the active CANN environment, report the first missing boundary, and close +owned handles idempotently. Do not link a TileXR OPP, depend on probe binaries, +or fall back to a custom `.aicpu` kernel. Do not source symbols from +`reference/` or introduce `devlib` runtime search paths. + +**Acceptance and verification:** Unit tests cover complete symbol resolution, +each required-symbol failure, idempotent close, and no-call behavior after a +failed load. `readelf -d/-l` and `ldd` later confirm no custom OPP dependency, +no CANN `devlib` RPATH/RUNPATH, and real-driver HAL resolution. + +**Artifacts and interfaces:** A private A5 runtime operation table plus scoped +wrappers for contexts, streams, tensors, device allocations, op workspaces, +RTSQ mappings, and library handles. Task 3 consumes this adapter and tests can +replace every external operation with deterministic fakes. + +## Task 3: Implement The 48-Channel A5 Host Backend + +**Objective and role:** Create, validate, publish, and destroy the complete A5 +SDMA resource set while containing all failures as best-effort capability +fallback. + +**Background and prerequisites:** Use the Task 1 ABI and Task 2 runtime adapter. +Create 48 `ACL_STREAM_DEVICE_USE_ONLY` streams in the communicator's device and +context. Present all stream records to a built-in query in one isolated context +first. Accept a complete successful result, or accept the `25.1.rc1` +compatibility path only for exact sync status `507018`, zero register-base, and +otherwise complete validated data; then query each remaining channel in a +fresh isolated context. + +**Modification scope:** Add private A5 backend/query implementation files under +`src/comm/sdma/`; refactor `src/comm/sdma/tilexr_sdma_transport.*` into a +runtime-SoC dispatcher; remove `PreflightAscend950SDMA`; update status values +only if diagnostics require a distinct A5 failure category. Preserve the +existing `TileXRComm::InitSDMA` publication flow. + +**Constraints and non-goals:** Validate every returned status, ID, SQ base, +64-byte SQE size, depth, head/tail range, Host HAL tail/SQE-size cross-check, +and RTSQ mapping length before publication. Restore the communicator context +and run a lightweight allocation/memset/sync/copy health check after every +expected partial failure. Publish only after all 48 channels succeed. On any +failure, unmap/destroy/free only TileXR-owned resources in reverse order, leave +the workspace null and SDMA flag clear, return TileXR success to communicator +initialization, and support repeated `Shutdown()`. + +**Acceptance and verification:** Fake-runtime tests cover full-query success, +expected partial batch plus 47 isolated queries, unexpected sync codes, +malformed/partial channel records, failed context restore or health check, +HAL mismatch, RTSQ failure, allocation/copy failure at every ownership stage, +reverse cleanup, reinitialization, and unchanged A2/A3 PTO selection. Hardware +validation later requires all 48 distinct streams/SQs/mappings alive together. + +**Artifacts and interfaces:** An A5 backend object owned by +`TileXRSDMATransport::Impl`, a fully initialized device workspace address, and +structured first-failure diagnostics. The transport exposes only its existing +availability/workspace/status interface to `TileXRComm`. + +## Task 4: Implement The AIV Direct Submission Backend + +**Objective and role:** Make the existing device calls submit and wait for A5 +SDMA directly on dav-3510 while leaving PTO calls unchanged on supported A2/A3 +kernels. + +**Background and prerequisites:** Use Task 1 layout and the approved two-SQE +protocol. Architecture selection follows the existing TileXR convention used +by UDMA: `__NPU_ARCH__ == 3510`, `CATLASS_ARCH == 3510`, plus a test-only force +macro if needed. The Host already selects and publishes the matching workspace. + +**Modification scope:** Add an installed A5 device adapter such as +`src/include/tilexr_sdma_a5.h`; update `src/include/tilexr_sdma.h`, install +wiring, and dav-3510 compile tests. Keep PTO-specific code in +`tilexr_sdma_compat.h`. + +**Constraints and non-goals:** Validate workspace magic/version/backend, +addresses, byte length, channel, depth, and queue capacity before writes. +Atomically claim the channel; a busy channel returns 0. Advance a nonzero +generation, clear completion, build data and 64-byte completion-copy SQEs with +type 11, `wrCqe=0`, credit 254, flush all touched cache lines, issue ordering, +and ring RTSQ only after both SQEs are ready. `SDMAWait(0)` succeeds; other +waits reject bad channel/generation/stale events, poll matching completion, and +release only their claimed channel. Never truncate a transfer length or ring a +doorbell after failed validation. + +**Acceptance and verification:** Host-pure helper tests cover invalid/busy/stale +paths and wraparound; dav-3510 compilation covers actual AIV intrinsics and SQE +writes; source guards ensure A5 does not include PTO headers and A2/A3 still do. +Hardware demo later verifies completion generation and every destination byte. + +**Artifacts and interfaces:** An A5 device adapter called by unchanged +`SDMACopyNbi`/`SDMAWait`, with unchanged event-zero semantics and explicit +compile-time routing between A5 and PTO implementations. + +## Task 5: Promote The Existing Demo To Production Acceptance + +**Objective and role:** Turn `tests/sdma/demo` into the single supported +data-plane validator for A2/A3 PTO and A5 direct SDMA, including boundary and +concurrency coverage. + +**Background and prerequisites:** Depends on Tasks 1-4. Current A5 build-target +selection changes are retained, but A5 demo availability must not depend on PTO +headers or `libnnopbase.so`. The demo must consume only installed TileXR APIs +and headers. + +**Modification scope:** Update `tests/sdma/CMakeLists.txt`, `tests/sdma/build.sh`, +`tests/sdma/run_tests.sh`, and `tests/sdma/demo/*`; add focused unit/source +tests for build selection and CLI validation. Integrate repeat/channel/block +options into the existing runner rather than creating a second probe target. + +**Constraints and non-goals:** Preserve the Ascend910B default and explicit +Ascend950 selection. Reject unsupported SoCs, channel IDs, sizes, or zero +iterations. Use distinct channels per AIV block, bounded Host waits, full byte +comparison, and generation checks. A skipped device or unavailable SDMA is not +an A5 data-plane pass. + +**Acceptance and verification:** On A5, run 64 B, 4 KiB, and 1 MiB on channel 0 +and channel 47; run a multi-block copy with distinct channels concurrently; +and repeat init/use/shutdown while confirming communicator health. On A2/A3 or +Host-only environments, existing unit/build behavior remains valid and claims +remain scoped. + +**Artifacts and interfaces:** One installed demo executable/kernel and runner +that prints SoC, driver, device, active backend, channel count, selected +channels, sizes, iterations, completion generations, and comparison results. + +## Task 6: Migrate The Driver Baseline To 25.1.rc1 + +**Objective and role:** Make repository guidance and CI enforcement agree with +the validated supported baseline, including Huawei RC syntax. + +**Background and prerequisites:** The minimum is `25.1.rc1`; build-suffixed +`25.1.rc1.b188` is supported. Comparison is semantic, not lexical: compare +major/minor first; for the same major/minor, RC numbers order before final +numeric patch releases; optional `.bN` is metadata and does not lower the base +version. A syntactically valid later major/minor release is supported even if +its final component is an RC. + +**Modification scope:** Update `scripts/ci/provision/common.sh`, `cann.sh`, and +`verify.sh`; extend `tests/ci/test_control_contract.py`; update minimum-version +claims in `AGENTS.md`, `README.md`, `docs/BUILD_VERIFICATION.md`, +`docs/SDMA_TRANSPORT.md`, and both architecture diagram sources. Preserve any +`25.5.0` text that is explicitly labeled as a historical observed environment. + +**Constraints and non-goals:** Accept case-insensitive `rc`, positive RC +numbers, optional numeric build suffixes, final numeric releases, and later +major/minor versions. Reject `25.1.rc0`, older releases, missing components, +extra arbitrary suffixes, and malformed strings. Do not use `sort -V` or make +driver validation exact-equality based. + +**Acceptance and verification:** The comparator matrix accepts at least +`25.1.rc1`, `25.1.RC1`, `25.1.rc1.b188`, `25.1.rc2`, `25.1.0`, `25.2.rc1`, and +`26.0.0`; it rejects at least `25.1.rc0`, `25.0.99`, `24.99.99`, `25.1`, +`25.1.rc`, `25.1.rc1.bad`, and empty input. Provision and verify messages both +name `25.1.rc1`. + +**Artifacts and interfaces:** A shared shell comparator with documented +ordering behavior, its Python contract tests, and consistent repository/diagram +minimum-version text. + +## Task 7: Remove PoC Paths And Consolidate Documentation + +**Objective and role:** Leave one production implementation and one accurate +set of operator/validation instructions after extracting all useful probe +evidence. + +**Background and prerequisites:** Execute only after Tasks 3-6 have tests and +the production demo covers the PoC's successful data path. Preserve the +approved design and this durable implementation plan. + +**Modification scope:** Delete `tests/sdma/a5_aicpu_probe/`, +`docs/plans/2026-07-27-a5-sdma-adaptation.md`, and +`docs/plans/2026-07-27-a5-aicpu-sdma-poc.md`; remove the temporary Host +preflight and its source-guard expectations; rewrite README/SDMA/build docs and +script catalog references around the production backend. + +**Constraints and non-goals:** Do not delete historical validation facts that +remain accurate; relabel them as evidence and state exact hardware. Remove all +claims that A5 production support is pending, requires 25.5.0, or needs a +TileXR OPP/custom OPP path. + +**Acceptance and verification:** Repository searches find no active reference +to `PreflightAscend950SDMA`, `a5_aicpu_probe`, custom TileXR AICPU artifacts, +`ASCEND_CUSTOM_OPP_PATH`, or a `25.5.0+` minimum. Installed manifests contain +only the production library and headers. Documentation commands match runnable +scripts. + +**Artifacts and interfaces:** Updated user/developer documentation, retained +approved design and implementation plan, and no PoC build/install artifacts. + +## Task 8: Execute Final Local And A5 Acceptance + +**Objective and role:** Produce proportionate evidence that the implementation +is compatible, failure-safe, dependency-clean, and functional on the supported +A5 baseline. + +**Background and prerequisites:** All earlier tasks complete. Synchronize a +clean workspace to the approved Ascend950 validation host, use physical device +5 and the validated CANN 9.1.0 installation, and confirm synchronization before +and after remote work. + +**Modification scope:** No feature expansion. Test-driven fixes may touch only +the owning modules above; any architecture, API, or acceptance change returns +to design approval. + +**Constraints and non-goals:** Do not reset unrelated devices, alter the CANN +installation, install an OPP, or treat a skipped/fallback run as a pass. Keep +remote temporary outputs outside tracked paths or remove them after evidence is +captured. Preserve unrelated local worktree changes. + +**Acceptance and verification:** Run CI contract tests, SDMA unit/integration +tests, C++14 core build/install, and Ascend910B/dav-3510 header/kernel builds. +On device 5 require 48 simultaneously live validated channels; copies at 64 B, +4 KiB, and 1 MiB on channels 0 and 47; a distinct-channel concurrent +multi-block run; repeated init/use/shutdown; full byte and generation matches; +and a healthy communicator context after each loop. Run dependency checks with +`readelf` and `ldd`, confirming real driver HAL and no `devlib` RPATH/RUNPATH. +Record exact SoC, driver, CANN, device, channel count, sizes, concurrency, and +iterations in `docs/SDMA_TRANSPORT.md`. + +**Artifacts and interfaces:** Test logs or concise recorded results, updated +validation documentation, and a final diff/status review demonstrating that +only scoped production, test, baseline, diagram, and cleanup files changed. + +## Key Risks And Controls + +- **Expected AICPU failure corrupts the active context:** every query runs in a + disposable isolated context; restoration and a runtime health operation are + mandatory before accepting partial data. +- **Partial workspace is accidentally trusted:** acceptance requires the exact + status and every independent field/ID/HAL/mapping cross-check; all other + partial states fail closed. +- **Queue overwrite or cross-block race:** reserve two slots, permit one + outstanding event per channel, atomically claim it, and assign distinct + channels to concurrent blocks. +- **Cache/order mismatch:** flush payload, completion record, and both SQEs; + issue the validated barrier sequence before the RTSQ tail write. +- **ABI drift between Host and AIV:** one installed fixed-width header and + static assertions are compiled in both modes. +- **Optional runtime becomes a hard dependency:** dynamically resolve the + built-in opapi functions and keep all failures best-effort. +- **False hardware confidence:** distinguish compile/source/fallback evidence + from the named A5 data-plane acceptance matrix. +- **Cleanup damages application state:** record ownership per resource, destroy + in reverse order, restore the prior context/device, and never reset a device. + +## Completion Gate + +The change is complete only when all eight tasks pass, A2/A3 PTO selection is +preserved, A5 device 5 passes the full hardware matrix with 48 live channels, +driver `25.1.rc1.b188` is accepted by the shared comparator, dependency checks +are clean, and the repository contains neither the PoC nor a TileXR OPP path. diff --git a/docs/specs/2026-07-28-a5-direct-sdma-design.md b/docs/specs/2026-07-28-a5-direct-sdma-design.md new file mode 100644 index 00000000..1d7375c7 --- /dev/null +++ b/docs/specs/2026-07-28-a5-direct-sdma-design.md @@ -0,0 +1,222 @@ +# A5 Direct SDMA Design + +Date: 2026-07-28 +Status: Approved on 2026-07-28 + +## Goal + +Add a production Ascend950/A5 backend for TileXR's opt-in local SDMA API. The +backend obtains STARS queue metadata through the CANN system AICPU, maps the +RTSQ doorbells on Host, and lets AIV write and submit A5 SDMA SQEs directly. + +The implementation must: + +- support CANN 9.1.0 and NPU driver `25.1.rc1` or later; +- preserve the existing A2/A3 PTO backend and public API; +- remain best-effort when any A5 capability is unavailable; +- ship no TileXR OPP package and require no custom OPP environment variable; +- keep `reference/` comparison-only. + +The validated baseline instance is driver `25.1.rc1.b188` on +`Ascend950PR_9589`. `25.1.rc1` is the project-wide minimum driver version, not +an exception specific to SDMA. + +## Delivery Boundary + +TileXR ships the A5 Host backend and AIV submission implementation in its +normal library and installed headers. It does not build, install, or load a +TileXR custom AICPU SO or OPP package. + +The backend dynamically resolves the public CANN opapi entry points for the +built-in `ShmemSdmaStarsQuery` operator. That operator and its system AICPU +kernel remain CANN-provided runtime dependencies. No active TileXR target +includes or links source from `reference/`, modifies the CANN installation, or +sets `ASCEND_CUSTOM_OPP_PATH`. + +An embedded `.aicpu` SO is not a substitute for the built-in query. CANN +registers binaries loaded through `aclrtBinaryLoadFromData` as custom AICPU +kernels. Hardware probing on the baseline host showed that this process cannot +see a Host-created STARS SQ: head, tail, depth, SQE size, and base queries all +returned status 3. The system AICPU query can see those resources. + +## Public Compatibility + +The following contracts remain unchanged: + +- opt-in enablement through `TILEXR_ENABLE_SDMA=1`; +- `CommArgs::sdmaWorkspacePtr` and `ExtraFlag::SDMA`; +- `TileXRSDMAAvailable` and `TileXRGetSDMAWorkspaceDev`; +- device calls `SDMACopyNbi` and `SDMAWait`; +- event handle 0 as the unavailable or invalid no-op result. + +Host selects the backend by runtime SoC: + +- A2/A3 continues to use PTO `SdmaWorkspaceManager` and PTO device intrinsics; +- Ascend950 uses the TileXR A5 backend described here; +- unsupported or failed initialization leaves SDMA unavailable without failing + communicator initialization. + +## Host Initialization + +### Owned resources + +One A5 transport owns, for the communicator lifetime: + +- 48 `ACL_STREAM_DEVICE_USE_ONLY` STARS streams; +- the stream, SQ, CQ, logical-CQ, and physical-die identifiers; +- one `PROCESS_CP1 + RES_ADDR_TYPE_STARS_RTSQ` mapping per SQ; +- the final TileXR A5 workspace in ordinary device memory; +- temporary query buffers and isolated query contexts used only during init. + +The CQ identifiers are validation and diagnostic fields. The data plane uses +`wrCqe=0` and completion-copy SQEs, so it does not consume a CQ base address. + +### Capability-driven query + +Host first presents all 48 streams to one built-in query in an isolated ACL +context. + +1. If launch and stream synchronization succeed, Host validates every channel + and takes the complete result as the fast path. +2. On driver `25.1.rc1`, the query fills the first channel and then fails while + requesting unsupported `DRV_SQCQ_PROP_SQ_REG_BASE`. Host may accept this + partial result only when synchronization returns exactly `507018`, the + register-base field is zero, and all earlier fields are complete. +3. After a partial batch result, Host queries every remaining stream separately + in a fresh isolated context. Each query writes into its own workspace slot. +4. After every expected partial failure, Host destroys the failed context, + restores the communicator context, and performs a lightweight runtime health + operation before continuing. + +The compatibility path is based on observed behavior, not on version-string +branching. A four-channel probe on `25.1.rc1.b188` initialized four simultaneously +live SQs successfully. The production acceptance gate extends this to all 48. + +### Validation and RTSQ mapping + +For each channel, Host requires: + +- successful local-device, head, tail, depth, SQE-size, and SQ-base statuses; +- matching stream/SQ/CQ/logical-CQ/device identifiers; +- a nonzero SQ base, 64-byte SQE size, and depth of at least three; +- head and tail inside the reported depth; +- Host HAL tail and SQE-size cross-checks matching the AICPU result; +- a nonzero RTSQ mapping of at least four bytes. + +The Host mapping replaces the unsupported query result for `sq_reg_base`. +Partial data is never accepted after an unexpected status, malformed field, +context restore failure, or failed health check. + +After all channels validate, Host copies a TileXR-owned workspace to GM and +publishes its address through `CommArgs`. Query intermediates are then freed; +STARS streams and RTSQ mappings remain alive until transport shutdown. + +## A5 Workspace ABI + +The A5 workspace is independent of PTO and uses fixed-width, 64-byte-aligned +structures shared by Host and AIV. It contains: + +- a header with magic, ABI version, backend kind, channel count, and SQE size; +- 48 channel records with SQ base, RTSQ address, depth, current tail, IDs, + generation, and outstanding state; +- one 64-byte completion payload and one 64-byte completion record per channel; +- reserved space for ABI-compatible diagnostics. + +Host and device compilation enforce sizes and key offsets with static +assertions. Every pointer stored in the workspace is a device-visible address. +The ABI version changes whenever an existing field's meaning or layout changes. + +Each channel allows one outstanding event. Different channels are independent +and may progress concurrently. The default channel group resolves from the AIV +block index and must be less than 48; explicit channel groups follow the same +range rule. + +## AIV Submission + +`SDMACopyNbi` on A5 performs these steps: + +1. Validate the workspace ABI, pointers, byte count, and channel group. +2. Atomically claim the channel's outstanding state and advance its nonzero + generation. A busy channel returns event 0. +3. Clear the completion record and write the generation into the channel's + completion payload. +4. Build two consecutive 64-byte A5 SQEs at the current tail: one data copy and + one 64-byte copy from the payload to the completion record. +5. Set SQE type 11, `wrCqe=0`, kernel credit 254, and the validated A5 address, + substream, length, and task fields. +6. Clean the written payload, completion record, and SQE cache lines, execute the + required store-ordering barrier, then write the new tail to the mapped RTSQ. +7. Return an event encoding the channel and generation. + +The tail advances modulo the queried depth. Completion of the second SQE makes +the channel safe for reuse. Transfer length handling must respect the A5 SQE +field width; unsupported lengths return event 0 instead of truncating. + +`SDMAWait` rejects malformed or stale events, polls the channel completion +record for the matching generation, then releases the outstanding state. A +zero event remains an immediate successful no-op. + +## Failure and Cleanup + +All A5 initialization failures are capability failures, not communicator +failures. TileXR logs the first failed boundary, releases only resources it +owns, leaves `sdmaWorkspacePtr` null, clears `ExtraFlag::SDMA`, and continues +with existing communication paths. + +Shutdown restores the owning device/context as needed, unmaps RTSQ mappings, +destroys STARS streams, frees the workspace, and is idempotent after partial +initialization. It never resets a device or destroys an application-owned +stream or context. + +Device submission failures are contained to the SDMA call. They must not write +a doorbell after failed validation or expose a nonzero event before both SQEs +are ready. + +## Driver Baseline Migration + +Repository requirements, architecture diagrams, build verification, and CI +provisioning change from `25.5.0 or later` to `25.1.rc1 or later`. + +The shared CI version comparator must recognize Huawei release-candidate forms: + +- accept `25.1.rc1`, `25.1.RC1`, and build-suffixed forms such as + `25.1.rc1.b188`; +- accept later RCs, final `25.1.x` releases, and later major/minor releases; +- reject versions earlier than RC1 and malformed or incomplete strings. + +Occurrences of `25.5.0` that record an actual 910B validation environment remain +historical facts. Only minimum-version claims and enforcement are changed. + +## Verification + +Implementation acceptance requires: + +- C++14 build with CANN 9.1.0 for the core library and Ascend950 demo kernel; +- unit coverage for A5 ABI layout, backend selection, partial-query validation, + cleanup, invalid events, busy channels, and best-effort fallback; +- CI tests for the `25.1.rc1` version-ordering matrix and provisioning messages; +- no TileXR OPP artifact, vendor install directory, or custom OPP environment; +- 48 distinct, simultaneously live, validated channels on the baseline host; +- direct copies of 64 B, 4 KiB, and 1 MiB on channel 0 and channel 47; +- a concurrent multi-block run using distinct channels, followed by full byte + comparisons and matching completion generations; +- repeated init/use/shutdown loops with a healthy communicator context; +- real driver `libascend_hal.so` resolution and no CANN `devlib` RPATH/RUNPATH. + +Host/source tests and simulator runs do not prove the A5 data plane. Hardware +claims must name the exercised SoC, driver, device, channel count, sizes, and +concurrency level. + +## Repository Cleanup + +The implementation change replaces the PoC-only material with this design and +the production backend. It removes: + +- `tests/sdma/a5_aicpu_probe/`; +- the A5 PoC and obsolete adaptation plans under `docs/plans/`; +- the temporary Ascend950 preflight that disables the PTO path; +- documentation that describes the production A5 path as pending or requires + driver 25.5.0. + +The existing SDMA demo becomes the single data-plane acceptance entry point for +both A2/A3 and A5. diff --git a/scripts/README.md b/scripts/README.md index c37f0f57..fe288997 100644 --- a/scripts/README.md +++ b/scripts/README.md @@ -75,6 +75,21 @@ bash scripts/download_open_source_deps.sh --check ## Testing +### SDMA + +The production SDMA build, unit-test, and hardware-demo entry points live under +`tests/sdma/`: + +```bash +bash tests/sdma/build.sh /path/to/cann Ascend950 +bash tests/sdma/run_tests.sh /path/to/cann +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 64 4096 1048576 +``` + +Omit `Ascend950` to compile the Ascend910B PTO demo. See +[`docs/SDMA_TRANSPORT.md`](../docs/SDMA_TRANSPORT.md) for channel, concurrency, +repeat, and dependency checks. + ### Pull-request CI The CI operator runbook is [docs/CI.md](../docs/CI.md). Primary entrypoints are: diff --git a/scripts/ci/control/build_blue.sh b/scripts/ci/control/build_blue.sh index 9880e917..59f17d5e 100755 --- a/scripts/ci/control/build_blue.sh +++ b/scripts/ci/control/build_blue.sh @@ -246,6 +246,8 @@ rm -rf \ "${SOURCE_DIR}/tests/udma/install" \ "${SOURCE_DIR}/tests/sdma/build" \ "${SOURCE_DIR}/tests/sdma/install" \ + "${SOURCE_DIR}/tests/sdma/build-a5" \ + "${SOURCE_DIR}/tests/sdma/install-a5" \ "${SOURCE_DIR}/tests/ep/build" \ "${SOURCE_DIR}/tests/ep/install" \ "${SOURCE_DIR}/tests/memory/build" \ @@ -278,12 +280,22 @@ run_case udma-source-guard "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udm run_logged_step "sdma-build" bash \ "${SOURCE_DIR}/tests/sdma/build.sh" "${ASCEND_HOME_PATH}" run_case sdma-metadata "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_metadata" +run_case sdma-a5-validation "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_a5_validation" run_case sdma-api-invalid "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_api_invalid" run_case sdma-transport-disabled "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_transport_disabled" run_case sdma-comm-wiring "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_comm_wiring" run_case sdma-source-guard "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_source_guard" run_case sdma-header-compile "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_header_compile" +run_logged_step "sdma-a5-configure" cmake \ + -S "${SOURCE_DIR}/tests/sdma" -B "${SOURCE_DIR}/tests/sdma/build-a5" \ + -DCMAKE_INSTALL_PREFIX="${SOURCE_DIR}/tests/sdma/install-a5" \ + -DBUILD_TILEXR_SDMA_DEMO=ON \ + -DTILEXR_SDMA_DEMO_SOC_TYPE=Ascend950 +run_logged_step "sdma-a5-build" cmake \ + --build "${SOURCE_DIR}/tests/sdma/build-a5" \ + --target tilexr_sdma_demo_kernel -j"${BUILD_JOBS}" + run_logged_step "ep-build" bash "${SOURCE_DIR}/tests/ep/build.sh" full run_case ep-layout "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_layout" run_case ep-api-sources "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_api_sources" @@ -319,7 +331,7 @@ for required_header in \ tilexr_api.h tilexr_types.h comm_args.h tilexr_sync.h \ tilexr_data_as_flag.h tilexr_perf_trace.h \ tilexr_udma.h tilexr_udma_reg.h tilexr_udma_types.h \ - tilexr_sdma_config.h tilexr_sdma_types.h tilexr_sdma.h tilexr_sdma_compat.h \ + tilexr_sdma_config.h tilexr_sdma_types.h tilexr_sdma_a5_types.h tilexr_sdma_a5.h tilexr_sdma.h tilexr_sdma_compat.h \ tilexr_ep.h tilexr_collectives.h tilexr_collectives_perf.h do require_regular_file "${SOURCE_DIR}/install/include/${required_header}" diff --git a/scripts/ci/provision/cann.sh b/scripts/ci/provision/cann.sh index af3947f4..7c453f22 100755 --- a/scripts/ci/provision/cann.sh +++ b/scripts/ci/provision/cann.sh @@ -143,8 +143,8 @@ check_driver_version() { local driver_version driver_version="$(awk -F= '$1 == "Version" {print $2; exit}' \ /usr/local/Ascend/driver/version.info 2>/dev/null)" || driver_version="" - if ! version_at_least "${driver_version}" 25.5.0; then - echo "ERROR: driver >= 25.5.0 is required, found ${driver_version:-unknown}" >&2 + if ! version_at_least "${driver_version}" 25.1.rc1; then + echo "ERROR: driver >= 25.1.rc1 is required, found ${driver_version:-unknown}" >&2 return 1 fi } diff --git a/scripts/ci/provision/common.sh b/scripts/ci/provision/common.sh index c53b2466..b0e49cc0 100755 --- a/scripts/ci/provision/common.sh +++ b/scripts/ci/provision/common.sh @@ -46,22 +46,45 @@ require_root() { version_at_least() { local actual="$1" local minimum="$2" - local actual_major actual_minor actual_patch - local minimum_major minimum_minor minimum_patch - - [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)$ ]] || return 1 - actual_major=$((10#${BASH_REMATCH[1]})) - actual_minor=$((10#${BASH_REMATCH[2]})) - actual_patch=$((10#${BASH_REMATCH[3]})) - [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)$ ]] || return 1 - minimum_major=$((10#${BASH_REMATCH[1]})) - minimum_minor=$((10#${BASH_REMATCH[2]})) - minimum_patch=$((10#${BASH_REMATCH[3]})) + local actual_major actual_minor actual_release actual_stage + local minimum_major minimum_minor minimum_release minimum_stage + + if [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.[rR][cC]([0-9]+)(\.b([0-9]+))?$ ]]; then + [[ "${BASH_REMATCH[3]}" != 0 ]] || return 1 + actual_major=$((10#${BASH_REMATCH[1]})) + actual_minor=$((10#${BASH_REMATCH[2]})) + actual_release=$((10#${BASH_REMATCH[3]})) + actual_stage=0 + elif [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)(\.b([0-9]+))?$ ]]; then + actual_major=$((10#${BASH_REMATCH[1]})) + actual_minor=$((10#${BASH_REMATCH[2]})) + actual_release=$((10#${BASH_REMATCH[3]})) + actual_stage=1 + else + return 1 + fi + + if [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.[rR][cC]([0-9]+)(\.b([0-9]+))?$ ]]; then + [[ "${BASH_REMATCH[3]}" != 0 ]] || return 1 + minimum_major=$((10#${BASH_REMATCH[1]})) + minimum_minor=$((10#${BASH_REMATCH[2]})) + minimum_release=$((10#${BASH_REMATCH[3]})) + minimum_stage=0 + elif [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)(\.b([0-9]+))?$ ]]; then + minimum_major=$((10#${BASH_REMATCH[1]})) + minimum_minor=$((10#${BASH_REMATCH[2]})) + minimum_release=$((10#${BASH_REMATCH[3]})) + minimum_stage=1 + else + return 1 + fi (( actual_major > minimum_major || (actual_major == minimum_major && actual_minor > minimum_minor) || (actual_major == minimum_major && actual_minor == minimum_minor && - actual_patch >= minimum_patch) )) + actual_stage > minimum_stage) || + (actual_major == minimum_major && actual_minor == minimum_minor && + actual_stage == minimum_stage && actual_release >= minimum_release) )) } ci_primary_group_has_non_root_gid() { diff --git a/scripts/ci/provision/verify.sh b/scripts/ci/provision/verify.sh index 6bc504d2..712f2f0d 100755 --- a/scripts/ci/provision/verify.sh +++ b/scripts/ci/provision/verify.sh @@ -33,8 +33,8 @@ check_driver_version() { local driver_version driver_version="$(awk -F= '$1 == "Version" {print $2; exit}' \ /usr/local/Ascend/driver/version.info 2>/dev/null)" || driver_version="" - if ! version_at_least "${driver_version}" 25.5.0; then - echo "ERROR: driver >= 25.5.0 is required, found ${driver_version:-unknown}" >&2 + if ! version_at_least "${driver_version}" 25.1.rc1; then + echo "ERROR: driver >= 25.1.rc1 is required, found ${driver_version:-unknown}" >&2 return 1 fi } diff --git a/src/comm/CMakeLists.txt b/src/comm/CMakeLists.txt index e0fd5253..202b2c24 100644 --- a/src/comm/CMakeLists.txt +++ b/src/comm/CMakeLists.txt @@ -118,6 +118,8 @@ set(TILEXR_SOURCE_FILE tilexr_comm.cpp udma/tilexr_udma_transport.cpp sdma/tilexr_sdma_transport.h sdma/tilexr_sdma_transport.cpp + sdma/tilexr_sdma_a5_backend.h + sdma/tilexr_sdma_a5_backend.cpp ) add_library(tile-comm SHARED ${TILEXR_SOURCE_FILE}) @@ -178,6 +180,8 @@ install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_udma_reg.h ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_udma_types.h ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sdma_types.h + ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sdma_a5_types.h + ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sdma_a5.h ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sdma.h ${CMAKE_CURRENT_SOURCE_DIR}/../include/tilexr_sdma_compat.h DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) diff --git a/src/comm/comm_wrap.cpp b/src/comm/comm_wrap.cpp index 27b18716..9c1d7878 100644 --- a/src/comm/comm_wrap.cpp +++ b/src/comm/comm_wrap.cpp @@ -328,7 +328,10 @@ int TileXRCommDestroy(TileXRCommPtr comm) return TILEXR_INVALID_VALUE; } auto *c = static_cast(comm); - + if (!c->PrepareDestroy()) { + TILEXR_LOG(ERROR) << "TileXR SDMA cleanup failed; communicator retained for retry"; + return TILEXR_ERROR_INTERNAL; + } delete c; return TILEXR_SUCCESS; } diff --git a/src/comm/sdma/tilexr_sdma_a5_backend.cpp b/src/comm/sdma/tilexr_sdma_a5_backend.cpp new file mode 100644 index 00000000..b9bb7357 --- /dev/null +++ b/src/comm/sdma/tilexr_sdma_a5_backend.cpp @@ -0,0 +1,710 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#include "sdma/tilexr_sdma_a5_backend.h" + +#include +#include +#include +#include +#include +#include + +#include "acl/acl.h" +#include "aclnn/aclnn_base.h" +#include "driver/ascend_hal.h" +#include "tilexr_log.h" + +#ifndef ACL_STREAM_DEVICE_USE_ONLY +#define ACL_STREAM_DEVICE_USE_ONLY 0x00000020U +#endif + +namespace TileXR { +namespace { + +constexpr size_t kBuiltinWorkspaceBytes = 16U * 1024U; +constexpr aclError kExpectedAicpuQueryFailure = + static_cast(detail::TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS); +constexpr int32_t kDeviceInfoModuleType = 0; +constexpr int32_t kPhysicalDieInfoType = 19; + +using RtGetDevicePhyIdByIndexFn = int32_t (*)(uint32_t, uint32_t*); +using RtStreamGetSqidFn = int32_t (*)(const void*, uint32_t*); +using RtStreamGetCqidFn = int32_t (*)(const void*, uint32_t*, uint32_t*); +using RtGetDeviceInfoFn = int32_t (*)(uint32_t, int32_t, int32_t, int64_t*); +using AclCreateTensorFn = aclTensor* (*)(const int64_t*, uint64_t, aclDataType, + const int64_t*, int64_t, aclFormat, + const int64_t*, uint64_t, void*); +using AclDestroyTensorFn = int32_t (*)(const aclTensor*); +using AclnnQueryWorkspaceFn = aclnnStatus (*)(const aclTensor*, aclTensor*, + uint64_t*, aclOpExecutor**); +using AclnnQueryFn = aclnnStatus (*)(void*, uint64_t, aclOpExecutor*, aclrtStream); + +template +bool LoadSymbol(void* handle, const char* name, T& symbol) +{ + symbol = reinterpret_cast(dlsym(handle, name)); + if (symbol != nullptr) { + return true; + } + TILEXR_LOG(WARN) << "TileXR A5 SDMA missing runtime symbol " << name; + return false; +} + +class A5RuntimeApi { +public: + ~A5RuntimeApi() + { + Close(); + } + + bool Load() + { + Close(); + runtimeHandle_ = dlopen("libruntime.so", RTLD_NOW | RTLD_LOCAL); + if (runtimeHandle_ == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not load libruntime.so: " << dlerror(); + return false; + } + if (!LoadSymbol(runtimeHandle_, "rtGetDevicePhyIdByIndex", getPhysicalDevice) || + !LoadSymbol(runtimeHandle_, "rtStreamGetSqid", getSqId) || + !LoadSymbol(runtimeHandle_, "rtStreamGetCqid", getCqId) || + !LoadSymbol(runtimeHandle_, "rtGetDeviceInfo", getDeviceInfo)) { + Close(); + return false; + } + + opapiHandle_ = dlopen("libopapi.so", RTLD_NOW | RTLD_LOCAL); + if (opapiHandle_ == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not load libopapi.so: " << dlerror(); + Close(); + return false; + } + if (!LoadSymbol(opapiHandle_, "aclCreateTensor", createTensor) || + !LoadSymbol(opapiHandle_, "aclDestroyTensor", destroyTensor) || + !LoadSymbol(opapiHandle_, "aclnnShmemSdmaStarsQueryGetWorkspaceSize", prepareQuery) || + !LoadSymbol(opapiHandle_, "aclnnShmemSdmaStarsQuery", executeQuery)) { + Close(); + return false; + } + return true; + } + + void Close() + { + getPhysicalDevice = nullptr; + getSqId = nullptr; + getCqId = nullptr; + getDeviceInfo = nullptr; + createTensor = nullptr; + destroyTensor = nullptr; + prepareQuery = nullptr; + executeQuery = nullptr; + if (opapiHandle_ != nullptr) { + (void)dlclose(opapiHandle_); + opapiHandle_ = nullptr; + } + if (runtimeHandle_ != nullptr) { + (void)dlclose(runtimeHandle_); + runtimeHandle_ = nullptr; + } + } + + RtGetDevicePhyIdByIndexFn getPhysicalDevice = nullptr; + RtStreamGetSqidFn getSqId = nullptr; + RtStreamGetCqidFn getCqId = nullptr; + RtGetDeviceInfoFn getDeviceInfo = nullptr; + AclCreateTensorFn createTensor = nullptr; + AclDestroyTensorFn destroyTensor = nullptr; + AclnnQueryWorkspaceFn prepareQuery = nullptr; + AclnnQueryFn executeQuery = nullptr; + +private: + void* runtimeHandle_ = nullptr; + void* opapiHandle_ = nullptr; +}; + +class DeviceBuffer { +public: + ~DeviceBuffer() + { + Reset(); + } + + bool Allocate(size_t bytes, bool zero = false) + { + Reset(); + if (aclrtMalloc(&address_, bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { + return false; + } + bytes_ = bytes; + if (zero && aclrtMemset(address_, bytes_, 0, bytes_) != ACL_SUCCESS) { + Reset(); + return false; + } + return true; + } + + void Reset() + { + if (address_ != nullptr) { + (void)aclrtFree(address_); + address_ = nullptr; + bytes_ = 0U; + } + } + + void* Get() const + { + return address_; + } + +private: + void* address_ = nullptr; + size_t bytes_ = 0U; +}; + +class TensorHandle { +public: + explicit TensorHandle(AclDestroyTensorFn destroy) : destroy_(destroy) {} + + ~TensorHandle() + { + Reset(); + } + + void Reset() + { + if (tensor_ != nullptr) { + (void)destroy_(tensor_); + tensor_ = nullptr; + } + } + + aclTensor*& Ref() + { + return tensor_; + } + + aclTensor* Get() const + { + return tensor_; + } + +private: + AclDestroyTensorFn destroy_; + aclTensor* tensor_ = nullptr; +}; + +struct QuerySnapshot { + uint32_t flag = 0U; + uint32_t totalQueueCount = 0U; + aclError syncStatus = ACL_SUCCESS; + std::vector channels; +}; + +bool CreateUint64Tensor(A5RuntimeApi& api, void* address, int64_t elements, + TensorHandle& tensor) +{ + const int64_t shape[] = {elements}; + const int64_t strides[] = {1}; + tensor.Ref() = api.createTensor(shape, 1U, ACL_UINT64, strides, 0, + ACL_FORMAT_ND, shape, 1U, address); + return tensor.Get() != nullptr; +} + +bool RestoreQueryContext(aclrtContext previous, aclrtContext& isolated) +{ + const aclError destroyStatus = isolated == nullptr + ? ACL_SUCCESS + : aclrtDestroyContext(isolated); + isolated = nullptr; + const aclError restoreStatus = aclrtSetCurrentContext(previous); + if (destroyStatus != ACL_SUCCESS || restoreStatus != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query context restore failed, destroy " + << destroyStatus << ", restore " << restoreStatus; + return false; + } + return true; +} + +bool CheckRuntimeHealth(DeviceBuffer& scratch) +{ + aclrtStream stream = nullptr; + if (aclrtCreateStream(&stream) != ACL_SUCCESS) { + return false; + } + const aclError memsetStatus = aclrtMemsetAsync( + scratch.Get(), sizeof(uint64_t), 0xA5, sizeof(uint64_t), stream); + const aclError syncStatus = memsetStatus == ACL_SUCCESS + ? aclrtSynchronizeStream(stream) + : memsetStatus; + uint64_t value = 0U; + const aclError copyStatus = syncStatus == ACL_SUCCESS + ? aclrtMemcpy(&value, sizeof(value), scratch.Get(), sizeof(value), ACL_MEMCPY_DEVICE_TO_HOST) + : syncStatus; + const aclError destroyStatus = aclrtDestroyStream(stream); + return memsetStatus == ACL_SUCCESS && syncStatus == ACL_SUCCESS && + copyStatus == ACL_SUCCESS && destroyStatus == ACL_SUCCESS && + value == 0xA5A5A5A5A5A5A5A5ULL; +} + +bool RunBuiltinQuery(A5RuntimeApi& api, + int32_t logicalDevice, + const std::vector& streams, + QuerySnapshot& snapshot) +{ + if (streams.empty() || streams.size() > detail::TILEXR_SDMA_A5_CHANNEL_COUNT) { + return false; + } + + const size_t streamsBytes = streams.size() * sizeof(detail::A5BuiltinStreamInfo); + DeviceBuffer streamsDev; + DeviceBuffer resourceDev; + DeviceBuffer builtinWorkspaceDev; + DeviceBuffer inputDev; + DeviceBuffer outputDev; + DeviceBuffer opWorkspaceDev; + if (!streamsDev.Allocate(streamsBytes) || + !resourceDev.Allocate(sizeof(detail::A5BuiltinOpResource), true) || + !builtinWorkspaceDev.Allocate(kBuiltinWorkspaceBytes, true) || + !inputDev.Allocate(2U * sizeof(uint64_t)) || + !outputDev.Allocate(sizeof(uint64_t), true) || + aclrtMemcpy(streamsDev.Get(), streamsBytes, streams.data(), streamsBytes, + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query buffer setup failed"; + return false; + } + + detail::A5BuiltinOpResource resource {}; + resource.size = streams.size(); + resource.streamsAddress = reinterpret_cast(streamsDev.Get()); + resource.workspaceAddress = reinterpret_cast(builtinWorkspaceDev.Get()); + const uint64_t inputs[] = { + reinterpret_cast(resourceDev.Get()), + reinterpret_cast(builtinWorkspaceDev.Get()), + }; + if (aclrtMemcpy(resourceDev.Get(), sizeof(resource), &resource, sizeof(resource), + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS || + aclrtMemcpy(inputDev.Get(), sizeof(inputs), inputs, sizeof(inputs), + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query resource upload failed"; + return false; + } + + aclrtContext previous = nullptr; + aclrtContext isolated = nullptr; + if (aclrtGetCurrentContext(&previous) != ACL_SUCCESS || + aclrtCreateContext(&isolated, logicalDevice) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA isolated query context creation failed"; + if (isolated != nullptr) { + (void)RestoreQueryContext(previous, isolated); + } + return false; + } + + TensorHandle inputTensor(api.destroyTensor); + TensorHandle outputTensor(api.destroyTensor); + aclrtStream queryStream = nullptr; + bool queryLaunched = false; + aclError syncStatus = ACL_SUCCESS; + uint64_t opWorkspaceBytes = 0U; + aclOpExecutor* executor = nullptr; + bool ok = CreateUint64Tensor(api, inputDev.Get(), 2, inputTensor) && + CreateUint64Tensor(api, outputDev.Get(), 1, outputTensor) && + api.prepareQuery(inputTensor.Get(), outputTensor.Get(), &opWorkspaceBytes, &executor) == ACL_SUCCESS; + if (ok && opWorkspaceBytes != 0U) { + ok = opWorkspaceDev.Allocate(static_cast(opWorkspaceBytes)); + } + if (ok) { + ok = aclrtCreateStreamWithConfig( + &queryStream, 0, ACL_STREAM_FAST_LAUNCH | ACL_STREAM_FAST_SYNC) == ACL_SUCCESS; + } + if (ok) { + aclrtStreamAttrValue failureMode {}; + failureMode.failureMode = 0; + ok = aclrtSetStreamAttribute( + queryStream, ACL_STREAM_ATTR_FAILURE_MODE, &failureMode) == ACL_SUCCESS; + } + if (ok) { + const aclnnStatus launchStatus = api.executeQuery( + opWorkspaceDev.Get(), opWorkspaceBytes, executor, queryStream); + queryLaunched = launchStatus == ACL_SUCCESS; + ok = queryLaunched; + if (queryLaunched) { + syncStatus = aclrtSynchronizeStream(queryStream); + } + } + + aclError streamDestroyStatus = ACL_SUCCESS; + if (queryStream != nullptr) { + streamDestroyStatus = aclrtDestroyStream(queryStream); + queryStream = nullptr; + } + inputTensor.Reset(); + outputTensor.Reset(); + const bool contextRestored = RestoreQueryContext(previous, isolated); + if (!ok || !queryLaunched || streamDestroyStatus != ACL_SUCCESS || !contextRestored) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA built-in query launch or cleanup failed"; + return false; + } + + const size_t snapshotBytes = sizeof(detail::A5BuiltinWorkspaceHeader) + + streams.size() * sizeof(detail::A5BuiltinChannelInfo); + std::vector bytes(snapshotBytes, 0U); + if (aclrtMemcpy(bytes.data(), bytes.size(), builtinWorkspaceDev.Get(), bytes.size(), + ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query workspace download failed"; + return false; + } + detail::A5BuiltinWorkspaceHeader header {}; + std::memcpy(&header, bytes.data(), sizeof(header)); + snapshot.flag = header.flag; + snapshot.totalQueueCount = header.totalQueueCount; + snapshot.syncStatus = syncStatus; + snapshot.channels.resize(streams.size()); + std::memcpy(snapshot.channels.data(), bytes.data() + sizeof(header), + snapshot.channels.size() * sizeof(snapshot.channels[0])); + + if (syncStatus == kExpectedAicpuQueryFailure && !CheckRuntimeHealth(outputDev)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA context health check failed after expected AICPU error"; + return false; + } + return true; +} + +int32_t QueryHostSq(uint32_t physicalDevice, uint32_t sqId, + drvSqCqPropType_t property, uint32_t (&values)[3]) +{ + halSqCqQueryInfo query {}; + query.type = DRV_NORMAL_TYPE; + query.tsId = 0U; + query.sqId = sqId; + query.cqId = 0U; + query.prop = property; + const drvError_t status = halSqCqQuery(physicalDevice, &query); + values[0] = query.value[0]; + values[1] = query.value[1]; + values[2] = query.value[2]; + return static_cast(status); +} + +} // namespace + +struct TileXRA5SDMABackend::Impl { + struct OwnedChannel { + aclrtStream stream = nullptr; + res_addr_info mapInfo {}; + bool mapped = false; + uint64_t rtsqAddress = 0U; + uint32_t rtsqLength = 0U; + detail::A5HostChannelIdentity identity {}; + detail::A5BuiltinChannelInfo query {}; + }; + + int32_t logicalDevice = -1; + uint32_t physicalDevice = 0U; + uint32_t physicalDieId = 0U; + aclrtContext ownerContext = nullptr; + aclrtContext restoreContext = nullptr; + bool restorePending = false; + void* workspaceDev = nullptr; + std::array channels {}; + + bool HasOwnedResources() const + { + if (workspaceDev != nullptr) { + return true; + } + for (const OwnedChannel& channel : channels) { + if (channel.mapped || channel.stream != nullptr) { + return true; + } + } + return false; + } +}; + +TileXRA5SDMABackend::TileXRA5SDMABackend() = default; + +TileXRA5SDMABackend::~TileXRA5SDMABackend() +{ + (void)Shutdown(); +} + +bool TileXRA5SDMABackend::Init(int32_t deviceId) +{ + if (!Shutdown()) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA previous resources could not be released"; + return false; + } + std::unique_ptr state(new (std::nothrow) Impl()); + if (state == nullptr) { + return false; + } + state->logicalDevice = deviceId; + if (aclrtGetCurrentContext(&state->ownerContext) != ACL_SUCCESS || + state->ownerContext == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA requires an active owner context"; + return false; + } + + A5RuntimeApi api; + if (!api.Load() || + api.getPhysicalDevice(static_cast(deviceId), &state->physicalDevice) != 0) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA runtime discovery failed"; + return false; + } + int64_t physicalDie = -1; + if (api.getDeviceInfo(static_cast(deviceId), kDeviceInfoModuleType, + kPhysicalDieInfoType, &physicalDie) != 0 || + physicalDie < 0 || static_cast(physicalDie) > + static_cast(std::numeric_limits::max())) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA physical die discovery failed"; + return false; + } + state->physicalDieId = static_cast(physicalDie); + impl_ = std::move(state); + + std::vector streamInfos; + streamInfos.reserve(detail::TILEXR_SDMA_A5_CHANNEL_COUNT); + for (uint32_t index = 0U; index < detail::TILEXR_SDMA_A5_CHANNEL_COUNT; ++index) { + Impl::OwnedChannel& owned = impl_->channels[index]; + if (aclrtCreateStreamWithConfig( + &owned.stream, 0, ACL_STREAM_DEVICE_USE_ONLY) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA STARS stream creation failed at channel " << index; + Shutdown(); + return false; + } + int32_t streamId = -1; + uint32_t sqId = 0U; + uint32_t cqId = 0U; + uint32_t logicalCqId = 0U; + if (aclrtStreamGetId(owned.stream, &streamId) != ACL_SUCCESS || streamId < 0 || + api.getSqId(owned.stream, &sqId) != 0 || + api.getCqId(owned.stream, &cqId, &logicalCqId) != 0) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA stream identifier query failed at channel " << index; + Shutdown(); + return false; + } + owned.identity = { + static_cast(streamId), sqId, cqId, logicalCqId, impl_->physicalDieId, + }; + owned.mapInfo.id = 0U; + owned.mapInfo.target_proc_type = PROCESS_CP1; + owned.mapInfo.res_type = RES_ADDR_TYPE_STARS_RTSQ; + owned.mapInfo.res_id = sqId; + unsigned long mappedAddress = 0UL; + unsigned int mappedLength = 0U; + const drvError_t mapStatus = halResAddrMap( + impl_->physicalDevice, &owned.mapInfo, &mappedAddress, &mappedLength); + owned.mapped = mapStatus == DRV_ERROR_NONE; + if (mapStatus != DRV_ERROR_NONE || mappedAddress == 0UL || + mappedLength < sizeof(uint32_t)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA RTSQ map failed at channel " << index + << ", status " << mapStatus; + Shutdown(); + return false; + } + owned.rtsqAddress = static_cast(mappedAddress); + owned.rtsqLength = mappedLength; + + detail::A5BuiltinStreamInfo info {}; + info.stream = reinterpret_cast(owned.stream); + info.context = reinterpret_cast(impl_->ownerContext); + info.streamId = streamId; + info.sqId = sqId; + info.cqId = cqId; + info.logicalCqId = logicalCqId; + info.deviceId = static_cast(impl_->physicalDieId); + streamInfos.push_back(info); + } + + QuerySnapshot batch; + if (!RunBuiltinQuery(api, deviceId, streamInfos, batch)) { + Shutdown(); + return false; + } + const detail::A5QueryResultKind batchKind = batch.channels.empty() + ? detail::A5QueryResultKind::INVALID + : detail::ClassifyA5QueryResult( + static_cast(batch.syncStatus), batch.flag, batch.totalQueueCount, + batch.channels[0], impl_->channels[0].identity); + if (batchKind == detail::A5QueryResultKind::COMPLETE) { + if (batch.channels.size() != detail::TILEXR_SDMA_A5_CHANNEL_COUNT) { + Shutdown(); + return false; + } + for (uint32_t index = 0U; index < detail::TILEXR_SDMA_A5_CHANNEL_COUNT; ++index) { + if (!detail::ValidateA5BuiltinChannel( + batch.channels[index], impl_->channels[index].identity, true)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA complete query validation failed at channel " << index; + Shutdown(); + return false; + } + impl_->channels[index].query = batch.channels[index]; + } + } else if (batchKind == detail::A5QueryResultKind::EXPECTED_PARTIAL) { + impl_->channels[0].query = batch.channels[0]; + for (uint32_t index = 1U; index < detail::TILEXR_SDMA_A5_CHANNEL_COUNT; ++index) { + std::vector oneStream(1U, streamInfos[index]); + QuerySnapshot isolated; + if (!RunBuiltinQuery(api, deviceId, oneStream, isolated) || + isolated.channels.size() != 1U || + detail::ClassifyA5QueryResult( + static_cast(isolated.syncStatus), isolated.flag, + isolated.totalQueueCount, isolated.channels[0], + impl_->channels[index].identity) != + detail::A5QueryResultKind::EXPECTED_PARTIAL) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA isolated query validation failed at channel " << index; + Shutdown(); + return false; + } + impl_->channels[index].query = isolated.channels[0]; + } + } else { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query returned unsupported status " << batch.syncStatus; + Shutdown(); + return false; + } + + detail::A5SdmaWorkspace hostWorkspace {}; + hostWorkspace.header.magic = detail::TILEXR_SDMA_A5_WORKSPACE_MAGIC; + hostWorkspace.header.abiVersion = detail::TILEXR_SDMA_A5_ABI_VERSION; + hostWorkspace.header.backendKind = detail::TILEXR_SDMA_A5_BACKEND_KIND; + hostWorkspace.header.channelCount = detail::TILEXR_SDMA_A5_CHANNEL_COUNT; + hostWorkspace.header.sqeSize = detail::TILEXR_SDMA_A5_SQE_BYTES; + hostWorkspace.header.channelStride = sizeof(detail::A5SdmaChannel); + hostWorkspace.header.workspaceSize = sizeof(hostWorkspace); + hostWorkspace.header.maxTransferBytes = + static_cast(detail::TILEXR_SDMA_A5_MAX_TRANSFER_BYTES); + + if (aclrtMalloc(&impl_->workspaceDev, sizeof(hostWorkspace), + ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA workspace allocation failed"; + Shutdown(); + return false; + } + const uint64_t workspaceBase = reinterpret_cast(impl_->workspaceDev); + for (uint32_t index = 0U; index < detail::TILEXR_SDMA_A5_CHANNEL_COUNT; ++index) { + const Impl::OwnedChannel& owned = impl_->channels[index]; + detail::A5SdmaChannel& channel = hostWorkspace.channels[index]; + channel.sqBase = owned.query.sqBase; + channel.rtsqAddress = owned.rtsqAddress; + channel.completionPayloadAddress = workspaceBase + + offsetof(detail::A5SdmaWorkspace, completionPayloads) + + index * sizeof(detail::A5SdmaCompletionLine); + channel.completionRecordAddress = workspaceBase + + offsetof(detail::A5SdmaWorkspace, completionRecords) + + index * sizeof(detail::A5SdmaCompletionLine); + channel.depth = owned.query.sqDepth; + channel.head = owned.query.sqHead; + channel.tail = owned.query.sqTail; + channel.taskId = detail::A5SdmaQueueDistance( + owned.query.sqHead, owned.query.sqTail, owned.query.sqDepth); + channel.rtsqLength = owned.rtsqLength; + channel.streamId = owned.identity.streamId; + channel.sqId = owned.identity.sqId; + channel.cqId = owned.identity.cqId; + channel.logicalCqId = owned.identity.logicalCqId; + channel.physicalDieId = owned.identity.physicalDieId; + + uint32_t tailValues[3] = {0U, 0U, 0U}; + uint32_t sqeSizeValues[3] = {0U, 0U, 0U}; + if (QueryHostSq(impl_->physicalDevice, channel.sqId, + DRV_SQCQ_PROP_SQ_TAIL, tailValues) != 0 || + QueryHostSq(impl_->physicalDevice, channel.sqId, + DRV_SQCQ_PROP_SQE_SIZE, sqeSizeValues) != 0 || + tailValues[0] != channel.tail || + sqeSizeValues[0] != detail::TILEXR_SDMA_A5_SQE_BYTES) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA Host SQ cross-check failed at channel " << index; + Shutdown(); + return false; + } + } + if (aclrtMemcpy(impl_->workspaceDev, sizeof(hostWorkspace), &hostWorkspace, + sizeof(hostWorkspace), ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA workspace upload failed"; + Shutdown(); + return false; + } + + TILEXR_LOG(INFO) << "TileXR A5 direct SDMA initialized on device " << deviceId + << " with " << detail::TILEXR_SDMA_A5_CHANNEL_COUNT << " channels"; + return true; +} + +bool TileXRA5SDMABackend::Shutdown() +{ + if (impl_ == nullptr) { + return true; + } + if (impl_->restorePending) { + if (aclrtSetCurrentContext(impl_->restoreContext) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not complete pending context restore"; + return false; + } + impl_->restoreContext = nullptr; + impl_->restorePending = false; + if (!impl_->HasOwnedResources()) { + impl_.reset(); + return true; + } + } + + aclrtContext previous = nullptr; + if (aclrtGetCurrentContext(&previous) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not capture current context for cleanup"; + return false; + } + if (impl_->ownerContext == nullptr || + aclrtSetCurrentContext(impl_->ownerContext) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not switch to owner context for cleanup"; + return false; + } + + if (impl_->workspaceDev != nullptr) { + if (aclrtFree(impl_->workspaceDev) == ACL_SUCCESS) { + impl_->workspaceDev = nullptr; + } + } + if (impl_->workspaceDev == nullptr) { + for (size_t reverse = detail::TILEXR_SDMA_A5_CHANNEL_COUNT; reverse > 0U; --reverse) { + Impl::OwnedChannel& owned = impl_->channels[reverse - 1U]; + bool mappingReleased = true; + if (owned.mapped) { + if (halResAddrUnmap(impl_->physicalDevice, &owned.mapInfo) == DRV_ERROR_NONE) { + owned.mapped = false; + } else { + mappingReleased = false; + } + } + if (mappingReleased && owned.stream != nullptr && + aclrtDestroyStream(owned.stream) == ACL_SUCCESS) { + owned.stream = nullptr; + } + } + } + if (previous != impl_->ownerContext && + aclrtSetCurrentContext(previous) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not restore previous context after cleanup"; + impl_->restoreContext = previous; + impl_->restorePending = true; + return false; + } + if (impl_->HasOwnedResources()) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA cleanup incomplete; retained resources for retry"; + return false; + } + impl_.reset(); + return true; +} + +GM_ADDR TileXRA5SDMABackend::GetWorkspaceDev() const +{ + return impl_ == nullptr ? nullptr : static_cast(impl_->workspaceDev); +} + +} // namespace TileXR diff --git a/src/comm/sdma/tilexr_sdma_a5_backend.h b/src/comm/sdma/tilexr_sdma_a5_backend.h new file mode 100644 index 00000000..d1b4a592 --- /dev/null +++ b/src/comm/sdma/tilexr_sdma_a5_backend.h @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_A5_BACKEND_H +#define TILEXR_SDMA_A5_BACKEND_H + +#include +#include +#include + +#include "comm_args.h" +#include "tilexr_sdma_a5_types.h" + +namespace TileXR { + +namespace detail { + +struct alignas(64) A5BuiltinStreamInfo { + uint64_t stream; + uint64_t context; + int32_t streamId; + uint32_t sqId; + uint32_t cqId; + uint32_t logicalCqId; + uint64_t cqeAddress; + int32_t deviceId; + uint8_t reserved[20]; +}; + +struct alignas(64) A5BuiltinOpResource { + uint64_t size; + uint64_t streamsAddress; + uint64_t workspaceAddress; + uint8_t reserved[40]; +}; + +struct alignas(64) A5BuiltinChannelInfo { + uint32_t sqHead; + uint32_t sqTail; + uint64_t sqBase; + uint64_t sqRegisterBase; + uint32_t sqDepth; + uint32_t sqId; + uint32_t cqId; + uint32_t logicalCqId; + uint64_t cqeAddress; + uint32_t reportCqeCount; + uint32_t streamId; + uint32_t deviceId; + uint8_t reserved[4]; +}; + +struct alignas(64) A5BuiltinWorkspaceHeader { + uint32_t flag; + uint32_t totalQueueCount; + uint8_t reserved[56]; +}; + +struct A5HostChannelIdentity { + uint32_t streamId; + uint32_t sqId; + uint32_t cqId; + uint32_t logicalCqId; + uint32_t physicalDieId; +}; + +constexpr int32_t TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS = 507018; + +enum class A5QueryResultKind : uint32_t { + INVALID = 0U, + COMPLETE = 1U, + EXPECTED_PARTIAL = 2U, +}; + +inline bool ValidateA5BuiltinChannel(const A5BuiltinChannelInfo& channel, + const A5HostChannelIdentity& expected, + bool requireRegisterBase) +{ + return channel.sqBase != 0U && + (!requireRegisterBase || channel.sqRegisterBase != 0U) && + A5SdmaQueueHasCapacity(channel.sqHead, channel.sqTail, channel.sqDepth) && + channel.streamId == expected.streamId && channel.sqId == expected.sqId && + channel.cqId == expected.cqId && channel.logicalCqId == expected.logicalCqId && + channel.deviceId == expected.physicalDieId; +} + +inline A5QueryResultKind ClassifyA5QueryResult(int32_t syncStatus, + uint32_t flag, + uint32_t totalQueueCount, + const A5BuiltinChannelInfo& firstChannel, + const A5HostChannelIdentity& expected) +{ + if (syncStatus == 0 && flag == 1U && + totalQueueCount == TILEXR_SDMA_A5_CHANNEL_COUNT && + ValidateA5BuiltinChannel(firstChannel, expected, true)) { + return A5QueryResultKind::COMPLETE; + } + if (syncStatus == TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS && flag == 0U && + totalQueueCount == 0U && firstChannel.sqRegisterBase == 0U && + ValidateA5BuiltinChannel(firstChannel, expected, false)) { + return A5QueryResultKind::EXPECTED_PARTIAL; + } + return A5QueryResultKind::INVALID; +} + +static_assert(sizeof(A5BuiltinStreamInfo) == 64U, "unexpected built-in stream ABI"); +static_assert(sizeof(A5BuiltinOpResource) == 64U, "unexpected built-in resource ABI"); +static_assert(sizeof(A5BuiltinChannelInfo) == 64U, "unexpected built-in channel ABI"); +static_assert(sizeof(A5BuiltinWorkspaceHeader) == 64U, "unexpected built-in header ABI"); +static_assert(offsetof(A5BuiltinChannelInfo, sqBase) == 8U, "unexpected built-in SQ base offset"); +static_assert(offsetof(A5BuiltinChannelInfo, sqRegisterBase) == 16U, + "unexpected built-in register offset"); +static_assert(offsetof(A5BuiltinChannelInfo, streamId) == 52U, + "unexpected built-in stream ID offset"); + +} // namespace detail + +class TileXRA5SDMABackend { +public: + TileXRA5SDMABackend(); + ~TileXRA5SDMABackend(); + TileXRA5SDMABackend(const TileXRA5SDMABackend&) = delete; + TileXRA5SDMABackend& operator=(const TileXRA5SDMABackend&) = delete; + + bool Init(int32_t deviceId); + bool Shutdown(); + GM_ADDR GetWorkspaceDev() const; + +private: + struct Impl; + std::unique_ptr impl_; +}; + +} // namespace TileXR + +#endif // TILEXR_SDMA_A5_BACKEND_H diff --git a/src/comm/sdma/tilexr_sdma_transport.cpp b/src/comm/sdma/tilexr_sdma_transport.cpp index 08ee0b6a..16fdb486 100644 --- a/src/comm/sdma/tilexr_sdma_transport.cpp +++ b/src/comm/sdma/tilexr_sdma_transport.cpp @@ -9,6 +9,8 @@ #include #include +#include "acl/acl_rt.h" +#include "sdma/tilexr_sdma_a5_backend.h" #include "tilexr_log.h" #include "tilexr_types.h" @@ -19,8 +21,10 @@ namespace TileXR { struct TileXRSDMATransport::Impl { + std::unique_ptr a5Backend; #if TILEXR_HAVE_PTO_SDMA pto::comm::sdma::SdmaWorkspaceManager workspaceManager; + bool ptoInitialized = false; #endif }; @@ -28,7 +32,7 @@ TileXRSDMATransport::TileXRSDMATransport() = default; TileXRSDMATransport::~TileXRSDMATransport() { - Shutdown(); + (void)Shutdown(); } bool TileXRSDMATransport::EnvEnabled() @@ -43,7 +47,10 @@ bool TileXRSDMATransport::EnvEnabled() int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) { - Shutdown(); + if (!Shutdown()) { + lastStatus_ = SDMAInitStatus::INIT_FAILED; + return TILEXR_SUCCESS; + } options_ = options; available_ = false; workspaceDev_ = nullptr; @@ -54,19 +61,53 @@ int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) return TILEXR_SUCCESS; } -#if TILEXR_HAVE_PTO_SDMA + const char* socName = aclrtGetSocName(); + const detail::SDMABackendKind backend = detail::ClassifySDMABackend(socName); + if (backend == detail::SDMABackendKind::UNSUPPORTED) { + lastStatus_ = SDMAInitStatus::PTO_UNAVAILABLE; + TILEXR_LOG(WARN) << "TileXR SDMA unsupported on SoC " + << (socName == nullptr ? "unknown" : socName); + return TILEXR_SUCCESS; + } + impl_.reset(new (std::nothrow) Impl()); if (impl_ == nullptr) { lastStatus_ = SDMAInitStatus::INIT_FAILED; - TILEXR_LOG(WARN) << "TileXR SDMA workspace manager allocation failed"; + TILEXR_LOG(WARN) << "TileXR SDMA implementation allocation failed"; + return TILEXR_SUCCESS; + } + + if (backend == detail::SDMABackendKind::A5_DIRECT) { + impl_->a5Backend.reset(new (std::nothrow) TileXRA5SDMABackend()); + if (impl_->a5Backend == nullptr) { + impl_.reset(); + lastStatus_ = SDMAInitStatus::INIT_FAILED; + return TILEXR_SUCCESS; + } + if (!impl_->a5Backend->Init(options_.devId)) { + lastStatus_ = SDMAInitStatus::INIT_FAILED; + TILEXR_LOG(WARN) << "TileXR A5 direct SDMA unavailable; communicator will continue without SDMA"; + return TILEXR_SUCCESS; + } + workspaceDev_ = impl_->a5Backend->GetWorkspaceDev(); + if (workspaceDev_ == nullptr) { + (void)Shutdown(); + lastStatus_ = SDMAInitStatus::NULL_WORKSPACE; + return TILEXR_SUCCESS; + } + available_ = true; + lastStatus_ = SDMAInitStatus::INITIALIZED; return TILEXR_SUCCESS; } + +#if TILEXR_HAVE_PTO_SDMA if (!impl_->workspaceManager.Init()) { lastStatus_ = SDMAInitStatus::INIT_FAILED; TILEXR_LOG(WARN) << "TileXR SDMA workspace manager init failed"; impl_.reset(); return TILEXR_SUCCESS; } + impl_->ptoInitialized = true; workspaceDev_ = static_cast(impl_->workspaceManager.GetWorkspaceAddr()); if (workspaceDev_ == nullptr) { lastStatus_ = SDMAInitStatus::NULL_WORKSPACE; @@ -83,20 +124,35 @@ int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) #else lastStatus_ = SDMAInitStatus::PTO_UNAVAILABLE; TILEXR_LOG(WARN) << "TileXR SDMA PTO headers unavailable at build time"; + impl_.reset(); return TILEXR_SUCCESS; #endif } -void TileXRSDMATransport::Shutdown() +bool TileXRSDMATransport::Shutdown() { -#if TILEXR_HAVE_PTO_SDMA + bool cleanupComplete = true; if (impl_ != nullptr) { - impl_->workspaceManager.Finalize(); - impl_.reset(); - } + if (impl_->a5Backend != nullptr) { + if (impl_->a5Backend->Shutdown()) { + impl_->a5Backend.reset(); + } else { + cleanupComplete = false; + } + } +#if TILEXR_HAVE_PTO_SDMA + if (impl_->ptoInitialized) { + impl_->workspaceManager.Finalize(); + impl_->ptoInitialized = false; + } #endif + if (cleanupComplete) { + impl_.reset(); + } + } available_ = false; workspaceDev_ = nullptr; + return cleanupComplete; } bool TileXRSDMATransport::IsAvailable() const diff --git a/src/comm/sdma/tilexr_sdma_transport.h b/src/comm/sdma/tilexr_sdma_transport.h index 2f8ed9f9..93f778c3 100644 --- a/src/comm/sdma/tilexr_sdma_transport.h +++ b/src/comm/sdma/tilexr_sdma_transport.h @@ -6,6 +6,7 @@ #ifndef TILEXR_SDMA_TRANSPORT_H #define TILEXR_SDMA_TRANSPORT_H +#include #include #include "comm_args.h" @@ -13,6 +14,35 @@ namespace TileXR { +namespace detail { + +enum class SDMABackendKind : uint32_t { + UNSUPPORTED = 0U, + PTO = 1U, + A5_DIRECT = 2U, +}; + +inline bool SDMASocHasPrefix(const char* socName, const char* prefix) +{ + return socName != nullptr && + std::strncmp(socName, prefix, std::strlen(prefix)) == 0; +} + +inline SDMABackendKind ClassifySDMABackend(const char* socName) +{ + if (SDMASocHasPrefix(socName, "Ascend950")) { + return SDMABackendKind::A5_DIRECT; + } + if (SDMASocHasPrefix(socName, "Ascend910B") || + SDMASocHasPrefix(socName, "Ascend910A") || + SDMASocHasPrefix(socName, "Ascend910_93")) { + return SDMABackendKind::PTO; + } + return SDMABackendKind::UNSUPPORTED; +} + +} // namespace detail + struct TileXRSDMATransportOptions { int devId = 0; }; @@ -25,7 +55,7 @@ class TileXRSDMATransport { TileXRSDMATransport& operator=(const TileXRSDMATransport&) = delete; int Init(const TileXRSDMATransportOptions& options); - void Shutdown(); + bool Shutdown(); bool IsAvailable() const; GM_ADDR GetWorkspaceDev() const; diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp index 69736d0d..72300c61 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -50,8 +50,6 @@ constexpr int TILEXR_INIT_TIMEOUT = 600; static map g_localPeerMemMap; static map g_devList; static std::mutex g_mtx; -static std::mutex g_sdmaMtx; -static bool g_sdmaUnavailable = false; // 如果是互联的链路，返回false； 对910B2C那些不互联的链路，返回true @@ -192,15 +190,14 @@ int TileXRComm::ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &st int TileXRComm::InitSDMA() { - { - lock_guard lock(g_sdmaMtx); - if (g_sdmaUnavailable) { - TILEXR_LOG(INFO) << "InitSDMA skipped after previous SDMA init failure"; - sdmaInitStatus_ = SDMAInitStatus::PTO_UNAVAILABLE; + if (sdmaTransport_ != nullptr) { + if (!sdmaTransport_->Shutdown()) { + TILEXR_LOG(WARN) << "TileXR previous SDMA resources are still pending cleanup"; + sdmaInitStatus_ = SDMAInitStatus::INIT_FAILED; return TILEXR_SUCCESS; } + sdmaTransport_.reset(); } - sdmaTransport_.reset(new (nothrow) TileXRSDMATransport()); if (sdmaTransport_ == nullptr) { TILEXR_LOG(WARN) << "TileXRSDMATransport allocation failed, SDMA disabled"; @@ -215,10 +212,10 @@ int TileXRComm::InitSDMA() if (ret != TILEXR_SUCCESS || !sdmaTransport_->IsAvailable()) { if (sdmaInitStatus_ != SDMAInitStatus::DISABLED_BY_ENV) { TILEXR_LOG(WARN) << "TileXR SDMA init unavailable, status " << static_cast(sdmaInitStatus_); - lock_guard lock(g_sdmaMtx); - g_sdmaUnavailable = true; } - sdmaTransport_.reset(); + if (sdmaTransport_->Shutdown()) { + sdmaTransport_.reset(); + } sdmaWorkspaceDev_ = nullptr; commArgs_.sdmaWorkspacePtr = nullptr; return TILEXR_SUCCESS; @@ -231,7 +228,9 @@ int TileXRComm::InitSDMA() commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; sdmaWorkspaceDev_ = nullptr; - sdmaTransport_.reset(); + if (sdmaTransport_->Shutdown()) { + sdmaTransport_.reset(); + } return TILEXR_SUCCESS; } @@ -242,16 +241,24 @@ int TileXRComm::InitSDMA() return TILEXR_SUCCESS; } -void TileXRComm::ResetSDMAState() +bool TileXRComm::ResetSDMAState() { commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; sdmaWorkspaceDev_ = nullptr; sdmaInitStatus_ = SDMAInitStatus::DISABLED_BY_ENV; if (sdmaTransport_ != nullptr) { - sdmaTransport_->Shutdown(); + if (!sdmaTransport_->Shutdown()) { + return false; + } sdmaTransport_.reset(); } + return true; +} + +bool TileXRComm::PrepareDestroy() +{ + return ResetSDMAState(); } bool TileXRComm::IsSDMAAvailable() const @@ -887,7 +894,7 @@ TileXRComm::~TileXRComm() FreePeerMem(commArgs_.dumpAddr); FreePeerMem(peerMem_[rank_]); FreePeerMem(commArgsPtr_); - ResetSDMAState(); + (void)ResetSDMAState(); } TileXRComm::TileXRComm(int rank, int rankSize) : rank_(rank), rankSize_(rankSize) diff --git a/src/comm/tilexr_comm.h b/src/comm/tilexr_comm.h index 8b8af295..5a4c5a42 100644 --- a/src/comm/tilexr_comm.h +++ b/src/comm/tilexr_comm.h @@ -51,6 +51,7 @@ class TileXRComm { bool IsSDMAAvailable() const; GM_ADDR GetSDMAWorkspacePtr() const; SDMAInitStatus GetSDMAInitStatus() const; + bool PrepareDestroy(); std::string PrintDFX(); friend class Lccl; friend class Lcoc; @@ -78,7 +79,7 @@ class TileXRComm { static int ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &state, void *userData); int InitSDMA(); int UpdateCommArgsDev(); - void ResetSDMAState(); + bool ResetSDMAState(); private: int rank_ = 0; // global rank id diff --git a/src/include/tilexr_sdma.h b/src/include/tilexr_sdma.h index db9fdf79..fd5b6035 100644 --- a/src/include/tilexr_sdma.h +++ b/src/include/tilexr_sdma.h @@ -19,7 +19,18 @@ #define TILEXR_HAVE_PTO_SDMA 0 #endif -#if TILEXR_ASCENDC_AICORE_COMPILE && defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA +#if TILEXR_ASCENDC_AICORE_COMPILE && \ + (((defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)) || \ + (defined(CATLASS_ARCH) && (CATLASS_ARCH == 3510))) || \ + defined(TILEXR_SDMA_FORCE_A5)) +#define TILEXR_SDMA_A5_AICORE_COMPILE 1 +#include "tilexr_sdma_a5.h" +#else +#define TILEXR_SDMA_A5_AICORE_COMPILE 0 +#endif + +#if TILEXR_ASCENDC_AICORE_COMPILE && !TILEXR_SDMA_A5_AICORE_COMPILE && \ + defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA #include "tilexr_sdma_compat.h" #endif @@ -47,7 +58,15 @@ __aicore__ inline uint64_t SDMACopyNbi( uint64_t bytes, uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP) { -#if defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA +#if TILEXR_SDMA_A5_AICORE_COMPILE + if (!SDMAEnabled(args) || dst == nullptr || src == nullptr || bytes == 0) { + return 0; + } + const uint32_t resolvedGroup = SDMAResolveChannelGroup(channelGroupIdx); + return detail::A5SdmaCopyNbi( + reinterpret_cast<__gm__ uint8_t*>(args->sdmaWorkspacePtr), + dst, src, bytes, resolvedGroup); +#elif defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA if (!SDMAEnabled(args) || dst == nullptr || src == nullptr || bytes == 0) { return 0; } @@ -75,7 +94,18 @@ __aicore__ inline bool SDMAWait( uint64_t eventHandle, uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP) { -#if defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA +#if TILEXR_SDMA_A5_AICORE_COMPILE + if (eventHandle == 0) { + return true; + } + if (!SDMAEnabled(args)) { + return false; + } + const uint32_t resolvedGroup = SDMAResolveChannelGroup(channelGroupIdx); + return detail::A5SdmaWaitEvent( + reinterpret_cast<__gm__ uint8_t*>(args->sdmaWorkspacePtr), + eventHandle, resolvedGroup); +#elif defined(TILEXR_HAVE_PTO_SDMA) && TILEXR_HAVE_PTO_SDMA if (eventHandle == 0) { return true; } diff --git a/src/include/tilexr_sdma_a5.h b/src/include/tilexr_sdma_a5.h new file mode 100644 index 00000000..65c19c6a --- /dev/null +++ b/src/include/tilexr_sdma_a5.h @@ -0,0 +1,206 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_A5_H +#define TILEXR_SDMA_A5_H + +#include "kernel_operator.h" +#include "tilexr_sdma_a5_types.h" + +namespace TileXR { +namespace detail { + +__aicore__ inline void A5SdmaCleanCacheLine(__gm__ uint8_t* address) +{ + AscendC::GlobalTensor line; + line.SetGlobalBuffer(address); + __asm__ __volatile__(""); + AscendC::DataCacheCleanAndInvalid< + uint8_t, AscendC::CacheLine::SINGLE_CACHE_LINE, + AscendC::DcciDst::CACHELINE_OUT>(line); + __asm__ __volatile__(""); +} + +__aicore__ inline void A5SdmaZeroSqe(__gm__ A5SdmaSqe* sqe) +{ + __gm__ uint32_t* words = reinterpret_cast<__gm__ uint32_t*>(sqe); + for (uint32_t index = 0U; index < TILEXR_SDMA_A5_SQE_BYTES / sizeof(uint32_t); ++index) { + words[index] = 0U; + } +} + +__aicore__ inline void A5SdmaBuildSqe(__gm__ A5SdmaSqe* sqe, + uint32_t streamId, + uint32_t taskId, + uint64_t source, + uint64_t destination, + uint32_t bytes) +{ + A5SdmaZeroSqe(sqe); + sqe->typeLockUnlock = static_cast(TILEXR_SDMA_A5_SQE_TYPE); + sqe->rtStreamId = static_cast(streamId); + sqe->taskId = static_cast(taskId); + sqe->kernelCredit = static_cast(TILEXR_SDMA_A5_KERNEL_CREDIT); + sqe->transferFlags = (1U << 8U) | (1U << 9U) | (1U << 10U) | (1U << 11U); + sqe->mpamFlags = static_cast(TILEXR_SDMA_A5_QOS << 3U); + sqe->srcAddressLow = static_cast(source & 0xFFFFFFFFULL); + sqe->srcAddressHigh = static_cast(source >> 32U); + sqe->dstAddressLow = static_cast(destination & 0xFFFFFFFFULL); + sqe->dstAddressHigh = static_cast(destination >> 32U); + sqe->length = bytes; +} + +__aicore__ inline void A5SdmaRingDoorbell(uint64_t address, uint32_t tail) +{ + __ubuf__ uint32_t value[8]; + value[0] = tail; + pipe_barrier(PIPE_ALL); + copy_ubuf_to_gm_align_v2(reinterpret_cast<__gm__ uint32_t*>(address), value, + 0, 1, sizeof(uint32_t), 0, 0, 0); + set_flag(PIPE_MTE3, PIPE_MTE2, static_cast(0)); + wait_flag(PIPE_MTE3, PIPE_MTE2, static_cast(0)); +} + +__aicore__ inline uint32_t A5SdmaReadCompletion(__gm__ A5SdmaCompletionLine* completion) +{ + __ubuf__ uint32_t line[TILEXR_SDMA_A5_COMPLETION_BYTES / sizeof(uint32_t)]; + copy_gm_to_ubuf_align_v2( + line, reinterpret_cast<__gm__ uint32_t*>(completion), + 0, 1, TILEXR_SDMA_A5_COMPLETION_BYTES, 0, 0, false, 0, 0, 0); + set_flag(PIPE_MTE2, PIPE_S, static_cast(0)); + wait_flag(PIPE_MTE2, PIPE_S, static_cast(0)); + return line[0]; +} + +__aicore__ inline bool A5SdmaWorkspaceValid(const __gm__ A5SdmaWorkspace* workspace) +{ + return workspace != nullptr && + workspace->header.magic == TILEXR_SDMA_A5_WORKSPACE_MAGIC && + workspace->header.abiVersion == TILEXR_SDMA_A5_ABI_VERSION && + workspace->header.backendKind == TILEXR_SDMA_A5_BACKEND_KIND && + workspace->header.channelCount == TILEXR_SDMA_A5_CHANNEL_COUNT && + workspace->header.sqeSize == TILEXR_SDMA_A5_SQE_BYTES && + workspace->header.channelStride == sizeof(A5SdmaChannel) && + workspace->header.workspaceSize == sizeof(A5SdmaWorkspace); +} + +__aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint64_t bytes, + uint32_t channelIndex) +{ + __gm__ A5SdmaWorkspace* workspace = + reinterpret_cast<__gm__ A5SdmaWorkspace*>(workspaceAddress); + if (!A5SdmaWorkspaceValid(workspace) || destination == nullptr || source == nullptr || + !A5SdmaTransferLengthValid(bytes) || !A5SdmaChannelValid(channelIndex)) { + return 0ULL; + } + + __gm__ A5SdmaChannel* channel = &workspace->channels[channelIndex]; + if (channel->sqBase == 0U || channel->rtsqAddress == 0U || + channel->completionPayloadAddress == 0U || channel->completionRecordAddress == 0U || + channel->rtsqLength < sizeof(uint32_t) || + !A5SdmaQueueHasCapacity(channel->head, channel->tail, channel->depth) || + channel->streamId > 0xFFFFU) { + return 0ULL; + } + if (AscendC::AtomicCas(&channel->outstanding, 0U, 1U) != 0U) { + return 0ULL; + } + + const uint32_t generation = A5SdmaNextGeneration(channel->generation); + const uint32_t dataIndex = channel->tail; + const uint32_t completionIndex = (dataIndex + 1U) % channel->depth; + const uint32_t newTail = A5SdmaAdvanceTail(dataIndex, channel->depth); + __gm__ A5SdmaCompletionLine* payload = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( + channel->completionPayloadAddress); + __gm__ A5SdmaCompletionLine* completion = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( + channel->completionRecordAddress); + __gm__ uint32_t* payloadWords = reinterpret_cast<__gm__ uint32_t*>(payload); + __gm__ uint32_t* completionWords = reinterpret_cast<__gm__ uint32_t*>(completion); + for (uint32_t index = 0U; index < TILEXR_SDMA_A5_COMPLETION_BYTES / sizeof(uint32_t); ++index) { + payloadWords[index] = 0U; + completionWords[index] = 0U; + } + payload->generation = generation; + + __gm__ A5SdmaSqe* sqBase = reinterpret_cast<__gm__ A5SdmaSqe*>(channel->sqBase); + A5SdmaBuildSqe(sqBase + dataIndex, channel->streamId, channel->taskId, + reinterpret_cast(source), + reinterpret_cast(destination), + static_cast(bytes)); + A5SdmaBuildSqe(sqBase + completionIndex, channel->streamId, channel->taskId + 1U, + channel->completionPayloadAddress, + channel->completionRecordAddress, + TILEXR_SDMA_A5_COMPLETION_BYTES); + + channel->generation = generation; + channel->tail = newTail; + channel->taskId = A5SdmaAdvanceTaskId(channel->taskId); + pipe_barrier(PIPE_ALL); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(payload)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(completion)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + dataIndex)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + completionIndex)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel) + 64U); + pipe_barrier(PIPE_ALL); + dsb(DSB_DDR); + A5SdmaRingDoorbell(channel->rtsqAddress, newTail); + return A5SdmaEncodeEvent(channelIndex, generation); +} + +__aicore__ inline bool A5SdmaWaitEvent(__gm__ uint8_t* workspaceAddress, + uint64_t event, + uint32_t expectedChannel) +{ + uint32_t channelIndex = 0U; + uint32_t generation = 0U; + if (!A5SdmaDecodeEvent(event, channelIndex, generation) || channelIndex != expectedChannel) { + return false; + } + __gm__ A5SdmaWorkspace* workspace = + reinterpret_cast<__gm__ A5SdmaWorkspace*>(workspaceAddress); + if (!A5SdmaWorkspaceValid(workspace)) { + return false; + } + __gm__ A5SdmaChannel* channel = &workspace->channels[channelIndex]; + if (channel->generation != generation || + AscendC::AtomicCas(&channel->outstanding, 1U, 1U) != 1U || + channel->completionRecordAddress == 0U) { + return false; + } + __gm__ A5SdmaCompletionLine* completion = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( + channel->completionRecordAddress); + bool completed = false; + for (uint32_t poll = 0U; poll < TILEXR_SDMA_A5_WAIT_MAX_POLLS; ++poll) { + if (A5SdmaReadCompletion(completion) == generation) { + completed = true; + break; + } + } + if (!completed) { + return false; + } + dsb(DSB_DDR); + if (channel->generation != generation) { + return false; + } + channel->head = channel->tail; + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel)); + pipe_barrier(PIPE_ALL); + dsb(DSB_DDR); + if (channel->generation != generation || + AscendC::AtomicCas(&channel->outstanding, 1U, 0U) != 1U) { + return false; + } + return true; +} + +} // namespace detail +} // namespace TileXR + +#endif // TILEXR_SDMA_A5_H diff --git a/src/include/tilexr_sdma_a5_types.h b/src/include/tilexr_sdma_a5_types.h new file mode 100644 index 00000000..2a75a084 --- /dev/null +++ b/src/include/tilexr_sdma_a5_types.h @@ -0,0 +1,200 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_A5_TYPES_H +#define TILEXR_SDMA_A5_TYPES_H + +#include +#include + +namespace TileXR { +namespace detail { + +#if defined(__CCE__) && defined(__CCE_IS_AICORE__) +#define TILEXR_SDMA_A5_HOST_DEVICE_INLINE __aicore__ inline +#else +#define TILEXR_SDMA_A5_HOST_DEVICE_INLINE inline +#endif + +constexpr uint32_t TILEXR_SDMA_A5_WORKSPACE_MAGIC = 0x41355344U; // "A5SD" +constexpr uint16_t TILEXR_SDMA_A5_ABI_VERSION = 1U; +constexpr uint16_t TILEXR_SDMA_A5_BACKEND_KIND = 2U; +constexpr uint32_t TILEXR_SDMA_A5_CHANNEL_COUNT = 48U; +constexpr uint32_t TILEXR_SDMA_A5_SQE_BYTES = 64U; +constexpr uint32_t TILEXR_SDMA_A5_COMPLETION_BYTES = 64U; +constexpr uint64_t TILEXR_SDMA_A5_MAX_TRANSFER_BYTES = 0xFFFFFFFFULL; +constexpr uint32_t TILEXR_SDMA_A5_SQE_TYPE = 11U; +constexpr uint32_t TILEXR_SDMA_A5_KERNEL_CREDIT = 254U; +constexpr uint32_t TILEXR_SDMA_A5_QOS = 6U; +constexpr uint32_t TILEXR_SDMA_A5_WAIT_MAX_POLLS = 1000000U; + +constexpr uint64_t TILEXR_SDMA_A5_EVENT_MAGIC = 0xA5D5ULL; +constexpr uint32_t TILEXR_SDMA_A5_EVENT_MAGIC_SHIFT = 48U; +constexpr uint32_t TILEXR_SDMA_A5_EVENT_CHANNEL_SHIFT = 40U; +constexpr uint64_t TILEXR_SDMA_A5_EVENT_RESERVED_MASK = 0x000000FF00000000ULL; +constexpr uint64_t TILEXR_SDMA_A5_EVENT_GENERATION_MASK = 0xFFFFFFFFULL; + +struct alignas(64) A5SdmaWorkspaceHeader { + uint32_t magic; + uint16_t abiVersion; + uint16_t backendKind; + uint32_t channelCount; + uint32_t sqeSize; + uint32_t channelStride; + uint32_t workspaceSize; + uint32_t maxTransferBytes; + uint32_t reserved[9]; +}; + +struct alignas(64) A5SdmaChannel { + uint64_t sqBase; + uint64_t rtsqAddress; + uint64_t completionPayloadAddress; + uint64_t completionRecordAddress; + uint32_t depth; + uint32_t head; + uint32_t tail; + uint32_t taskId; + uint32_t rtsqLength; + uint32_t streamId; + uint32_t sqId; + uint32_t cqId; + uint32_t logicalCqId; + uint32_t physicalDieId; + uint32_t generation; + uint32_t reserved0[13]; + uint32_t outstanding; + uint32_t reserved1[15]; +}; + +struct alignas(64) A5SdmaCompletionLine { + uint32_t generation; + uint32_t reserved[15]; +}; + +struct A5SdmaSqe { + uint8_t typeLockUnlock; + uint8_t controlFlags; + uint16_t numBlocks; + uint16_t rtStreamId; + uint16_t taskId; + uint32_t reserved0; + uint16_t reserved1; + uint8_t kernelCredit; + uint8_t reserved2; + uint32_t transferFlags; + uint16_t sqeId; + uint8_t mpamPartId; + uint8_t mpamFlags; + uint16_t srcStreamId; + uint16_t srcSubStreamId; + uint16_t dstStreamId; + uint16_t dstSubStreamId; + uint32_t srcAddressLow; + uint32_t srcAddressHigh; + uint32_t dstAddressLow; + uint32_t dstAddressHigh; + uint32_t length; + uint32_t srcOffsetLow; + uint32_t dstOffsetLow; + uint16_t srcOffsetHigh; + uint16_t dstOffsetHigh; +}; + +struct alignas(64) A5SdmaWorkspace { + A5SdmaWorkspaceHeader header; + A5SdmaChannel channels[TILEXR_SDMA_A5_CHANNEL_COUNT]; + A5SdmaCompletionLine completionPayloads[TILEXR_SDMA_A5_CHANNEL_COUNT]; + A5SdmaCompletionLine completionRecords[TILEXR_SDMA_A5_CHANNEL_COUNT]; +}; + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaChannelValid(uint32_t channel) +{ + return channel < TILEXR_SDMA_A5_CHANNEL_COUNT; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaTransferLengthValid(uint64_t bytes) +{ + return bytes != 0U && bytes <= TILEXR_SDMA_A5_MAX_TRANSFER_BYTES; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueStateValid(uint32_t tail, uint32_t depth) +{ + return depth >= 3U && tail < depth; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTail(uint32_t tail, uint32_t depth) +{ + return (tail + 2U) % depth; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaQueueDistance( + uint32_t head, uint32_t tail, uint32_t depth) +{ + return (tail + depth - head) % depth; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasCapacity( + uint32_t head, uint32_t tail, uint32_t depth) +{ + if (depth < 3U || head >= depth || tail >= depth) { + return false; + } + return A5SdmaQueueDistance(head, tail, depth) <= depth - 3U; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskId(uint32_t taskId) +{ + return (taskId + 2U) & 0xFFFFU; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaNextGeneration(uint32_t generation) +{ + return generation == 0xFFFFFFFFU ? 1U : generation + 1U; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint64_t A5SdmaEncodeEvent( + uint32_t channel, uint32_t generation) +{ + return !A5SdmaChannelValid(channel) || generation == 0U + ? 0ULL + : (TILEXR_SDMA_A5_EVENT_MAGIC << TILEXR_SDMA_A5_EVENT_MAGIC_SHIFT) | + (static_cast(channel) << TILEXR_SDMA_A5_EVENT_CHANNEL_SHIFT) | + static_cast(generation); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaDecodeEvent( + uint64_t event, uint32_t& channel, uint32_t& generation) +{ + if ((event >> TILEXR_SDMA_A5_EVENT_MAGIC_SHIFT) != TILEXR_SDMA_A5_EVENT_MAGIC || + (event & TILEXR_SDMA_A5_EVENT_RESERVED_MASK) != 0U) { + return false; + } + channel = static_cast((event >> TILEXR_SDMA_A5_EVENT_CHANNEL_SHIFT) & 0xFFULL); + generation = static_cast(event & TILEXR_SDMA_A5_EVENT_GENERATION_MASK); + return A5SdmaChannelValid(channel) && generation != 0U; +} + +static_assert(sizeof(A5SdmaWorkspaceHeader) == 64U, "A5 workspace header must be 64 bytes"); +static_assert(sizeof(A5SdmaChannel) == 192U, "A5 channel ABI must be 192 bytes"); +static_assert(sizeof(A5SdmaCompletionLine) == 64U, "A5 completion line must be 64 bytes"); +static_assert(sizeof(A5SdmaSqe) == TILEXR_SDMA_A5_SQE_BYTES, "A5 SQE must be 64 bytes"); +static_assert(alignof(A5SdmaWorkspace) == 64U, "A5 workspace must be cache-line aligned"); +static_assert(sizeof(A5SdmaWorkspace) % 64U == 0U, "A5 workspace size must be cache-line aligned"); +static_assert(offsetof(A5SdmaChannel, sqBase) == 0U, "unexpected A5 SQ base offset"); +static_assert(offsetof(A5SdmaChannel, completionPayloadAddress) == 16U, + "unexpected A5 completion payload offset"); +static_assert(offsetof(A5SdmaChannel, generation) == 72U, "unexpected A5 generation offset"); +static_assert(offsetof(A5SdmaChannel, outstanding) == 128U, "unexpected A5 outstanding offset"); +static_assert(offsetof(A5SdmaSqe, srcAddressLow) == 32U, "unexpected A5 source address offset"); +static_assert(offsetof(A5SdmaSqe, length) == 48U, "unexpected A5 length offset"); +static_assert(offsetof(A5SdmaWorkspace, channels) == 64U, "unexpected A5 channel array offset"); + +#undef TILEXR_SDMA_A5_HOST_DEVICE_INLINE + +} // namespace detail +} // namespace TileXR + +#endif // TILEXR_SDMA_A5_TYPES_H diff --git a/tests/ci/test_control_contract.py b/tests/ci/test_control_contract.py index adec07f3..ee39d880 100644 --- a/tests/ci/test_control_contract.py +++ b/tests/ci/test_control_contract.py @@ -80,11 +80,14 @@ def test_build_manifest_enables_and_runs_all_non_hardware_coverage(self): "test_tilexr_udma_demo_sources", "test_tilexr_udma_source_guard", "test_tilexr_sdma_metadata", + "test_tilexr_sdma_a5_validation", "test_tilexr_sdma_api_invalid", "test_tilexr_sdma_transport_disabled", "test_tilexr_sdma_comm_wiring", "test_tilexr_sdma_source_guard", "test_tilexr_sdma_header_compile", + "-DTILEXR_SDMA_DEMO_SOC_TYPE=Ascend950", + "--target tilexr_sdma_demo_kernel", "test_tilexr_memory_demo_sources", "test_tilexr_ep_layout", "test_tilexr_ep_api_sources", @@ -194,9 +197,9 @@ def test_provisioning_enforces_the_documented_minimum_driver_version(self): self.assertIn("version_at_least()", common) for text in [cann, verify]: - self.assertIn('version_at_least "${driver_version}" 25.5.0', text) - self.assertNotIn('"${driver_version}" != 25.5.0', text) - self.assertNotIn("grep -Fx Version=25.5.0", cann) + self.assertIn('version_at_least "${driver_version}" 25.1.rc1', text) + self.assertNotIn('"${driver_version}" != 25.1.rc1', text) + self.assertNotIn("grep -Fx Version=25.1.rc1", cann) def test_cann_paths_use_installer_required_permissions(self): common = self.read("scripts/ci/provision/common.sh") @@ -493,6 +496,8 @@ def test_build_manifest_captures_configuration_and_build_logs(self): "comm-build", "udma-build", "sdma-build", + "sdma-a5-configure", + "sdma-a5-build", "ep-build", "memory-configure", "memory-build", @@ -583,7 +588,7 @@ def test_runner_dry_run_does_not_require_registration_token_input(self): def test_version_at_least_compares_numeric_release_components(self): common = ROOT / "scripts/ci/provision/common.sh" - harness = 'source "$1"; version_at_least "$2" 25.5.0' + harness = 'source "$1"; version_at_least "$2" 25.1.rc1' def supported(version): return subprocess.run( @@ -593,10 +598,16 @@ def supported(version): stderr=subprocess.PIPE, ) - for version in ["25.5.0", "25.5.1", "25.6.0", "26.0.0"]: + for version in [ + "25.1.rc1", "25.1.RC1", "25.1.rc1.b188", "25.1.rc2", + "25.1.0", "25.2.rc1", "25.5.0", "26.0.0", + ]: with self.subTest(version=version): self.assertEqual(0, supported(version).returncode) - for version in ["25.4.99", "24.99.99", "25.5", "25.5.0.1", "25.5.RC1", ""]: + for version in [ + "25.1.rc0", "25.0.99", "24.99.99", "25.1", "25.1.rc", + "25.1.rc1.bad", "25.1.0.1", "", + ]: with self.subTest(version=version): self.assertNotEqual(0, supported(version).returncode) diff --git a/tests/sdma/CMakeLists.txt b/tests/sdma/CMakeLists.txt index 17c85a30..b0c332bb 100644 --- a/tests/sdma/CMakeLists.txt +++ b/tests/sdma/CMakeLists.txt @@ -5,6 +5,14 @@ set(CMAKE_CXX_STANDARD 14) set(CMAKE_CXX_STANDARD_REQUIRED ON) option(BUILD_TILEXR_SDMA_DEMO "Build TileXR SDMA data-plane demo with Ascend C kernel" ON) set(TILEXR_SDMA_DEMO_SOC_TYPE "Ascend910B" CACHE STRING "SOC type used for the TileXR SDMA demo kernel") +set_property(CACHE TILEXR_SDMA_DEMO_SOC_TYPE PROPERTY STRINGS Ascend910B Ascend950) +if(NOT TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend910B" AND + NOT TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend950") + message(FATAL_ERROR + "Unsupported TILEXR_SDMA_DEMO_SOC_TYPE=${TILEXR_SDMA_DEMO_SOC_TYPE}; " + "expected Ascend910B or Ascend950") +endif() +message(STATUS "TILEXR_SDMA_DEMO_SOC_TYPE: ${TILEXR_SDMA_DEMO_SOC_TYPE}") set(ASCEND_HOME_PATH $ENV{ASCEND_HOME_PATH}) set(ARCH $ENV{ARCH}) @@ -77,6 +85,18 @@ target_link_libraries(test_tilexr_sdma_api_invalid add_executable(test_tilexr_sdma_transport_disabled unit/test_tilexr_sdma_transport_disabled.cpp ${TILEXR_ROOT}/src/comm/sdma/tilexr_sdma_transport.cpp + ${TILEXR_ROOT}/src/comm/sdma/tilexr_sdma_a5_backend.cpp +) + +add_executable(test_tilexr_sdma_a5_validation + unit/test_tilexr_sdma_a5_validation.cpp +) +target_include_directories(test_tilexr_sdma_a5_validation PRIVATE + ${TILEXR_ROOT}/src/comm + ${TILEXR_ROOT}/src/include + ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/ + ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ + ${ASCEND_DRIVER_PATH}/kernel/inc ) target_include_directories(test_tilexr_sdma_transport_disabled PRIVATE ${TILEXR_ROOT}/src/comm @@ -86,6 +106,12 @@ target_include_directories(test_tilexr_sdma_transport_disabled PRIVATE ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ ${ASCEND_DRIVER_PATH}/kernel/inc ) +target_link_libraries(test_tilexr_sdma_transport_disabled + ascendcl + runtime + ascend_hal + dl +) add_executable(test_tilexr_sdma_comm_wiring unit/test_tilexr_sdma_comm_wiring.cpp @@ -132,13 +158,16 @@ set(INSTALL_TARGETS test_tilexr_sdma_metadata test_tilexr_sdma_api_invalid test_tilexr_sdma_transport_disabled + test_tilexr_sdma_a5_validation test_tilexr_sdma_comm_wiring test_tilexr_sdma_source_guard test_tilexr_sdma_header_compile test_tilexr_sdma_disabled_comm ) -if(BUILD_TILEXR_SDMA_DEMO AND NOT TILEXR_TEST_HAVE_PTO_SDMA) +if(BUILD_TILEXR_SDMA_DEMO AND + TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend910B" AND + NOT TILEXR_TEST_HAVE_PTO_SDMA) message(WARNING "installed tilexr_sdma_config.h disables PTO SDMA; skip tilexr_sdma_demo") elseif(BUILD_TILEXR_SDMA_DEMO) find_program(BISHENG_EXECUTABLE bisheng) @@ -149,10 +178,14 @@ elseif(BUILD_TILEXR_SDMA_DEMO) set(TILEXR_SDMA_NPU_ARCH "dav-3510") set(TILEXR_SDMA_AICORE_ARCH "--cce-aicore-arch=dav-c310-vec") set(TILEXR_SDMA_CATLASS_ARCH "3510") - else() + set(TILEXR_SDMA_KERNEL_DEFINITIONS -DTILEXR_HAVE_PTO_SDMA=0) + set(TILEXR_SDMA_KERNEL_EXTRA_LIBS "") + elseif(TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend910B") set(TILEXR_SDMA_NPU_ARCH "dav-2201") set(TILEXR_SDMA_AICORE_ARCH "--cce-aicore-arch=dav-c220-vec") set(TILEXR_SDMA_CATLASS_ARCH "2201") + set(TILEXR_SDMA_KERNEL_DEFINITIONS -DTILEXR_HAVE_PTO_SDMA=1) + set(TILEXR_SDMA_KERNEL_EXTRA_LIBS -lnnopbase) endif() execute_process( @@ -209,11 +242,12 @@ elseif(BUILD_TILEXR_SDMA_DEMO) COMMAND ${BISHENG_EXECUTABLE} ${TILEXR_SDMA_KERNEL_COMPILE_OPTIONS} -std=gnu++17 + -O2 -fPIC -shared ${TILEXR_SDMA_KERNEL_LINK_OPTIONS} -DCATLASS_ARCH=${TILEXR_SDMA_CATLASS_ARCH} - -DTILEXR_HAVE_PTO_SDMA=1 + ${TILEXR_SDMA_KERNEL_DEFINITIONS} ${TILEXR_SDMA_DEMO_KERNEL_INCLUDES} "${CMAKE_CURRENT_SOURCE_DIR}/demo/tilexr_sdma_demo_kernel.cpp" -L${ASCEND_DRIVER_PATH}/lib64 @@ -225,12 +259,14 @@ elseif(BUILD_TILEXR_SDMA_DEMO) -lstdc++ -lm -ldl - -lnnopbase + ${TILEXR_SDMA_KERNEL_EXTRA_LIBS} -lpthread -o "${TILEXR_SDMA_DEMO_KERNEL_SO}" DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/demo/tilexr_sdma_demo_kernel.cpp" "${TILEXR_ROOT}/src/include/tilexr_sdma.h" + "${TILEXR_ROOT}/src/include/tilexr_sdma_a5.h" + "${TILEXR_ROOT}/src/include/tilexr_sdma_a5_types.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_compat.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_types.h" VERBATIM @@ -241,6 +277,11 @@ elseif(BUILD_TILEXR_SDMA_DEMO) add_executable(tilexr_sdma_demo demo/tilexr_sdma_demo.cpp ) + if(TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend950") + target_compile_definitions(tilexr_sdma_demo PRIVATE TILEXR_SDMA_DEMO_A5=1) + else() + target_compile_definitions(tilexr_sdma_demo PRIVATE TILEXR_SDMA_DEMO_A5=0) + endif() add_dependencies(tilexr_sdma_demo tilexr_sdma_demo_kernel) target_include_directories(tilexr_sdma_demo PRIVATE ${TILEXR_ROOT}/install/include diff --git a/tests/sdma/build.sh b/tests/sdma/build.sh index 9f83acd5..8839634f 100755 --- a/tests/sdma/build.sh +++ b/tests/sdma/build.sh @@ -4,6 +4,7 @@ set -euo pipefail SCRIPT_DIR=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd) TILEXR_ROOT=$(cd "${SCRIPT_DIR}/../.." && pwd) CANN_HOME="${1:-${ASCEND_HOME_PATH:-}}" +SDMA_SOC_TYPE="${2:-${TILEXR_SDMA_DEMO_SOC_TYPE:-Ascend910B}}" if [ -z "${CANN_HOME}" ]; then set +u @@ -39,7 +40,7 @@ cmake -S "${TILEXR_ROOT}" -B "${ROOT_BUILD}" \ -DTILEXR_BUILD_TESTS=OFF cmake --build "${ROOT_BUILD}" --target install -j"$(nproc)" -for header in comm_args.h tilexr_sdma_types.h tilexr_sdma_config.h; do +for header in comm_args.h tilexr_sdma_types.h tilexr_sdma_a5_types.h tilexr_sdma_a5.h tilexr_sdma_config.h; do if [ ! -f "${ROOT_INSTALL}/include/${header}" ]; then echo "ERROR: expected installed header missing: ${ROOT_INSTALL}/include/${header}" >&2 exit 1 @@ -55,7 +56,9 @@ fi cmake -S "${SCRIPT_DIR}" -B "${TEST_BUILD}" \ -DCMAKE_INSTALL_PREFIX="${TEST_INSTALL}" \ + -DTILEXR_SDMA_DEMO_SOC_TYPE="${SDMA_SOC_TYPE}" \ ${DEMO_OPTION} cmake --build "${TEST_BUILD}" --target install -j"$(nproc)" +echo "SDMA demo SOC type: ${SDMA_SOC_TYPE}" echo "SDMA tests installed to ${TEST_INSTALL}/bin" diff --git a/tests/sdma/demo/run_tilexr_sdma_demo.sh b/tests/sdma/demo/run_tilexr_sdma_demo.sh index d9be2a49..59d5e6f0 100755 --- a/tests/sdma/demo/run_tilexr_sdma_demo.sh +++ b/tests/sdma/demo/run_tilexr_sdma_demo.sh @@ -11,6 +11,10 @@ DEVICE_ID="${2:-0}" shift $(( $# > 0 ? 1 : 0 )) || true shift $(( $# > 0 ? 1 : 0 )) || true SIZES=("$@") +FIRST_CHANNEL="${TILEXR_SDMA_DEMO_CHANNEL:-0}" +BLOCKS="${TILEXR_SDMA_DEMO_BLOCKS:-1}" +ITERATIONS="${TILEXR_SDMA_DEMO_ITERATIONS:-1}" +REPEATS="${TILEXR_SDMA_DEMO_REPEATS:-1}" if [ ${#SIZES[@]} -eq 0 ]; then SIZES=(64 4096 1048576) fi @@ -80,11 +84,17 @@ echo " TileXR SDMA Demo" echo "==========================================" echo "CANN_HOME: ${CANN_HOME}" echo "DEVICE_ID: ${DEVICE_ID}" +echo "FIRST_CHANNEL: ${FIRST_CHANNEL}" +echo "BLOCKS: ${BLOCKS}" +echo "ITERATIONS: ${ITERATIONS}" +echo "REPEATS: ${REPEATS}" echo "Sizes: ${SIZES[*]}" echo "Binary: ${bin}" echo "==========================================" -for bytes in "${SIZES[@]}"; do - echo "---- bytes=${bytes} ----" - "${bin}" "${bytes}" +for ((repeat = 1; repeat <= REPEATS; ++repeat)); do + for bytes in "${SIZES[@]}"; do + echo "---- repeat=${repeat} bytes=${bytes} ----" + "${bin}" "${bytes}" "${FIRST_CHANNEL}" "${BLOCKS}" "${ITERATIONS}" + done done diff --git a/tests/sdma/demo/tilexr_sdma_demo.cpp b/tests/sdma/demo/tilexr_sdma_demo.cpp index cd0f0620..a7087cbc 100644 --- a/tests/sdma/demo/tilexr_sdma_demo.cpp +++ b/tests/sdma/demo/tilexr_sdma_demo.cpp @@ -15,6 +15,7 @@ #include "acl/acl.h" #include "tilexr_api.h" +#include "tilexr_sdma_a5_types.h" #include "tilexr_sdma_types.h" #include "tilexr_types.h" @@ -25,13 +26,20 @@ extern "C" void launch_tilexr_sdma_copy( GM_ADDR dst, GM_ADDR src, GM_ADDR debug, - uint32_t bytes); + uint32_t bytes, + uint32_t firstChannel, + uint32_t iterations); namespace { constexpr uint32_t kDefaultBytes = 4096; constexpr uint32_t kAlignmentBytes = 64; -constexpr size_t kDebugWords = 6; +constexpr size_t kDebugWordsPerBlock = 16; constexpr int kDeviceId = 0; +constexpr int32_t kStreamTimeoutMs = 60000; + +#ifndef TILEXR_SDMA_DEMO_A5 +#define TILEXR_SDMA_DEMO_A5 0 +#endif bool CheckAcl(const std::string& label, aclError ret) { @@ -69,6 +77,22 @@ bool ParseBytes(const char* text, uint32_t* bytes) return true; } +bool ParseUint32(const char* text, uint32_t* value) +{ + if (text == nullptr || value == nullptr || text[0] == '\0') { + return false; + } + errno = 0; + char* end = nullptr; + const unsigned long long parsed = std::strtoull(text, &end, 10); + if (errno != 0 || end == text || *end != '\0' || + parsed > std::numeric_limits::max()) { + return false; + } + *value = static_cast(parsed); + return true; +} + void FillPattern(std::vector& data) { for (size_t i = 0; i < data.size(); ++i) { @@ -120,28 +144,24 @@ bool CopyDeviceToHost(void* dst, size_t dstSize, const void* src, size_t srcSize aclrtMemcpy(dst, dstSize, src, srcSize, ACL_MEMCPY_DEVICE_TO_HOST)); } -bool VerifyDebug(const std::vector& debug, uint32_t bytes) +bool VerifyDebug(const std::vector& debug, uint32_t bytes, + uint32_t firstChannel, uint32_t blocks, uint32_t iterations) { bool ok = true; - if (debug[0] != TileXR::TILEXR_SDMA_DEMO_MAGIC) { - std::cerr << "ERROR: debug magic mismatch: got " << debug[0] << std::endl; - ok = false; - } - if (debug[2] != static_cast(bytes)) { - std::cerr << "ERROR: debug bytes mismatch: got " << debug[2] << std::endl; - ok = false; - } - if (debug[3] != 1) { - std::cerr << "ERROR: SDMA was not enabled in kernel debug word" << std::endl; - ok = false; - } - if (debug[4] != 1) { - std::cerr << "ERROR: SDMA event was not posted" << std::endl; - ok = false; - } - if (debug[5] != 1) { - std::cerr << "ERROR: SDMA wait did not report success" << std::endl; - ok = false; + for (uint32_t block = 0U; block < blocks; ++block) { + const size_t base = static_cast(block) * kDebugWordsPerBlock; + if (debug[base] != TileXR::TILEXR_SDMA_DEMO_MAGIC || + debug[base + 1U] != static_cast(block) || + debug[base + 2U] != static_cast(bytes) || + debug[base + 3U] != 1 || debug[base + 4U] != 1 || + debug[base + 5U] != 1 || + debug[base + 6U] != static_cast(firstChannel + block) || + debug[base + 7U] != static_cast(iterations) || + debug[base + 8U] != (TILEXR_SDMA_DEMO_A5 ? 1 : -1) || + debug[base + 9U] != static_cast(iterations)) { + std::cerr << "ERROR: SDMA debug validation failed for block " << block << std::endl; + ok = false; + } } return ok; } @@ -150,10 +170,22 @@ bool VerifyDebug(const std::vector& debug, uint32_t bytes) int main(int argc, char** argv) { uint32_t bytes = kDefaultBytes; - if (argc > 2 || (argc == 2 && !ParseBytes(argv[1], &bytes))) { - std::cerr << "ERROR: bytes must be a non-zero 64-byte aligned unsigned integer" << std::endl; + uint32_t firstChannel = 0U; + uint32_t blocks = 1U; + uint32_t iterations = 1U; + if (argc > 5 || (argc >= 2 && !ParseBytes(argv[1], &bytes)) || + (argc >= 3 && !ParseUint32(argv[2], &firstChannel)) || + (argc >= 4 && !ParseUint32(argv[3], &blocks)) || + (argc == 5 && !ParseUint32(argv[4], &iterations)) || + blocks == 0U || iterations == 0U || + firstChannel >= TileXR::detail::TILEXR_SDMA_A5_CHANNEL_COUNT || + blocks > TileXR::detail::TILEXR_SDMA_A5_CHANNEL_COUNT - firstChannel) { + std::cerr << "ERROR: usage: tilexr_sdma_demo " + "[aligned-bytes [first-channel [blocks [iterations]]]]" << std::endl; return 1; } + const size_t totalBytes = static_cast(bytes) * blocks; + const size_t debugWords = static_cast(blocks) * kDebugWordsPerBlock; (void)setenv("TILEXR_ENABLE_SDMA", "1", 1); @@ -207,48 +239,63 @@ int main(int argc, char** argv) } if (!CheckAcl("aclrtMalloc src", - aclrtMalloc(reinterpret_cast(&src), bytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + aclrtMalloc(reinterpret_cast(&src), totalBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || !CheckAcl("aclrtMalloc dst", - aclrtMalloc(reinterpret_cast(&dst), bytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + aclrtMalloc(reinterpret_cast(&dst), totalBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || !CheckAcl("aclrtMalloc debug", - aclrtMalloc(reinterpret_cast(&debug), kDebugWords * sizeof(int32_t), ACL_MEM_MALLOC_HUGE_FIRST))) { + aclrtMalloc(reinterpret_cast(&debug), debugWords * sizeof(int32_t), ACL_MEM_MALLOC_HUGE_FIRST))) { Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); return 1; } - std::vector hostSrc(bytes); - std::vector hostDst(bytes, 0); - std::vector hostDebug(kDebugWords, 0); + std::vector hostSrc(totalBytes); + std::vector hostDst(totalBytes, 0); + std::vector hostDebug(debugWords, 0); FillPattern(hostSrc); - if (!CopyHostToDevice(src, bytes, hostSrc.data(), hostSrc.size(), "src") || - !CopyHostToDevice(dst, bytes, hostDst.data(), hostDst.size(), "dst") || - !CopyHostToDevice(debug, kDebugWords * sizeof(int32_t), hostDebug.data(), + if (!CopyHostToDevice(src, totalBytes, hostSrc.data(), hostSrc.size(), "src") || + !CopyHostToDevice(dst, totalBytes, hostDst.data(), hostDst.size(), "dst") || + !CopyHostToDevice(debug, debugWords * sizeof(int32_t), hostDebug.data(), hostDebug.size() * sizeof(int32_t), "debug")) { Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); return 1; } launch_tilexr_sdma_copy( - 1, stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), - reinterpret_cast(debug), bytes); - if (!CheckAcl("aclrtSynchronizeStream", aclrtSynchronizeStream(stream))) { + blocks, stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), + reinterpret_cast(debug), bytes, firstChannel, iterations); + if (!CheckAcl("aclrtSynchronizeStreamWithTimeout", + aclrtSynchronizeStreamWithTimeout(stream, kStreamTimeoutMs))) { Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); return 1; } - if (!CopyDeviceToHost(hostDst.data(), hostDst.size(), dst, bytes, "dst") || + if (!CopyDeviceToHost(hostDst.data(), hostDst.size(), dst, totalBytes, "dst") || !CopyDeviceToHost(hostDebug.data(), hostDebug.size() * sizeof(int32_t), debug, - kDebugWords * sizeof(int32_t), "debug")) { + debugWords * sizeof(int32_t), "debug")) { Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); return 1; } - std::cout << "debug words:"; - for (size_t i = 0; i < hostDebug.size(); ++i) { - std::cout << " d" << i << "=" << hostDebug[i]; + for (uint32_t block = 0U; block < blocks; ++block) { + const size_t base = static_cast(block) * kDebugWordsPerBlock; + std::cout << "block=" << block + << " channel=" << hostDebug[base + 6U] + << " generation=" << hostDebug[base + 7U] +#if TILEXR_SDMA_DEMO_A5 + << " busy_rejected=" << hostDebug[base + 8U] +#else + << " busy_check=skipped" +#endif + << " event=" << hostDebug[base + 4U] + << " wait=" << hostDebug[base + 5U] + << " workspace_valid=" << hostDebug[base + 10U] + << " sq_valid=" << hostDebug[base + 11U] + << " rtsq_valid=" << hostDebug[base + 12U] + << " queue_valid=" << hostDebug[base + 13U] + << " stream_valid=" << hostDebug[base + 14U] + << " first_wait_or_claim=" << hostDebug[base + 15U] << std::endl; } - std::cout << std::endl; const bool dataOk = (hostDst == hostSrc); if (!dataOk) { @@ -258,7 +305,7 @@ int main(int argc, char** argv) << " dst=" << static_cast(*mismatch.first) << " src=" << static_cast(*mismatch.second) << std::endl; } - const bool debugOk = VerifyDebug(hostDebug, bytes); + const bool debugOk = VerifyDebug(hostDebug, bytes, firstChannel, blocks, iterations); Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); if (!dataOk || !debugOk) { @@ -266,6 +313,8 @@ int main(int argc, char** argv) return 1; } - std::cout << "PASS TileXR SDMA copied " << bytes << " bytes correctly" << std::endl; + std::cout << "PASS TileXR SDMA copied " << bytes << " bytes on " << blocks + << " block(s), channels " << firstChannel << ".." + << (firstChannel + blocks - 1U) << ", iterations " << iterations << std::endl; return 0; } diff --git a/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp b/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp index eab37b11..1ec0ac62 100644 --- a/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp +++ b/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp @@ -18,32 +18,82 @@ extern "C" __global__ __aicore__ void tilexr_sdma_copy_kernel( GM_ADDR dstGM, GM_ADDR srcGM, GM_ADDR debugGM, - uint32_t bytes) + uint32_t bytes, + uint32_t firstChannel, + uint32_t iterations) { auto args = reinterpret_cast<__gm__ TileXR::CommArgs*>(commArgsGM); - auto dst = reinterpret_cast<__gm__ uint8_t*>(dstGM); - auto src = reinterpret_cast<__gm__ uint8_t*>(srcGM); - auto debug = reinterpret_cast<__gm__ int32_t*>(debugGM); + const uint32_t block = static_cast(AscendC::GetBlockIdx()); + const uint32_t channel = firstChannel + block; + auto dst = reinterpret_cast<__gm__ uint8_t*>(dstGM) + static_cast(block) * bytes; + auto src = reinterpret_cast<__gm__ uint8_t*>(srcGM) + static_cast(block) * bytes; + auto debug = reinterpret_cast<__gm__ int32_t*>(debugGM) + block * 16U; if ASCEND_IS_AIV { if (debug != nullptr) { debug[0] = TileXR::TILEXR_SDMA_DEMO_MAGIC; - debug[1] = static_cast(AscendC::GetBlockIdx()); + debug[1] = static_cast(block); debug[2] = static_cast(bytes); debug[3] = TileXR::SDMAEnabled(args) ? 1 : 0; debug[4] = 0; debug[5] = 0; + debug[6] = static_cast(channel); + debug[7] = 0; + debug[8] = 0; + debug[9] = static_cast(iterations); } - if (AscendC::GetBlockIdx() != 0) { - return; - } - uint64_t event = TileXR::SDMACopyNbi(args, dst, src, static_cast(bytes), 0); + uint64_t event = TileXR::SDMACopyNbi( + args, dst, src, static_cast(bytes), channel); + uint64_t busyEvent = 0U; +#if TILEXR_SDMA_A5_AICORE_COMPILE + busyEvent = TileXR::SDMACopyNbi( + args, dst, src, static_cast(bytes), channel); +#endif if (debug != nullptr) { debug[4] = event == 0 ? 0 : 1; +#if TILEXR_SDMA_A5_AICORE_COMPILE + debug[7] = static_cast(event & 0xFFFFFFFFULL); + debug[8] = busyEvent == 0 ? 1 : 0; +#else + debug[7] = event == 0 ? 0 : 1; + debug[8] = -1; +#endif + } + bool waitOk = TileXR::SDMAWait(args, event, channel); +#if TILEXR_SDMA_A5_AICORE_COMPILE + auto workspace = reinterpret_cast<__gm__ TileXR::detail::A5SdmaWorkspace*>( + args->sdmaWorkspacePtr); + auto channelState = &workspace->channels[channel]; + if (debug != nullptr) { + debug[10] = TileXR::detail::A5SdmaWorkspaceValid(workspace) ? 1 : 0; + debug[11] = channelState->sqBase != 0U ? 1 : 0; + debug[12] = channelState->rtsqAddress != 0U ? 1 : 0; + debug[13] = TileXR::detail::A5SdmaQueueStateValid( + channelState->tail, channelState->depth) ? 1 : 0; + debug[14] = channelState->streamId <= 0xFFFFU ? 1 : 0; + debug[15] = waitOk ? 1 : 0; + } +#endif + for (uint32_t iteration = 1U; iteration < iterations && waitOk; ++iteration) { + const uint64_t nextEvent = TileXR::SDMACopyNbi( + args, dst, src, static_cast(bytes), channel); + if (debug != nullptr) { +#if TILEXR_SDMA_A5_AICORE_COMPILE + debug[7] = static_cast(nextEvent & 0xFFFFFFFFULL); +#else + debug[7] = nextEvent == 0 ? debug[7] : static_cast(iteration + 1U); +#endif + } + waitOk = nextEvent != 0U && TileXR::SDMAWait(args, nextEvent, channel); + event = nextEvent; } - bool waitOk = TileXR::SDMAWait(args, event, 0); if (debug != nullptr) { debug[5] = waitOk ? 1 : 0; +#if TILEXR_SDMA_A5_AICORE_COMPILE + pipe_barrier(PIPE_ALL); + TileXR::detail::A5SdmaCleanCacheLine( + reinterpret_cast<__gm__ uint8_t*>(debug)); +#endif } } } @@ -55,7 +105,10 @@ extern "C" void launch_tilexr_sdma_copy( GM_ADDR dst, GM_ADDR src, GM_ADDR debug, - uint32_t bytes) + uint32_t bytes, + uint32_t firstChannel, + uint32_t iterations) { - tilexr_sdma_copy_kernel<<>>(commArgs, dst, src, debug, bytes); + tilexr_sdma_copy_kernel<<>>( + commArgs, dst, src, debug, bytes, firstChannel, iterations); } diff --git a/tests/sdma/run_tests.sh b/tests/sdma/run_tests.sh index 793cb9bb..3591e4ba 100755 --- a/tests/sdma/run_tests.sh +++ b/tests/sdma/run_tests.sh @@ -63,6 +63,7 @@ fi "${INSTALL_DIR}/bin/test_tilexr_sdma_metadata" "${INSTALL_DIR}/bin/test_tilexr_sdma_transport_disabled" +"${INSTALL_DIR}/bin/test_tilexr_sdma_a5_validation" "${INSTALL_DIR}/bin/test_tilexr_sdma_comm_wiring" "${INSTALL_DIR}/bin/test_tilexr_sdma_source_guard" "${INSTALL_DIR}/bin/test_tilexr_sdma_header_compile" diff --git a/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp b/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp new file mode 100644 index 00000000..572af880 --- /dev/null +++ b/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp @@ -0,0 +1,106 @@ +#include +#include + +#include "sdma/tilexr_sdma_a5_backend.h" + +namespace { + +int g_failures = 0; + +#define CHECK_TRUE(expr) \ + do { \ + if (!(expr)) { \ + std::cerr << "CHECK_TRUE failed at line " << __LINE__ << ": " #expr << std::endl; \ + ++g_failures; \ + } \ + } while (0) + +TileXR::detail::A5BuiltinChannelInfo ValidChannel() +{ + TileXR::detail::A5BuiltinChannelInfo channel {}; + channel.sqHead = 1U; + channel.sqTail = 2U; + channel.sqBase = 0x1000U; + channel.sqRegisterBase = 0x2000U; + channel.sqDepth = 8U; + channel.sqId = 3U; + channel.cqId = 4U; + channel.logicalCqId = 5U; + channel.streamId = 6U; + channel.deviceId = 7U; + return channel; +} + +TileXR::detail::A5HostChannelIdentity ValidIdentity() +{ + return {6U, 3U, 4U, 5U, 7U}; +} + +void TestCompleteAndExpectedPartialClassification() +{ + using namespace TileXR::detail; + A5BuiltinChannelInfo channel = ValidChannel(); + const A5HostChannelIdentity identity = ValidIdentity(); + CHECK_TRUE(ClassifyA5QueryResult(0, 1U, 48U, channel, identity) == + A5QueryResultKind::COMPLETE); + channel.sqRegisterBase = 0U; + CHECK_TRUE(ClassifyA5QueryResult(TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS, + 0U, 0U, channel, identity) == + A5QueryResultKind::EXPECTED_PARTIAL); +} + +void TestPartialClassificationFailsClosed() +{ + using namespace TileXR::detail; + A5BuiltinChannelInfo channel = ValidChannel(); + channel.sqRegisterBase = 0U; + const A5HostChannelIdentity identity = ValidIdentity(); + CHECK_TRUE(ClassifyA5QueryResult(507019, 0U, 0U, channel, identity) == + A5QueryResultKind::INVALID); + CHECK_TRUE(ClassifyA5QueryResult(TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS, + 1U, 0U, channel, identity) == + A5QueryResultKind::INVALID); + channel.sqTail = channel.sqDepth; + CHECK_TRUE(ClassifyA5QueryResult(TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS, + 0U, 0U, channel, identity) == + A5QueryResultKind::INVALID); + channel = ValidChannel(); + channel.sqRegisterBase = 0U; + channel.streamId += 1U; + CHECK_TRUE(ClassifyA5QueryResult(TILEXR_SDMA_A5_EXPECTED_QUERY_STATUS, + 0U, 0U, channel, identity) == + A5QueryResultKind::INVALID); + channel = ValidChannel(); + channel.sqHead = 0U; + channel.sqTail = channel.sqDepth - 1U; + CHECK_TRUE(ClassifyA5QueryResult(0, 1U, TILEXR_SDMA_A5_CHANNEL_COUNT, + channel, identity) == A5QueryResultKind::INVALID); +} + +void TestCompleteClassificationRequiresFinishedHeader() +{ + using namespace TileXR::detail; + const A5BuiltinChannelInfo channel = ValidChannel(); + const A5HostChannelIdentity identity = ValidIdentity(); + CHECK_TRUE(ClassifyA5QueryResult(0, 0U, TILEXR_SDMA_A5_CHANNEL_COUNT, + channel, identity) == A5QueryResultKind::INVALID); + CHECK_TRUE(ClassifyA5QueryResult(0, 1U, TILEXR_SDMA_A5_CHANNEL_COUNT - 1U, + channel, identity) == A5QueryResultKind::INVALID); + CHECK_TRUE(ClassifyA5QueryResult(0, 1U, 0U, channel, identity) == + A5QueryResultKind::INVALID); +} + +} // namespace + +int main() +{ + TestCompleteAndExpectedPartialClassification(); + TestPartialClassificationFailsClosed(); + TestCompleteClassificationRequiresFinishedHeader(); + if (g_failures != 0) { + std::cerr << g_failures << " A5 SDMA validation checks failed" << std::endl; + return 1; + } + std::cout << "TileXR A5 SDMA validation checks passed" << std::endl; + return 0; +} diff --git a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp index d6bd3016..be247c78 100644 --- a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp @@ -3,6 +3,7 @@ #include #include "comm_args.h" +#include "tilexr_sdma_a5_types.h" #include "tilexr_sdma_types.h" namespace { @@ -51,6 +52,65 @@ void TestSdmaConstants() CHECK_EQ(TileXR::TILEXR_SDMA_SCRATCH_BYTES, 256U); } +void TestA5WorkspaceAbi() +{ + using namespace TileXR::detail; + CHECK_EQ(TILEXR_SDMA_A5_CHANNEL_COUNT, 48U); + CHECK_TRUE(TILEXR_SDMA_A5_WAIT_MAX_POLLS > 0U); + CHECK_EQ(sizeof(A5SdmaWorkspaceHeader), 64U); + CHECK_EQ(sizeof(A5SdmaChannel), 192U); + CHECK_EQ(sizeof(A5SdmaCompletionLine), 64U); + CHECK_EQ(sizeof(A5SdmaSqe), 64U); + CHECK_EQ(offsetof(A5SdmaChannel, generation), 72U); + CHECK_EQ(offsetof(A5SdmaChannel, outstanding), 128U); + CHECK_EQ(offsetof(A5SdmaSqe, srcAddressLow), 32U); + CHECK_EQ(offsetof(A5SdmaSqe, length), 48U); + CHECK_EQ(offsetof(A5SdmaWorkspace, channels), 64U); + CHECK_EQ(sizeof(A5SdmaWorkspace) % 64U, 0U); +} + +void TestA5QueueAndTransferHelpers() +{ + using namespace TileXR::detail; + CHECK_TRUE(!A5SdmaQueueStateValid(0U, 2U)); + CHECK_TRUE(A5SdmaQueueStateValid(2U, 3U)); + CHECK_TRUE(!A5SdmaQueueStateValid(3U, 3U)); + CHECK_EQ(A5SdmaAdvanceTail(0U, 3U), 2U); + CHECK_EQ(A5SdmaAdvanceTail(2U, 3U), 1U); + CHECK_EQ(A5SdmaQueueDistance(2U, 1U, 3U), 2U); + CHECK_TRUE(A5SdmaQueueHasCapacity(0U, 0U, 3U)); + CHECK_TRUE(A5SdmaQueueHasCapacity(1U, 2U, 8U)); + CHECK_TRUE(!A5SdmaQueueHasCapacity(0U, 2U, 3U)); + CHECK_TRUE(!A5SdmaQueueHasCapacity(3U, 0U, 3U)); + CHECK_EQ(A5SdmaAdvanceTaskId(0xFFFFU), 1U); + CHECK_TRUE(!A5SdmaTransferLengthValid(0U)); + CHECK_TRUE(A5SdmaTransferLengthValid(1U)); + CHECK_TRUE(A5SdmaTransferLengthValid(TILEXR_SDMA_A5_MAX_TRANSFER_BYTES)); + CHECK_TRUE(!A5SdmaTransferLengthValid(TILEXR_SDMA_A5_MAX_TRANSFER_BYTES + 1U)); + CHECK_EQ(A5SdmaNextGeneration(0U), 1U); + CHECK_EQ(A5SdmaNextGeneration(1U), 2U); + CHECK_EQ(A5SdmaNextGeneration(0xFFFFFFFFU), 1U); +} + +void TestA5EventHelpers() +{ + using namespace TileXR::detail; + uint32_t channel = 0U; + uint32_t generation = 0U; + const uint64_t event = A5SdmaEncodeEvent(47U, 0x12345678U); + CHECK_TRUE(event != 0U); + CHECK_TRUE(A5SdmaDecodeEvent(event, channel, generation)); + CHECK_EQ(channel, 47U); + CHECK_EQ(generation, 0x12345678U); + CHECK_EQ(A5SdmaEncodeEvent(48U, 1U), 0U); + CHECK_EQ(A5SdmaEncodeEvent(0U, 0U), 0U); + CHECK_TRUE(!A5SdmaDecodeEvent(0U, channel, generation)); + CHECK_TRUE(!A5SdmaDecodeEvent(event | TILEXR_SDMA_A5_EVENT_RESERVED_MASK, + channel, generation)); + CHECK_TRUE(!A5SdmaDecodeEvent(event & ~TILEXR_SDMA_A5_EVENT_GENERATION_MASK, + channel, generation)); +} + } // namespace int main() @@ -58,6 +118,9 @@ int main() TestSdmaFlagDoesNotOverlapExistingFlags(); TestCommArgsHasSdmaWorkspace(); TestSdmaConstants(); + TestA5WorkspaceAbi(); + TestA5QueueAndTransferHelpers(); + TestA5EventHelpers(); if (g_failures != 0) { std::cerr << g_failures << " SDMA metadata checks failed" << std::endl; return 1; diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index cfe96e35..cce95ba2 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -108,12 +108,73 @@ void TestOnlyCompatIncludesSdmaIntrinsics() } } +void TestBuildSelectsExplicitSoc() +{ + const std::string buildPath = "tests/sdma/build.sh"; + const auto buildText = ReadFile(buildPath); + CheckNeedle(buildPath, buildText, + "SDMA_SOC_TYPE=\"${2:-${TILEXR_SDMA_DEMO_SOC_TYPE:-Ascend910B}}\""); + CheckNeedle(buildPath, buildText, + "-DTILEXR_SDMA_DEMO_SOC_TYPE=\"${SDMA_SOC_TYPE}\""); + + const std::string cmakePath = "tests/sdma/CMakeLists.txt"; + const auto cmakeText = ReadFile(cmakePath); + CheckNeedle(cmakePath, cmakeText, + "TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL \"Ascend950\""); + CheckNeedle(cmakePath, cmakeText, + "TILEXR_SDMA_NPU_ARCH \"dav-3510\""); + CheckNeedle(cmakePath, cmakeText, + "TILEXR_SDMA_AICORE_ARCH \"--cce-aicore-arch=dav-c310-vec\""); + CheckNeedle(cmakePath, cmakeText, + "Unsupported TILEXR_SDMA_DEMO_SOC_TYPE="); +} + +void TestAscend950UsesOwnedDirectBackend() +{ + const std::string transportPath = "src/comm/sdma/tilexr_sdma_transport.cpp"; + const auto transport = ReadFile(transportPath); + CheckNeedle(transportPath, transport, "TileXRA5SDMABackend"); + CheckNeedle(transportPath, transport, "ClassifySDMABackend(socName)"); + + const std::string backendPath = "src/comm/sdma/tilexr_sdma_a5_backend.cpp"; + const auto backend = ReadFile(backendPath); + CheckNeedle(backendPath, backend, "aclnnShmemSdmaStarsQuery"); + CheckNeedle(backendPath, backend, "RES_ADDR_TYPE_STARS_RTSQ"); + CheckNeedle(backendPath, backend, "TILEXR_SDMA_A5_CHANNEL_COUNT"); + CheckNeedle(backendPath, backend, "kExpectedAicpuQueryFailure"); + + const std::string devicePath = "src/include/tilexr_sdma_a5.h"; + const auto device = ReadFile(devicePath); + CheckNeedle(devicePath, device, "TILEXR_SDMA_A5_WAIT_MAX_POLLS"); + + const std::string demoPath = "tests/sdma/demo/tilexr_sdma_demo.cpp"; + const auto demo = ReadFile(demoPath); + CheckNeedle(demoPath, demo, "aclrtSynchronizeStreamWithTimeout"); +} + +void TestSdmaFailureIsScopedToOneCommunicator() +{ + const std::string commPath = "src/comm/tilexr_comm.cpp"; + const auto comm = ReadFile(commPath); + CheckNoNeedle(commPath, comm, "g_sdmaUnavailable"); + CheckNeedle(commPath, comm, "sdmaTransport_->Init(options)"); + CheckNeedle(commPath, comm, "bool TileXRComm::PrepareDestroy()"); + + const std::string wrapPath = "src/comm/comm_wrap.cpp"; + const auto wrap = ReadFile(wrapPath); + CheckNeedle(wrapPath, wrap, "if (!c->PrepareDestroy())"); + CheckNeedle(wrapPath, wrap, "return TILEXR_ERROR_INTERNAL;"); +} + } // namespace int main() { TestCommSourcesDoNotUseShmem(); TestOnlyCompatIncludesSdmaIntrinsics(); + TestBuildSelectsExplicitSoc(); + TestAscend950UsesOwnedDirectBackend(); + TestSdmaFailureIsScopedToOneCommunicator(); if (g_failures != 0) { std::cerr << g_failures << " SDMA source guard checks failed" << std::endl; return 1; diff --git a/tests/sdma/unit/test_tilexr_sdma_transport_disabled.cpp b/tests/sdma/unit/test_tilexr_sdma_transport_disabled.cpp index 7e85967d..2504b22b 100644 --- a/tests/sdma/unit/test_tilexr_sdma_transport_disabled.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_transport_disabled.cpp @@ -2,6 +2,7 @@ #include #include +#include "acl/acl_rt.h" #include "sdma/tilexr_sdma_transport.h" #include "tilexr_sdma_types.h" #include "tilexr_types.h" @@ -84,7 +85,27 @@ void TestEnvZeroSkipsInitialization() CHECK_EQ(transport.GetLastStatus(), TileXR::SDMAInitStatus::DISABLED_BY_ENV); } -void TestEnvOneReportsPTOUnavailable() +TileXR::SDMAInitStatus ExpectedUnavailableStatus() +{ + return TileXR::detail::ClassifySDMABackend(aclrtGetSocName()) == + TileXR::detail::SDMABackendKind::A5_DIRECT + ? TileXR::SDMAInitStatus::INIT_FAILED + : TileXR::SDMAInitStatus::PTO_UNAVAILABLE; +} + +void TestBackendClassification() +{ + using TileXR::detail::ClassifySDMABackend; + using TileXR::detail::SDMABackendKind; + CHECK_EQ(ClassifySDMABackend("Ascend950PR_9589"), SDMABackendKind::A5_DIRECT); + CHECK_EQ(ClassifySDMABackend("Ascend910B3"), SDMABackendKind::PTO); + CHECK_EQ(ClassifySDMABackend("Ascend910A"), SDMABackendKind::PTO); + CHECK_EQ(ClassifySDMABackend("Ascend910_9391"), SDMABackendKind::PTO); + CHECK_EQ(ClassifySDMABackend("Ascend310P"), SDMABackendKind::UNSUPPORTED); + CHECK_EQ(ClassifySDMABackend(nullptr), SDMABackendKind::UNSUPPORTED); +} + +void TestEnvOneReportsBackendUnavailable() { EnvGuard env; setenv("TILEXR_ENABLE_SDMA", "1", 1); @@ -94,7 +115,7 @@ void TestEnvOneReportsPTOUnavailable() CHECK_EQ(transport.Init(options), TileXR::TILEXR_SUCCESS); CHECK_TRUE(!transport.IsAvailable()); CHECK_TRUE(transport.GetWorkspaceDev() == nullptr); - CHECK_EQ(transport.GetLastStatus(), TileXR::SDMAInitStatus::PTO_UNAVAILABLE); + CHECK_EQ(transport.GetLastStatus(), ExpectedUnavailableStatus()); } void TestSameInstanceTransitionsResetState() @@ -114,12 +135,12 @@ void TestSameInstanceTransitionsResetState() CHECK_EQ(transport.Init(options), TileXR::TILEXR_SUCCESS); CHECK_TRUE(!transport.IsAvailable()); CHECK_TRUE(transport.GetWorkspaceDev() == nullptr); - CHECK_EQ(transport.GetLastStatus(), TileXR::SDMAInitStatus::PTO_UNAVAILABLE); + CHECK_EQ(transport.GetLastStatus(), ExpectedUnavailableStatus()); transport.Shutdown(); CHECK_TRUE(!transport.IsAvailable()); CHECK_TRUE(transport.GetWorkspaceDev() == nullptr); - CHECK_EQ(transport.GetLastStatus(), TileXR::SDMAInitStatus::PTO_UNAVAILABLE); + CHECK_EQ(transport.GetLastStatus(), ExpectedUnavailableStatus()); unsetenv("TILEXR_ENABLE_SDMA"); CHECK_EQ(transport.Init(options), TileXR::TILEXR_SUCCESS); @@ -134,7 +155,8 @@ int main() { TestEnvDisabledSkipsInitialization(); TestEnvZeroSkipsInitialization(); - TestEnvOneReportsPTOUnavailable(); + TestBackendClassification(); + TestEnvOneReportsBackendUnavailable(); TestSameInstanceTransitionsResetState(); if (g_failures != 0) { std::cerr << g_failures << " SDMA transport disabled checks failed" << std::endl; From 6996191f9d74611f7e92b319a68a503ba2a0d85e Mon Sep 17 00:00:00 2001 From: Kur0x Date: Wed, 29 Jul 2026 15:08:18 +0800 Subject: [PATCH 2/7] fix: make A5 SDMA cleanup retryable --- src/comm/CMakeLists.txt | 1 + src/comm/sdma/tilexr_sdma_a5_backend.cpp | 366 ++++++++++-------- src/comm/sdma/tilexr_sdma_a5_cleanup.h | 173 +++++++++ .../unit/test_tilexr_sdma_a5_validation.cpp | 159 ++++++++ .../unit/test_tilexr_sdma_source_guard.cpp | 5 + 5 files changed, 549 insertions(+), 155 deletions(-) create mode 100644 src/comm/sdma/tilexr_sdma_a5_cleanup.h diff --git a/src/comm/CMakeLists.txt b/src/comm/CMakeLists.txt index 202b2c24..91e8e0ab 100644 --- a/src/comm/CMakeLists.txt +++ b/src/comm/CMakeLists.txt @@ -118,6 +118,7 @@ set(TILEXR_SOURCE_FILE tilexr_comm.cpp udma/tilexr_udma_transport.cpp sdma/tilexr_sdma_transport.h sdma/tilexr_sdma_transport.cpp + sdma/tilexr_sdma_a5_cleanup.h sdma/tilexr_sdma_a5_backend.h sdma/tilexr_sdma_a5_backend.cpp ) diff --git a/src/comm/sdma/tilexr_sdma_a5_backend.cpp b/src/comm/sdma/tilexr_sdma_a5_backend.cpp index b9bb7357..9873d8f1 100644 --- a/src/comm/sdma/tilexr_sdma_a5_backend.cpp +++ b/src/comm/sdma/tilexr_sdma_a5_backend.cpp @@ -10,11 +10,13 @@ #include #include #include +#include #include #include "acl/acl.h" #include "aclnn/aclnn_base.h" #include "driver/ascend_hal.h" +#include "sdma/tilexr_sdma_a5_cleanup.h" #include "tilexr_log.h" #ifndef ACL_STREAM_DEVICE_USE_ONLY @@ -34,6 +36,9 @@ using RtGetDevicePhyIdByIndexFn = int32_t (*)(uint32_t, uint32_t*); using RtStreamGetSqidFn = int32_t (*)(const void*, uint32_t*); using RtStreamGetCqidFn = int32_t (*)(const void*, uint32_t*, uint32_t*); using RtGetDeviceInfoFn = int32_t (*)(uint32_t, int32_t, int32_t, int64_t*); +using HalResAddrMapFn = drvError_t (*)(unsigned int, res_addr_info*, + unsigned long*, unsigned int*); +using HalResAddrUnmapFn = drvError_t (*)(unsigned int, res_addr_info*); using AclCreateTensorFn = aclTensor* (*)(const int64_t*, uint64_t, aclDataType, const int64_t*, int64_t, aclFormat, const int64_t*, uint64_t, void*); @@ -89,6 +94,12 @@ class A5RuntimeApi { Close(); return false; } + + if (!LoadSymbol(RTLD_DEFAULT, "halResAddrMap", mapResource) || + !LoadSymbol(RTLD_DEFAULT, "halResAddrUnmap", unmapResource)) { + Close(); + return false; + } return true; } @@ -102,6 +113,8 @@ class A5RuntimeApi { destroyTensor = nullptr; prepareQuery = nullptr; executeQuery = nullptr; + mapResource = nullptr; + unmapResource = nullptr; if (opapiHandle_ != nullptr) { (void)dlclose(opapiHandle_); opapiHandle_ = nullptr; @@ -120,207 +133,211 @@ class A5RuntimeApi { AclDestroyTensorFn destroyTensor = nullptr; AclnnQueryWorkspaceFn prepareQuery = nullptr; AclnnQueryFn executeQuery = nullptr; + HalResAddrMapFn mapResource = nullptr; + HalResAddrUnmapFn unmapResource = nullptr; private: void* runtimeHandle_ = nullptr; void* opapiHandle_ = nullptr; }; -class DeviceBuffer { -public: - ~DeviceBuffer() - { - Reset(); - } - - bool Allocate(size_t bytes, bool zero = false) - { - Reset(); - if (aclrtMalloc(&address_, bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { - return false; - } - bytes_ = bytes; - if (zero && aclrtMemset(address_, bytes_, 0, bytes_) != ACL_SUCCESS) { - Reset(); - return false; - } - return true; - } - - void Reset() - { - if (address_ != nullptr) { - (void)aclrtFree(address_); - address_ = nullptr; - bytes_ = 0U; - } - } - - void* Get() const - { - return address_; - } - -private: - void* address_ = nullptr; - size_t bytes_ = 0U; +struct QuerySnapshot { + uint32_t flag = 0U; + uint32_t totalQueueCount = 0U; + aclError syncStatus = ACL_SUCCESS; + std::vector channels; }; -class TensorHandle { -public: - explicit TensorHandle(AclDestroyTensorFn destroy) : destroy_(destroy) {} - - ~TensorHandle() - { - Reset(); - } +int CleanupSetCurrentContext(void*, void* context) +{ + return static_cast( + aclrtSetCurrentContext(static_cast(context))); +} - void Reset() - { - if (tensor_ != nullptr) { - (void)destroy_(tensor_); - tensor_ = nullptr; - } - } +int CleanupDestroyStream(void*, void* stream) +{ + return static_cast(aclrtDestroyStream(static_cast(stream))); +} - aclTensor*& Ref() - { - return tensor_; - } +int CleanupDestroyContext(void*, void* context) +{ + return static_cast(aclrtDestroyContext(static_cast(context))); +} - aclTensor* Get() const - { - return tensor_; - } +int CleanupFreeDevice(void*, void* address) +{ + return static_cast(aclrtFree(address)); +} -private: - AclDestroyTensorFn destroy_; - aclTensor* tensor_ = nullptr; -}; +int CleanupDestroyTensor(void* opaque, const void* tensor) +{ + A5RuntimeApi* api = static_cast(opaque); + return api == nullptr || api->destroyTensor == nullptr + ? -1 + : api->destroyTensor(static_cast(tensor)); +} -struct QuerySnapshot { - uint32_t flag = 0U; - uint32_t totalQueueCount = 0U; - aclError syncStatus = ACL_SUCCESS; - std::vector channels; -}; +detail::A5QueryCleanupOps MakeQueryCleanupOps(A5RuntimeApi& api) +{ + detail::A5QueryCleanupOps ops; + ops.opaque = &api; + ops.setCurrentContext = CleanupSetCurrentContext; + ops.destroyStream = CleanupDestroyStream; + ops.destroyContext = CleanupDestroyContext; + ops.freeDevice = CleanupFreeDevice; + ops.destroyTensor = CleanupDestroyTensor; + return ops; +} -bool CreateUint64Tensor(A5RuntimeApi& api, void* address, int64_t elements, - TensorHandle& tensor) +bool AllocateTrackedBuffer(std::vector& buffers, size_t bytes, + bool zero, void*& address) { - const int64_t shape[] = {elements}; - const int64_t strides[] = {1}; - tensor.Ref() = api.createTensor(shape, 1U, ACL_UINT64, strides, 0, - ACL_FORMAT_ND, shape, 1U, address); - return tensor.Get() != nullptr; + address = nullptr; + if (aclrtMalloc(&address, bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { + return false; + } + buffers.push_back(address); + return !zero || aclrtMemset(address, bytes, 0, bytes) == ACL_SUCCESS; } -bool RestoreQueryContext(aclrtContext previous, aclrtContext& isolated) +bool CreateTrackedUint64Tensor(A5RuntimeApi& api, + detail::A5PendingQueryCleanup& cleanup, + void* address, int64_t elements, + aclTensor*& tensor) { - const aclError destroyStatus = isolated == nullptr - ? ACL_SUCCESS - : aclrtDestroyContext(isolated); - isolated = nullptr; - const aclError restoreStatus = aclrtSetCurrentContext(previous); - if (destroyStatus != ACL_SUCCESS || restoreStatus != ACL_SUCCESS) { - TILEXR_LOG(WARN) << "TileXR A5 SDMA query context restore failed, destroy " - << destroyStatus << ", restore " << restoreStatus; + const int64_t shape[] = {elements}; + const int64_t strides[] = {1}; + tensor = api.createTensor(shape, 1U, ACL_UINT64, strides, 0, + ACL_FORMAT_ND, shape, 1U, address); + if (tensor == nullptr) { return false; } + cleanup.tensors.push_back(tensor); return true; } -bool CheckRuntimeHealth(DeviceBuffer& scratch) +bool CheckRuntimeHealth(detail::A5PendingQueryCleanup& cleanup, void* scratch) { aclrtStream stream = nullptr; - if (aclrtCreateStream(&stream) != ACL_SUCCESS) { + const aclError createStatus = aclrtCreateStream(&stream); + cleanup.healthStream = stream; + if (createStatus != ACL_SUCCESS || stream == nullptr) { return false; } const aclError memsetStatus = aclrtMemsetAsync( - scratch.Get(), sizeof(uint64_t), 0xA5, sizeof(uint64_t), stream); + scratch, sizeof(uint64_t), 0xA5, sizeof(uint64_t), stream); const aclError syncStatus = memsetStatus == ACL_SUCCESS ? aclrtSynchronizeStream(stream) : memsetStatus; uint64_t value = 0U; const aclError copyStatus = syncStatus == ACL_SUCCESS - ? aclrtMemcpy(&value, sizeof(value), scratch.Get(), sizeof(value), ACL_MEMCPY_DEVICE_TO_HOST) + ? aclrtMemcpy(&value, sizeof(value), scratch, sizeof(value), + ACL_MEMCPY_DEVICE_TO_HOST) : syncStatus; - const aclError destroyStatus = aclrtDestroyStream(stream); return memsetStatus == ACL_SUCCESS && syncStatus == ACL_SUCCESS && - copyStatus == ACL_SUCCESS && destroyStatus == ACL_SUCCESS && - value == 0xA5A5A5A5A5A5A5A5ULL; + copyStatus == ACL_SUCCESS && value == 0xA5A5A5A5A5A5A5A5ULL; +} + +bool FinishQuery(A5RuntimeApi& api, + detail::A5PendingQueryCleanup& cleanup, + std::vector& pending, + bool result) +{ + const bool released = detail::CleanupA5QueryResources( + cleanup, MakeQueryCleanupOps(api), cleanup.ownerContext); + if (!cleanup.Empty()) { + pending.push_back(std::move(cleanup)); + } + return result && released; } bool RunBuiltinQuery(A5RuntimeApi& api, int32_t logicalDevice, const std::vector& streams, + std::vector& pending, QuerySnapshot& snapshot) { if (streams.empty() || streams.size() > detail::TILEXR_SDMA_A5_CHANNEL_COUNT) { return false; } + detail::A5PendingQueryCleanup cleanup; + aclrtContext ownerContext = nullptr; + if (aclrtGetCurrentContext(&ownerContext) != ACL_SUCCESS || + ownerContext == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query requires an active owner context"; + return false; + } + cleanup.ownerContext = ownerContext; + const size_t streamsBytes = streams.size() * sizeof(detail::A5BuiltinStreamInfo); - DeviceBuffer streamsDev; - DeviceBuffer resourceDev; - DeviceBuffer builtinWorkspaceDev; - DeviceBuffer inputDev; - DeviceBuffer outputDev; - DeviceBuffer opWorkspaceDev; - if (!streamsDev.Allocate(streamsBytes) || - !resourceDev.Allocate(sizeof(detail::A5BuiltinOpResource), true) || - !builtinWorkspaceDev.Allocate(kBuiltinWorkspaceBytes, true) || - !inputDev.Allocate(2U * sizeof(uint64_t)) || - !outputDev.Allocate(sizeof(uint64_t), true) || - aclrtMemcpy(streamsDev.Get(), streamsBytes, streams.data(), streamsBytes, + void* streamsDev = nullptr; + void* resourceDev = nullptr; + void* builtinWorkspaceDev = nullptr; + void* inputDev = nullptr; + void* outputDev = nullptr; + void* opWorkspaceDev = nullptr; + if (!AllocateTrackedBuffer(cleanup.ownerBuffers, streamsBytes, false, streamsDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + sizeof(detail::A5BuiltinOpResource), true, resourceDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + kBuiltinWorkspaceBytes, true, builtinWorkspaceDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + 2U * sizeof(uint64_t), false, inputDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + sizeof(uint64_t), true, outputDev) || + aclrtMemcpy(streamsDev, streamsBytes, streams.data(), streamsBytes, ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { TILEXR_LOG(WARN) << "TileXR A5 SDMA query buffer setup failed"; - return false; + return FinishQuery(api, cleanup, pending, false); } detail::A5BuiltinOpResource resource {}; resource.size = streams.size(); - resource.streamsAddress = reinterpret_cast(streamsDev.Get()); - resource.workspaceAddress = reinterpret_cast(builtinWorkspaceDev.Get()); + resource.streamsAddress = reinterpret_cast(streamsDev); + resource.workspaceAddress = reinterpret_cast(builtinWorkspaceDev); const uint64_t inputs[] = { - reinterpret_cast(resourceDev.Get()), - reinterpret_cast(builtinWorkspaceDev.Get()), + reinterpret_cast(resourceDev), + reinterpret_cast(builtinWorkspaceDev), }; - if (aclrtMemcpy(resourceDev.Get(), sizeof(resource), &resource, sizeof(resource), + if (aclrtMemcpy(resourceDev, sizeof(resource), &resource, sizeof(resource), ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS || - aclrtMemcpy(inputDev.Get(), sizeof(inputs), inputs, sizeof(inputs), + aclrtMemcpy(inputDev, sizeof(inputs), inputs, sizeof(inputs), ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { TILEXR_LOG(WARN) << "TileXR A5 SDMA query resource upload failed"; - return false; + return FinishQuery(api, cleanup, pending, false); } - aclrtContext previous = nullptr; aclrtContext isolated = nullptr; - if (aclrtGetCurrentContext(&previous) != ACL_SUCCESS || - aclrtCreateContext(&isolated, logicalDevice) != ACL_SUCCESS) { + const aclError createContextStatus = aclrtCreateContext(&isolated, logicalDevice); + cleanup.isolatedContext = isolated; + if (createContextStatus != ACL_SUCCESS || isolated == nullptr) { TILEXR_LOG(WARN) << "TileXR A5 SDMA isolated query context creation failed"; - if (isolated != nullptr) { - (void)RestoreQueryContext(previous, isolated); - } - return false; + return FinishQuery(api, cleanup, pending, false); } - TensorHandle inputTensor(api.destroyTensor); - TensorHandle outputTensor(api.destroyTensor); + aclTensor* inputTensor = nullptr; + aclTensor* outputTensor = nullptr; aclrtStream queryStream = nullptr; bool queryLaunched = false; aclError syncStatus = ACL_SUCCESS; uint64_t opWorkspaceBytes = 0U; aclOpExecutor* executor = nullptr; - bool ok = CreateUint64Tensor(api, inputDev.Get(), 2, inputTensor) && - CreateUint64Tensor(api, outputDev.Get(), 1, outputTensor) && - api.prepareQuery(inputTensor.Get(), outputTensor.Get(), &opWorkspaceBytes, &executor) == ACL_SUCCESS; + bool ok = CreateTrackedUint64Tensor( + api, cleanup, inputDev, 2, inputTensor) && + CreateTrackedUint64Tensor(api, cleanup, outputDev, 1, outputTensor) && + api.prepareQuery(inputTensor, outputTensor, + &opWorkspaceBytes, &executor) == ACL_SUCCESS; if (ok && opWorkspaceBytes != 0U) { - ok = opWorkspaceDev.Allocate(static_cast(opWorkspaceBytes)); + ok = AllocateTrackedBuffer(cleanup.isolatedBuffers, + static_cast(opWorkspaceBytes), false, + opWorkspaceDev); } if (ok) { - ok = aclrtCreateStreamWithConfig( - &queryStream, 0, ACL_STREAM_FAST_LAUNCH | ACL_STREAM_FAST_SYNC) == ACL_SUCCESS; + const aclError createStreamStatus = aclrtCreateStreamWithConfig( + &queryStream, 0, ACL_STREAM_FAST_LAUNCH | ACL_STREAM_FAST_SYNC); + cleanup.queryStream = queryStream; + ok = createStreamStatus == ACL_SUCCESS && queryStream != nullptr; } if (ok) { aclrtStreamAttrValue failureMode {}; @@ -330,7 +347,7 @@ bool RunBuiltinQuery(A5RuntimeApi& api, } if (ok) { const aclnnStatus launchStatus = api.executeQuery( - opWorkspaceDev.Get(), opWorkspaceBytes, executor, queryStream); + opWorkspaceDev, opWorkspaceBytes, executor, queryStream); queryLaunched = launchStatus == ACL_SUCCESS; ok = queryLaunched; if (queryLaunched) { @@ -338,26 +355,22 @@ bool RunBuiltinQuery(A5RuntimeApi& api, } } - aclError streamDestroyStatus = ACL_SUCCESS; - if (queryStream != nullptr) { - streamDestroyStatus = aclrtDestroyStream(queryStream); - queryStream = nullptr; + if (aclrtSetCurrentContext(ownerContext) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not restore query owner context"; + return FinishQuery(api, cleanup, pending, false); } - inputTensor.Reset(); - outputTensor.Reset(); - const bool contextRestored = RestoreQueryContext(previous, isolated); - if (!ok || !queryLaunched || streamDestroyStatus != ACL_SUCCESS || !contextRestored) { - TILEXR_LOG(WARN) << "TileXR A5 SDMA built-in query launch or cleanup failed"; - return false; + if (!ok || !queryLaunched) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA built-in query launch failed"; + return FinishQuery(api, cleanup, pending, false); } const size_t snapshotBytes = sizeof(detail::A5BuiltinWorkspaceHeader) + streams.size() * sizeof(detail::A5BuiltinChannelInfo); std::vector bytes(snapshotBytes, 0U); - if (aclrtMemcpy(bytes.data(), bytes.size(), builtinWorkspaceDev.Get(), bytes.size(), + if (aclrtMemcpy(bytes.data(), bytes.size(), builtinWorkspaceDev, bytes.size(), ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) { TILEXR_LOG(WARN) << "TileXR A5 SDMA query workspace download failed"; - return false; + return FinishQuery(api, cleanup, pending, false); } detail::A5BuiltinWorkspaceHeader header {}; std::memcpy(&header, bytes.data(), sizeof(header)); @@ -368,11 +381,12 @@ bool RunBuiltinQuery(A5RuntimeApi& api, std::memcpy(snapshot.channels.data(), bytes.data() + sizeof(header), snapshot.channels.size() * sizeof(snapshot.channels[0])); - if (syncStatus == kExpectedAicpuQueryFailure && !CheckRuntimeHealth(outputDev)) { + if (syncStatus == kExpectedAicpuQueryFailure && + !CheckRuntimeHealth(cleanup, outputDev)) { TILEXR_LOG(WARN) << "TileXR A5 SDMA context health check failed after expected AICPU error"; - return false; + return FinishQuery(api, cleanup, pending, false); } - return true; + return FinishQuery(api, cleanup, pending, true); } int32_t QueryHostSq(uint32_t physicalDevice, uint32_t sqId, @@ -407,15 +421,17 @@ struct TileXRA5SDMABackend::Impl { int32_t logicalDevice = -1; uint32_t physicalDevice = 0U; uint32_t physicalDieId = 0U; + A5RuntimeApi api; aclrtContext ownerContext = nullptr; aclrtContext restoreContext = nullptr; bool restorePending = false; void* workspaceDev = nullptr; + std::vector pendingQueries; std::array channels {}; bool HasOwnedResources() const { - if (workspaceDev != nullptr) { + if (workspaceDev != nullptr || !pendingQueries.empty()) { return true; } for (const OwnedChannel& channel : channels) { @@ -425,6 +441,18 @@ struct TileXRA5SDMABackend::Impl { } return false; } + + void EraseCompletedQueries() + { + auto query = pendingQueries.begin(); + while (query != pendingQueries.end()) { + if (query->Empty()) { + query = pendingQueries.erase(query); + } else { + ++query; + } + } + } }; TileXRA5SDMABackend::TileXRA5SDMABackend() = default; @@ -451,15 +479,16 @@ bool TileXRA5SDMABackend::Init(int32_t deviceId) return false; } - A5RuntimeApi api; - if (!api.Load() || - api.getPhysicalDevice(static_cast(deviceId), &state->physicalDevice) != 0) { + if (!state->api.Load() || + state->api.getPhysicalDevice( + static_cast(deviceId), &state->physicalDevice) != 0) { TILEXR_LOG(WARN) << "TileXR A5 SDMA runtime discovery failed"; return false; } int64_t physicalDie = -1; - if (api.getDeviceInfo(static_cast(deviceId), kDeviceInfoModuleType, - kPhysicalDieInfoType, &physicalDie) != 0 || + if (state->api.getDeviceInfo( + static_cast(deviceId), kDeviceInfoModuleType, + kPhysicalDieInfoType, &physicalDie) != 0 || physicalDie < 0 || static_cast(physicalDie) > static_cast(std::numeric_limits::max())) { TILEXR_LOG(WARN) << "TileXR A5 SDMA physical die discovery failed"; @@ -467,6 +496,7 @@ bool TileXRA5SDMABackend::Init(int32_t deviceId) } state->physicalDieId = static_cast(physicalDie); impl_ = std::move(state); + A5RuntimeApi& api = impl_->api; std::vector streamInfos; streamInfos.reserve(detail::TILEXR_SDMA_A5_CHANNEL_COUNT); @@ -498,7 +528,7 @@ bool TileXRA5SDMABackend::Init(int32_t deviceId) owned.mapInfo.res_id = sqId; unsigned long mappedAddress = 0UL; unsigned int mappedLength = 0U; - const drvError_t mapStatus = halResAddrMap( + const drvError_t mapStatus = api.mapResource( impl_->physicalDevice, &owned.mapInfo, &mappedAddress, &mappedLength); owned.mapped = mapStatus == DRV_ERROR_NONE; if (mapStatus != DRV_ERROR_NONE || mappedAddress == 0UL || @@ -523,7 +553,8 @@ bool TileXRA5SDMABackend::Init(int32_t deviceId) } QuerySnapshot batch; - if (!RunBuiltinQuery(api, deviceId, streamInfos, batch)) { + if (!RunBuiltinQuery( + api, deviceId, streamInfos, impl_->pendingQueries, batch)) { Shutdown(); return false; } @@ -551,7 +582,8 @@ bool TileXRA5SDMABackend::Init(int32_t deviceId) for (uint32_t index = 1U; index < detail::TILEXR_SDMA_A5_CHANNEL_COUNT; ++index) { std::vector oneStream(1U, streamInfos[index]); QuerySnapshot isolated; - if (!RunBuiltinQuery(api, deviceId, oneStream, isolated) || + if (!RunBuiltinQuery( + api, deviceId, oneStream, impl_->pendingQueries, isolated) || isolated.channels.size() != 1U || detail::ClassifyA5QueryResult( static_cast(isolated.syncStatus), isolated.flag, @@ -654,11 +686,34 @@ bool TileXRA5SDMABackend::Shutdown() } } + const detail::A5QueryCleanupOps queryCleanupOps = + MakeQueryCleanupOps(impl_->api); + for (size_t reverse = impl_->pendingQueries.size(); reverse > 0U; --reverse) { + detail::A5PendingQueryCleanup& query = impl_->pendingQueries[reverse - 1U]; + if (query.restorePending && + !detail::CleanupA5QueryResources( + query, queryCleanupOps, query.restoreContext)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query context restore remains pending"; + return false; + } + } + impl_->EraseCompletedQueries(); + aclrtContext previous = nullptr; if (aclrtGetCurrentContext(&previous) != ACL_SUCCESS) { TILEXR_LOG(WARN) << "TileXR A5 SDMA could not capture current context for cleanup"; return false; } + for (size_t reverse = impl_->pendingQueries.size(); reverse > 0U; --reverse) { + detail::A5PendingQueryCleanup& query = impl_->pendingQueries[reverse - 1U]; + if (!detail::CleanupA5QueryResources(query, queryCleanupOps, previous)) { + impl_->EraseCompletedQueries(); + TILEXR_LOG(WARN) << "TileXR A5 SDMA query cleanup incomplete; retained for retry"; + return false; + } + } + impl_->EraseCompletedQueries(); + if (impl_->ownerContext == nullptr || aclrtSetCurrentContext(impl_->ownerContext) != ACL_SUCCESS) { TILEXR_LOG(WARN) << "TileXR A5 SDMA could not switch to owner context for cleanup"; @@ -675,7 +730,8 @@ bool TileXRA5SDMABackend::Shutdown() Impl::OwnedChannel& owned = impl_->channels[reverse - 1U]; bool mappingReleased = true; if (owned.mapped) { - if (halResAddrUnmap(impl_->physicalDevice, &owned.mapInfo) == DRV_ERROR_NONE) { + if (impl_->api.unmapResource( + impl_->physicalDevice, &owned.mapInfo) == DRV_ERROR_NONE) { owned.mapped = false; } else { mappingReleased = false; diff --git a/src/comm/sdma/tilexr_sdma_a5_cleanup.h b/src/comm/sdma/tilexr_sdma_a5_cleanup.h new file mode 100644 index 00000000..b5098021 --- /dev/null +++ b/src/comm/sdma/tilexr_sdma_a5_cleanup.h @@ -0,0 +1,173 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_A5_CLEANUP_H +#define TILEXR_SDMA_A5_CLEANUP_H + +#include +#include + +namespace TileXR { +namespace detail { + +using A5CleanupFn = int (*)(void*, void*); +using A5ConstCleanupFn = int (*)(void*, const void*); + +struct A5QueryCleanupOps { + void* opaque = nullptr; + A5CleanupFn setCurrentContext = nullptr; + A5CleanupFn destroyStream = nullptr; + A5CleanupFn destroyContext = nullptr; + A5CleanupFn freeDevice = nullptr; + A5ConstCleanupFn destroyTensor = nullptr; +}; + +struct A5PendingQueryCleanup { + void* ownerContext = nullptr; + void* isolatedContext = nullptr; + void* queryStream = nullptr; + void* healthStream = nullptr; + std::vector ownerBuffers; + std::vector isolatedBuffers; + std::vector tensors; + void* restoreContext = nullptr; + bool restorePending = false; + + bool Empty() const + { + return isolatedContext == nullptr && queryStream == nullptr && + healthStream == nullptr && ownerBuffers.empty() && + isolatedBuffers.empty() && tensors.empty() && !restorePending; + } +}; + +inline bool A5SetCleanupContext(const A5QueryCleanupOps& ops, void* context) +{ + return ops.setCurrentContext != nullptr && + ops.setCurrentContext(ops.opaque, context) == 0; +} + +inline bool A5ReleaseCleanupHandle(void*& handle, A5CleanupFn cleanup, + const A5QueryCleanupOps& ops) +{ + if (handle == nullptr) { + return true; + } + if (cleanup != nullptr && cleanup(ops.opaque, handle) == 0) { + handle = nullptr; + return true; + } + return false; +} + +inline bool A5ReleaseCleanupHandles(std::vector& handles, + A5CleanupFn cleanup, + const A5QueryCleanupOps& ops) +{ + bool released = true; + size_t reverse = handles.size(); + while (reverse > 0U) { + --reverse; + if (handles[reverse] != nullptr && cleanup != nullptr && + cleanup(ops.opaque, handles[reverse]) == 0) { + handles.erase(handles.begin() + static_cast(reverse)); + } else { + released = false; + } + } + return released; +} + +inline bool A5ReleaseCleanupTensors(std::vector& tensors, + const A5QueryCleanupOps& ops) +{ + bool released = true; + size_t reverse = tensors.size(); + while (reverse > 0U) { + --reverse; + if (tensors[reverse] != nullptr && ops.destroyTensor != nullptr && + ops.destroyTensor(ops.opaque, tensors[reverse]) == 0) { + tensors.erase(tensors.begin() + static_cast(reverse)); + } else { + released = false; + } + } + return released; +} + +inline bool CleanupA5QueryResources(A5PendingQueryCleanup& state, + const A5QueryCleanupOps& ops, + void* callerContext) +{ + bool released = true; + if (state.restorePending) { + if (!A5SetCleanupContext(ops, state.restoreContext)) { + return false; + } + state.restoreContext = nullptr; + state.restorePending = false; + } + + if (state.healthStream != nullptr) { + if (state.ownerContext == nullptr || + !A5SetCleanupContext(ops, state.ownerContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandle( + state.healthStream, ops.destroyStream, ops) && released; + } + } + + const bool hasIsolatedResources = state.queryStream != nullptr || + !state.isolatedBuffers.empty() || !state.tensors.empty(); + if (state.isolatedContext != nullptr) { + if (!A5SetCleanupContext(ops, state.isolatedContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandle( + state.queryStream, ops.destroyStream, ops) && released; + if (state.queryStream == nullptr) { + released = A5ReleaseCleanupHandles( + state.isolatedBuffers, ops.freeDevice, ops) && released; + released = A5ReleaseCleanupTensors(state.tensors, ops) && released; + } + if (state.queryStream == nullptr && state.tensors.empty() && + state.isolatedBuffers.empty()) { + released = A5ReleaseCleanupHandle( + state.isolatedContext, ops.destroyContext, ops) && released; + } + } + } else if (hasIsolatedResources) { + released = false; + } + + const bool isolatedReleased = state.isolatedContext == nullptr && + state.queryStream == nullptr && state.isolatedBuffers.empty() && + state.tensors.empty(); + if (!state.ownerBuffers.empty() && state.healthStream == nullptr && + isolatedReleased) { + if (state.ownerContext == nullptr || + !A5SetCleanupContext(ops, state.ownerContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandles( + state.ownerBuffers, ops.freeDevice, ops) && released; + } + } else if (!state.ownerBuffers.empty()) { + released = false; + } + + if (!A5SetCleanupContext(ops, callerContext)) { + state.restoreContext = callerContext; + state.restorePending = true; + released = false; + } + return released && state.Empty(); +} + +} // namespace detail +} // namespace TileXR + +#endif // TILEXR_SDMA_A5_CLEANUP_H diff --git a/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp b/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp index 572af880..b76c9a92 100644 --- a/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp @@ -1,6 +1,9 @@ #include #include +#include +#include +#include "sdma/tilexr_sdma_a5_cleanup.h" #include "sdma/tilexr_sdma_a5_backend.h" namespace { @@ -36,6 +39,105 @@ TileXR::detail::A5HostChannelIdentity ValidIdentity() return {6U, 3U, 4U, 5U, 7U}; } +void* Handle(uintptr_t value) +{ + return reinterpret_cast(value); +} + +struct FakeCleanupRuntime { + std::string failedOperation; + void* failedHandle = nullptr; + bool failureConsumed = false; + std::vector calls; + + int Call(const char* operation, void* handle) + { + calls.push_back(operation); + if (!failureConsumed && failedOperation == operation && + failedHandle == handle) { + failureConsumed = true; + return 1; + } + return 0; + } +}; + +int FakeSetContext(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("set", handle); +} + +int FakeDestroyStream(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("stream", handle); +} + +int FakeDestroyContext(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("context", handle); +} + +int FakeFreeDevice(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("free", handle); +} + +int FakeDestroyTensor(void* opaque, const void* handle) +{ + return static_cast(opaque)->Call( + "tensor", const_cast(handle)); +} + +TileXR::detail::A5QueryCleanupOps FakeCleanupOps(FakeCleanupRuntime& runtime) +{ + TileXR::detail::A5QueryCleanupOps ops; + ops.opaque = &runtime; + ops.setCurrentContext = FakeSetContext; + ops.destroyStream = FakeDestroyStream; + ops.destroyContext = FakeDestroyContext; + ops.freeDevice = FakeFreeDevice; + ops.destroyTensor = FakeDestroyTensor; + return ops; +} + +TileXR::detail::A5PendingQueryCleanup FullCleanupState() +{ + TileXR::detail::A5PendingQueryCleanup state; + state.ownerContext = Handle(1U); + state.isolatedContext = Handle(2U); + state.queryStream = Handle(3U); + state.healthStream = Handle(4U); + state.ownerBuffers = {Handle(5U), Handle(6U)}; + state.isolatedBuffers = {Handle(7U)}; + state.tensors = {Handle(8U), Handle(9U)}; + return state; +} + +bool ContainsHandle(const TileXR::detail::A5PendingQueryCleanup& state, + void* handle) +{ + if (state.isolatedContext == handle || state.queryStream == handle || + state.healthStream == handle || state.restoreContext == handle) { + return true; + } + for (void* buffer : state.ownerBuffers) { + if (buffer == handle) { + return true; + } + } + for (void* buffer : state.isolatedBuffers) { + if (buffer == handle) { + return true; + } + } + for (const void* tensor : state.tensors) { + if (tensor == handle) { + return true; + } + } + return false; +} + void TestCompleteAndExpectedPartialClassification() { using namespace TileXR::detail; @@ -90,6 +192,61 @@ void TestCompleteClassificationRequiresFinishedHeader() A5QueryResultKind::INVALID); } +void TestCleanupFailuresRetainHandlesForRetry() +{ + struct FailureCase { + const char* operation; + void* handle; + }; + const std::vector failures = { + {"set", Handle(2U)}, + {"stream", Handle(4U)}, + {"stream", Handle(3U)}, + {"free", Handle(7U)}, + {"tensor", Handle(9U)}, + {"context", Handle(2U)}, + {"free", Handle(6U)}, + }; + for (const FailureCase& failure : failures) { + FakeCleanupRuntime runtime; + runtime.failedOperation = failure.operation; + runtime.failedHandle = failure.handle; + TileXR::detail::A5PendingQueryCleanup state = FullCleanupState(); + const TileXR::detail::A5QueryCleanupOps ops = FakeCleanupOps(runtime); + CHECK_TRUE(!TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(runtime.failureConsumed); + CHECK_TRUE(ContainsHandle(state, failure.handle)); + CHECK_TRUE(!state.Empty()); + + CHECK_TRUE(TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(state.Empty()); + } +} + +void TestCleanupRestoreFailureIsRetryable() +{ + FakeCleanupRuntime runtime; + runtime.failedOperation = "set"; + runtime.failedHandle = Handle(1U); + TileXR::detail::A5PendingQueryCleanup state; + state.ownerContext = Handle(1U); + state.isolatedContext = Handle(2U); + state.queryStream = Handle(3U); + const TileXR::detail::A5QueryCleanupOps ops = FakeCleanupOps(runtime); + + CHECK_TRUE(!TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(state.restorePending); + CHECK_TRUE(state.restoreContext == state.ownerContext); + CHECK_TRUE(!state.Empty()); + + CHECK_TRUE(TileXR::detail::CleanupA5QueryResources( + state, ops, state.restoreContext)); + CHECK_TRUE(state.Empty()); +} + } // namespace int main() @@ -97,6 +254,8 @@ int main() TestCompleteAndExpectedPartialClassification(); TestPartialClassificationFailsClosed(); TestCompleteClassificationRequiresFinishedHeader(); + TestCleanupFailuresRetainHandlesForRetry(); + TestCleanupRestoreFailureIsRetryable(); if (g_failures != 0) { std::cerr << g_failures << " A5 SDMA validation checks failed" << std::endl; return 1; diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index cce95ba2..15825a6f 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -142,6 +142,11 @@ void TestAscend950UsesOwnedDirectBackend() CheckNeedle(backendPath, backend, "RES_ADDR_TYPE_STARS_RTSQ"); CheckNeedle(backendPath, backend, "TILEXR_SDMA_A5_CHANNEL_COUNT"); CheckNeedle(backendPath, backend, "kExpectedAicpuQueryFailure"); + CheckNeedle(backendPath, backend, + "LoadSymbol(RTLD_DEFAULT, \"halResAddrMap\", mapResource)"); + CheckNeedle(backendPath, backend, "pendingQueries"); + CheckNoNeedle(backendPath, backend, "halResAddrMap("); + CheckNoNeedle(backendPath, backend, "halResAddrUnmap("); const std::string devicePath = "src/include/tilexr_sdma_a5.h"; const auto device = ReadFile(devicePath); From 42ec9c76e0b6bccd49733520bd4608699e940bd9 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Wed, 29 Jul 2026 15:33:36 +0800 Subject: [PATCH 3/7] fix: make A5 SDMA cleanup safe --- docs/SDMA_TRANSPORT.md | 5 ++--- src/comm/tilexr_comm.cpp | 5 ++++- src/include/tilexr_sdma_a5.h | 10 +--------- src/include/tilexr_sdma_a5_types.h | 1 - tests/sdma/demo/tilexr_sdma_demo.cpp | 4 +--- tests/sdma/unit/test_tilexr_sdma_metadata.cpp | 1 - tests/sdma/unit/test_tilexr_sdma_source_guard.cpp | 5 +++-- 7 files changed, 11 insertions(+), 20 deletions(-) diff --git a/docs/SDMA_TRANSPORT.md b/docs/SDMA_TRANSPORT.md index 6bb86e10..89d7d8dd 100644 --- a/docs/SDMA_TRANSPORT.md +++ b/docs/SDMA_TRANSPORT.md @@ -92,9 +92,8 @@ During communicator initialization, the Host backend: The AIV backend atomically claims a channel, writes a data-copy SQE followed by a 64-byte completion-copy SQE, cleans the touched cache lines, orders the stores, and updates the mapped RTSQ tail. `SDMAWait` validates the event generation, -polls the completion record with a fixed upper bound, and releases the channel -only after completion. A timed-out channel remains busy so later submissions -cannot reuse an uncertain queue. +waits for the matching completion record, and releases the channel only after +completion. No TileXR target builds or installs a custom OPP, and this feature needs no custom OPP environment setting. `reference/` remains comparison-only and is not diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp index 72300c61..d06696c4 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -245,6 +245,9 @@ bool TileXRComm::ResetSDMAState() { commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; + if (UpdateCommArgsDev() != TILEXR_SUCCESS) { + return false; + } sdmaWorkspaceDev_ = nullptr; sdmaInitStatus_ = SDMAInitStatus::DISABLED_BY_ENV; if (sdmaTransport_ != nullptr) { @@ -893,8 +896,8 @@ TileXRComm::~TileXRComm() } FreePeerMem(commArgs_.dumpAddr); FreePeerMem(peerMem_[rank_]); - FreePeerMem(commArgsPtr_); (void)ResetSDMAState(); + FreePeerMem(commArgsPtr_); } TileXRComm::TileXRComm(int rank, int rankSize) : rank_(rank), rankSize_(rankSize) diff --git a/src/include/tilexr_sdma_a5.h b/src/include/tilexr_sdma_a5.h index 65c19c6a..63c564eb 100644 --- a/src/include/tilexr_sdma_a5.h +++ b/src/include/tilexr_sdma_a5.h @@ -175,15 +175,7 @@ __aicore__ inline bool A5SdmaWaitEvent(__gm__ uint8_t* workspaceAddress, } __gm__ A5SdmaCompletionLine* completion = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( channel->completionRecordAddress); - bool completed = false; - for (uint32_t poll = 0U; poll < TILEXR_SDMA_A5_WAIT_MAX_POLLS; ++poll) { - if (A5SdmaReadCompletion(completion) == generation) { - completed = true; - break; - } - } - if (!completed) { - return false; + while (A5SdmaReadCompletion(completion) != generation) { } dsb(DSB_DDR); if (channel->generation != generation) { diff --git a/src/include/tilexr_sdma_a5_types.h b/src/include/tilexr_sdma_a5_types.h index 2a75a084..507f082d 100644 --- a/src/include/tilexr_sdma_a5_types.h +++ b/src/include/tilexr_sdma_a5_types.h @@ -28,7 +28,6 @@ constexpr uint64_t TILEXR_SDMA_A5_MAX_TRANSFER_BYTES = 0xFFFFFFFFULL; constexpr uint32_t TILEXR_SDMA_A5_SQE_TYPE = 11U; constexpr uint32_t TILEXR_SDMA_A5_KERNEL_CREDIT = 254U; constexpr uint32_t TILEXR_SDMA_A5_QOS = 6U; -constexpr uint32_t TILEXR_SDMA_A5_WAIT_MAX_POLLS = 1000000U; constexpr uint64_t TILEXR_SDMA_A5_EVENT_MAGIC = 0xA5D5ULL; constexpr uint32_t TILEXR_SDMA_A5_EVENT_MAGIC_SHIFT = 48U; diff --git a/tests/sdma/demo/tilexr_sdma_demo.cpp b/tests/sdma/demo/tilexr_sdma_demo.cpp index a7087cbc..cadbf0e1 100644 --- a/tests/sdma/demo/tilexr_sdma_demo.cpp +++ b/tests/sdma/demo/tilexr_sdma_demo.cpp @@ -35,7 +35,6 @@ constexpr uint32_t kDefaultBytes = 4096; constexpr uint32_t kAlignmentBytes = 64; constexpr size_t kDebugWordsPerBlock = 16; constexpr int kDeviceId = 0; -constexpr int32_t kStreamTimeoutMs = 60000; #ifndef TILEXR_SDMA_DEMO_A5 #define TILEXR_SDMA_DEMO_A5 0 @@ -264,8 +263,7 @@ int main(int argc, char** argv) launch_tilexr_sdma_copy( blocks, stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), reinterpret_cast(debug), bytes, firstChannel, iterations); - if (!CheckAcl("aclrtSynchronizeStreamWithTimeout", - aclrtSynchronizeStreamWithTimeout(stream, kStreamTimeoutMs))) { + if (!CheckAcl("aclrtSynchronizeStream", aclrtSynchronizeStream(stream))) { Cleanup(src, dst, debug, comm, stream, deviceSet, aclInitialized); return 1; } diff --git a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp index be247c78..c4b7db1f 100644 --- a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp @@ -56,7 +56,6 @@ void TestA5WorkspaceAbi() { using namespace TileXR::detail; CHECK_EQ(TILEXR_SDMA_A5_CHANNEL_COUNT, 48U); - CHECK_TRUE(TILEXR_SDMA_A5_WAIT_MAX_POLLS > 0U); CHECK_EQ(sizeof(A5SdmaWorkspaceHeader), 64U); CHECK_EQ(sizeof(A5SdmaChannel), 192U); CHECK_EQ(sizeof(A5SdmaCompletionLine), 64U); diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index 15825a6f..c3a0b03b 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -150,11 +150,11 @@ void TestAscend950UsesOwnedDirectBackend() const std::string devicePath = "src/include/tilexr_sdma_a5.h"; const auto device = ReadFile(devicePath); - CheckNeedle(devicePath, device, "TILEXR_SDMA_A5_WAIT_MAX_POLLS"); + CheckNeedle(devicePath, device, "while (A5SdmaReadCompletion(completion) != generation)"); const std::string demoPath = "tests/sdma/demo/tilexr_sdma_demo.cpp"; const auto demo = ReadFile(demoPath); - CheckNeedle(demoPath, demo, "aclrtSynchronizeStreamWithTimeout"); + CheckNeedle(demoPath, demo, "aclrtSynchronizeStream(stream)"); } void TestSdmaFailureIsScopedToOneCommunicator() @@ -164,6 +164,7 @@ void TestSdmaFailureIsScopedToOneCommunicator() CheckNoNeedle(commPath, comm, "g_sdmaUnavailable"); CheckNeedle(commPath, comm, "sdmaTransport_->Init(options)"); CheckNeedle(commPath, comm, "bool TileXRComm::PrepareDestroy()"); + CheckNeedle(commPath, comm, "if (UpdateCommArgsDev() != TILEXR_SUCCESS)"); const std::string wrapPath = "src/comm/comm_wrap.cpp"; const auto wrap = ReadFile(wrapPath); From df0db58b56b255c3dd12b154586186c86afb614f Mon Sep 17 00:00:00 2001 From: Kur0x Date: Thu, 30 Jul 2026 09:51:35 +0800 Subject: [PATCH 4/7] feat: add A5 SDMA strided batch submission --- docs/SDMA_TRANSPORT.md | 54 +++ ...026-07-28-a5-direct-sdma-implementation.md | 71 +++- .../specs/2026-07-28-a5-direct-sdma-design.md | 102 ++++- src/include/tilexr_sdma.h | 35 ++ src/include/tilexr_sdma_a5.h | 98 +++-- src/include/tilexr_sdma_a5_types.h | 56 ++- tests/sdma/CMakeLists.txt | 29 +- tests/sdma/demo/tilexr_sdma_benchmark.cpp | 362 ++++++++++++++++++ tests/sdma/demo/tilexr_sdma_benchmark_types.h | 38 ++ tests/sdma/demo/tilexr_sdma_demo_kernel.cpp | 173 +++++++++ tests/sdma/unit/test_tilexr_sdma_metadata.cpp | 18 + .../unit/test_tilexr_sdma_source_guard.cpp | 11 + 12 files changed, 983 insertions(+), 64 deletions(-) create mode 100644 tests/sdma/demo/tilexr_sdma_benchmark.cpp create mode 100644 tests/sdma/demo/tilexr_sdma_benchmark_types.h diff --git a/docs/SDMA_TRANSPORT.md b/docs/SDMA_TRANSPORT.md index 89d7d8dd..7d3b9069 100644 --- a/docs/SDMA_TRANSPORT.md +++ b/docs/SDMA_TRANSPORT.md @@ -51,6 +51,10 @@ mappings. uint64_t event = TileXR::SDMACopyNbi(args, dst, src, bytes, channel); bool ok = TileXR::SDMAWait(args, event, channel); + +uint64_t batchEvent = TileXR::SDMACopyStridedNbi( + args, dst, src, bytes, copyCount, dstStrideBytes, srcStrideBytes, channel); +bool batchOk = TileXR::SDMAWait(args, batchEvent, channel); ``` The API accepts same-device GM pointers. It does not register memory or validate @@ -61,6 +65,12 @@ buffer ownership. Event 0 is the unavailable or invalid no-op result, and provides 48 channels numbered 0 through 47. Each channel permits one outstanding event; different channels can progress concurrently. +`SDMACopyStridedNbi` submits equal-sized, non-overlapping strided slices. On A5, +`copyCount > 1` writes one data SQE per slice, appends one completion SQE, and +rings one doorbell. The returned event completes the whole batch. A2/A3 retains +the existing PTO single-copy path; multi-copy strided submission returns event +0 there. `copyCount == 1` delegates to `SDMACopyNbi` on every backend. + ## A2/A3 PTO Backend On A2/A3, `TileXRSDMATransport` owns a PTO @@ -95,6 +105,12 @@ and updates the mapped RTSQ tail. `SDMAWait` validates the event generation, waits for the matching completion record, and releases the channel only after completion. +For a strided batch, the AIV backend reserves `copyCount + 1` SQ entries, +advances consecutive task IDs, and reuses the same completion/wait path. Queue +capacity, count, length, stride, and 64-bit address arithmetic are validated +before the channel is claimed. Multi-copy source and destination strides must +each be at least the copy length. + No TileXR target builds or installs a custom OPP, and this feature needs no custom OPP environment setting. `reference/` remains comparison-only and is not an include, source, or link dependency. @@ -151,6 +167,37 @@ Every run must report a nonzero event, a successful wait, the expected generation, and a full byte-for-byte destination comparison. A skipped or unavailable backend is not an A5 data-plane pass. +### Batch performance and selection + +The A5-only benchmark accepts: + +```text +tilexr_sdma_benchmark \ + [aligned-bytes [warmup-batches [measured-batches [samples \ + [channel [phase-profile [batch-copies]]]]]]] +``` + +The following device-cycle measurements use a rotating 64 MiB working set and +the same `ACL_MEM_MALLOC_HUGE_FIRST` backing size for single and batch runs. +Each result includes SQE construction, cache maintenance, doorbell, device +completion, and channel release, while excluding Host launch and allocation. + +| Bytes per copy | Single completion | Batch 16 amortized | Speedup | +| ---: | ---: | ---: | ---: | +| 4 KiB | 3.351 us | 0.954 us | 3.51x | +| 8 KiB | 3.372 us | 0.975 us | 3.46x | +| 16 KiB | 3.415 us | 0.972 us | 3.51x | +| 32 KiB | 4.133 us | 1.031 us | 4.01x | +| 64 KiB | 3.453 us | 1.012 us | 3.41x | +| 1 MiB | 4.314 us | 1.962 us | 2.20x | +| 4 MiB | 8.752 us | 6.313 us | 1.39x | + +Use batching for naturally available 4-64 KiB work. From 64 KiB through +1 MiB it remains throughput-beneficial, but do not wait to fill a batch. Above +1 MiB, batch only already-available parallel copies; send latency-sensitive +isolated copies immediately. Batch values are amortized: all copies share one +event that becomes visible only after the complete batch. + ## Validation Record Production A5 acceptance was exercised on: @@ -164,6 +211,13 @@ Production A5 acceptance was exercised on: - channels 0 through 3: four concurrent AIV blocks at 4 KiB, three iterations; - channel 1: three init/use/destroy repetitions at 4 KiB, three iterations each; - busy-channel behavior: a second outstanding submission returned event 0; +- formal `SDMACopyStridedNbi` batches: counts 1, 2, 4, 8, 16, and 32 at + 4 KiB on channel 0, with 20,000 warmup batches and repeated SQ-tail/task-ID + wraparound; +- channel 47: batch 16 at 4 KiB completed at 0.907 us/copy amortized; +- representative batch 16 results through the production API: 4 KiB at + 0.943 us/copy, 64 KiB at 1.044 us/copy, 1 MiB at 1.909 us/copy, and 4 MiB + at 6.315 us/copy; - validation: every successful run compared the full source and destination. The A2/A3 PTO path was previously exercised on Ascend 910B3 with CANN 9.0.0 and diff --git a/docs/plans/2026-07-28-a5-direct-sdma-implementation.md b/docs/plans/2026-07-28-a5-direct-sdma-implementation.md index e1931c47..7786f2d9 100644 --- a/docs/plans/2026-07-28-a5-direct-sdma-implementation.md +++ b/docs/plans/2026-07-28-a5-direct-sdma-implementation.md @@ -3,6 +3,7 @@ Date: 2026-07-28 Status: Approved for execution Design: `docs/specs/2026-07-28-a5-direct-sdma-design.md` +Updated: 2026-07-29 (Task 9 adds approved strided batch submission) ## Goal @@ -10,7 +11,7 @@ Implement the approved Ascend950/A5 direct-SDMA backend inside TileXR while preserving the existing A2/A3 PTO backend and public API. The production path must use CANN's built-in system AICPU `ShmemSdmaStarsQuery`, map Host RTSQ doorbells, publish a TileXR-owned 48-channel workspace, and let AIV submit the -two-SQE data/completion sequence directly. +single-copy two-SQE or strided-batch `N + 1` sequence directly. The project-wide supported baseline becomes CANN 9.1.0 with driver `25.1.rc1` or later. The validated hardware baseline is @@ -22,8 +23,9 @@ The project-wide supported baseline becomes CANN 9.1.0 with driver - Add a production A5 Host backend to `libtile-comm.so`, including dynamic built-in-op loading, 48 STARS streams, query compatibility, RTSQ mappings, workspace publication, and complete failure cleanup. -- Dispatch `SDMACopyNbi` and `SDMAWait` to A5 direct submission for dav-3510 - kernels while keeping PTO behavior on A2/A3. +- Dispatch `SDMACopyNbi`, `SDMACopyStridedNbi`, and `SDMAWait` to A5 direct + submission for dav-3510 kernels while keeping PTO single-copy behavior on + A2/A3. - Extend SDMA unit, integration, build, demo, dependency, and hardware checks. - Change requirements and CI provisioning from driver `25.5.0+` to `25.1.rc1+`, including Huawei RC/build-suffix comparison. @@ -37,7 +39,7 @@ The project-wide supported baseline becomes CANN 9.1.0 with driver - Do not include or link any source under `reference/`; it remains evidence only. - Do not change `CommArgs`, public Host API signatures, the opt-in environment - variable, or event-zero semantics. + variable, workspace ABI, or event-zero semantics. - Do not replace PTO SDMA on A2/A3, make SDMA mandatory for communicator init, add more than one outstanding event per channel, or claim Host/simulator checks prove the A5 data plane. @@ -67,8 +69,10 @@ paths, sources, or link inputs. Tasks 1 and 2 establish contracts needed by Tasks 3 and 4. Task 5 depends on the production Host and device paths. Task 6 is independent of Tasks 2-5 at the file level but must land before final validation. Task 7 follows functional -implementation and baseline migration. Task 8 is the final acceptance gate; -PoC removal is not considered complete until its evidence is covered there. +implementation and baseline migration. Task 8 is the original backend +acceptance gate. Task 9 depends on Tasks 1, 4, 5, and 8 and is the final batch +extension gate; PoC removal is not considered complete until its evidence is +covered there. ## Task 1: Define The A5 ABI And Test Seams @@ -197,7 +201,7 @@ wiring, and dav-3510 compile tests. Keep PTO-specific code in **Constraints and non-goals:** Validate workspace magic/version/backend, addresses, byte length, channel, depth, and queue capacity before writes. Atomically claim the channel; a busy channel returns 0. Advance a nonzero -generation, clear completion, build data and 64-byte completion-copy SQEs with +generation, update the completion payload, build data and 64-byte completion-copy SQEs with type 11, `wrCqe=0`, credit 254, flush all touched cache lines, issue ordering, and ring RTSQ only after both SQEs are ready. `SDMAWait(0)` succeeds; other waits reject bad channel/generation/stale events, poll matching completion, and @@ -343,6 +347,47 @@ iterations in `docs/SDMA_TRANSPORT.md`. validation documentation, and a final diff/status review demonstrating that only scoped production, test, baseline, diagram, and cleanup files changed. +## Task 9: Promote Strided Batch Submission + +**Objective and role:** Move the hardware-proven batch prototype into the +installed TileXR device API so multiple independent same-sized copies can share +one A5 doorbell and one completion event. + +**Background and prerequisites:** Representative A5 measurements show +3.4-4.0x amortized improvement for 4-64 KiB, 2.2x at 1 MiB, and 1.39x at +4 MiB. The accepted design adds `SDMACopyStridedNbi` without changing Host or +workspace ABI. One batch uses `N` data SQEs plus one completion-copy SQE and +retains one outstanding event per channel. + +**Modification scope:** Update `src/include/tilexr_sdma_a5_types.h` with pure +multi-entry queue helpers, implement the A5 strided batch in +`src/include/tilexr_sdma_a5.h`, expose the architecture-routed wrapper in +`src/include/tilexr_sdma.h`, and replace the private implementation in +`tests/sdma/demo/tilexr_sdma_demo_kernel.cpp`. Extend metadata, header/source, +benchmark, and transport documentation tests as appropriate. + +**Constraints and non-goals:** Preserve C++14 and CANN 9.1, the A5 workspace +layout/version, Host APIs, event format, existing single-copy behavior, A2/A3 +PTO behavior, and one outstanding event per channel. Validate all inputs before +claiming a channel. Require non-overlapping strides for multi-copy batches, +reject address arithmetic overflow and insufficient queue capacity, leave one +SQ entry unused, and ring exactly one doorbell. Do not add descriptor arrays, +begin/append/commit state, CQE polling, broadcast/overlap semantics, or a +multi-copy PTO implementation. + +**Acceptance and verification:** Unit-test arbitrary-entry capacity, tail and +task-ID wrap, invalid count/stride/overflow inputs, and unchanged two-entry +single-copy helpers. Compile the installed header and benchmark kernel for +dav-3510. On A5 run counts 1, 2, 4, 8, 16, and 32 across queue/task-ID wraps, +verify every destination slice, repeat init/use/destroy, and rerun 4 KiB, +64 KiB, 1 MiB, and 4 MiB representative timings with equal 64 MiB backing and +a rotating working set. Run the eight focused SDMA test executables and inspect +the final dependency/RPATH state. + +**Artifacts and interfaces:** Installed `SDMACopyStridedNbi`, reusable A5 +multi-entry queue helpers, benchmark coverage using only the production API, +and recorded hardware semantics/performance evidence. + ## Key Risks And Controls - **Expected AICPU failure corrupts the active context:** every query runs in a @@ -351,11 +396,11 @@ only scoped production, test, baseline, diagram, and cleanup files changed. - **Partial workspace is accidentally trusted:** acceptance requires the exact status and every independent field/ID/HAL/mapping cross-check; all other partial states fail closed. -- **Queue overwrite or cross-block race:** reserve two slots, permit one - outstanding event per channel, atomically claim it, and assign distinct - channels to concurrent blocks. -- **Cache/order mismatch:** flush payload, completion record, and both SQEs; - issue the validated barrier sequence before the RTSQ tail write. +- **Queue overwrite or cross-block race:** reserve `N + 1` slots while leaving + one unused, permit one outstanding event per channel, atomically claim it, + and assign distinct channels to concurrent blocks. +- **Cache/order mismatch:** flush the payload and every written SQE; issue the + validated barrier sequence before the RTSQ tail write. - **ABI drift between Host and AIV:** one installed fixed-width header and static assertions are compiled in both modes. - **Optional runtime becomes a hard dependency:** dynamically resolve the @@ -367,7 +412,7 @@ only scoped production, test, baseline, diagram, and cleanup files changed. ## Completion Gate -The change is complete only when all eight tasks pass, A2/A3 PTO selection is +The change is complete only when all nine tasks pass, A2/A3 PTO selection is preserved, A5 device 5 passes the full hardware matrix with 48 live channels, driver `25.1.rc1.b188` is accepted by the shared comparator, dependency checks are clean, and the repository contains neither the PoC nor a TileXR OPP path. diff --git a/docs/specs/2026-07-28-a5-direct-sdma-design.md b/docs/specs/2026-07-28-a5-direct-sdma-design.md index 1d7375c7..907c3fe0 100644 --- a/docs/specs/2026-07-28-a5-direct-sdma-design.md +++ b/docs/specs/2026-07-28-a5-direct-sdma-design.md @@ -2,6 +2,7 @@ Date: 2026-07-28 Status: Approved on 2026-07-28 +Updated: 2026-07-29 (strided batch submission approved for implementation) ## Goal @@ -12,7 +13,8 @@ RTSQ doorbells on Host, and lets AIV write and submit A5 SDMA SQEs directly. The implementation must: - support CANN 9.1.0 and NPU driver `25.1.rc1` or later; -- preserve the existing A2/A3 PTO backend and public API; +- preserve the existing A2/A3 PTO backend and existing API behavior; +- add an A5 strided batch device API without changing Host or workspace ABI; - remain best-effort when any A5 capability is unavailable; - ship no TileXR OPP package and require no custom OPP environment variable; - keep `reference/` comparison-only. @@ -49,6 +51,32 @@ The following contracts remain unchanged: - device calls `SDMACopyNbi` and `SDMAWait`; - event handle 0 as the unavailable or invalid no-op result. +The installed device header additionally exposes: + +```cpp +uint64_t SDMACopyStridedNbi( + const __gm__ CommArgs* args, + __gm__ uint8_t* dst, + __gm__ uint8_t* src, + uint64_t bytes, + uint32_t copyCount, + uint64_t dstStrideBytes, + uint64_t srcStrideBytes, + uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP); +``` + +For copy `i`, the source and destination are `src + i * srcStrideBytes` and +`dst + i * dstStrideBytes`. `copyCount == 1` delegates to `SDMACopyNbi` and +preserves its behavior. On A5, `copyCount > 1` submits one ordered batch. On +A2/A3, multi-copy batch submission is unavailable and returns event 0; the +existing PTO single-copy path is unchanged. + +For `copyCount > 1`, both strides must be at least `bytes`, every individual +length must fit the A5 SQE, address arithmetic must not overflow, and the +channel must have room for `copyCount + 1` SQEs. These conservative rules keep +all source and destination slices distinct. The API does not promise broadcast +or overlapping-copy semantics. + Host selects the backend by runtime SoC: - A2/A3 continues to use PTO `SdmaWorkspaceManager` and PTO device intrinsics; @@ -138,14 +166,15 @@ range rule. 1. Validate the workspace ABI, pointers, byte count, and channel group. 2. Atomically claim the channel's outstanding state and advance its nonzero generation. A busy channel returns event 0. -3. Clear the completion record and write the generation into the channel's - completion payload. +3. Write the generation into the channel's completion payload. The completion + record is not cleared because generations are monotonic and the completion + SQE overwrites the full record before it can match the new generation. 4. Build two consecutive 64-byte A5 SQEs at the current tail: one data copy and one 64-byte copy from the payload to the completion record. 5. Set SQE type 11, `wrCqe=0`, kernel credit 254, and the validated A5 address, substream, length, and task fields. -6. Clean the written payload, completion record, and SQE cache lines, execute the - required store-ordering barrier, then write the new tail to the mapped RTSQ. +6. Clean the written payload and SQE cache lines, execute the required + store-ordering barrier, then write the new tail to the mapped RTSQ. 7. Return an event encoding the channel and generation. The tail advances modulo the queried depth. Completion of the second SQE makes @@ -153,8 +182,59 @@ the channel safe for reuse. Transfer length handling must respect the A5 SQE field width; unsupported lengths return event 0 instead of truncating. `SDMAWait` rejects malformed or stale events, polls the channel completion -record for the matching generation, then releases the outstanding state. A -zero event remains an immediate successful no-op. +generation through `ReadGmByPassDCache`, then releases the outstanding state. +A zero event remains an immediate successful no-op. + +### Strided batch submission + +`SDMACopyStridedNbi` uses the same channel claim, event encoding, completion +record, and wait path. After all validation succeeds, it: + +1. reserves `copyCount + 1` entries while leaving one SQ slot unused; +2. writes `copyCount` ordered data SQEs with consecutive task IDs; +3. appends one 64-byte completion-copy SQE; +4. cleans the payload and every written SQE cache line; +5. advances the tail and task ID by `copyCount + 1` modulo their hardware + widths; +6. rings one doorbell and returns one event for the complete batch. + +No validation failure after the channel claim is allowed. All pointer, stride, +length, count, queue-capacity, and workspace checks therefore occur before the +atomic claim. `SDMAWait` releases the channel only after the final completion +SQE, so every data SQE in the batch is complete when the event becomes visible. + +The returned event represents the whole batch. Reported per-copy latency is +total submit-through-wait time divided by `copyCount`; callers cannot observe +or wait for an individual copy through this event. + +## Batch Performance Evidence + +The batch decision was validated on `Ascend950PR_9589`, CANN 9.1.0, driver +`25.1.rc1.b188`, physical device 5. The benchmark used AIV +`GetSystemCycle()` at 1000 cycles/us, a rotating 64 MiB source/destination +working set with the same allocation policy for single and batch runs, device +warmup before sampling, ten samples, and full working-set byte comparison. +Host launch, initialization, allocation, H2D/D2H, and stream synchronization +were outside the timed region. SQE construction, cache maintenance, doorbell, +device completion, and channel release were included. + +| Bytes per copy | Single completion | Batch 16 amortized | Speedup | +| ---: | ---: | ---: | ---: | +| 4 KiB | 3.351 us | 0.954 us | 3.51x | +| 8 KiB | 3.372 us | 0.975 us | 3.46x | +| 16 KiB | 3.415 us | 0.972 us | 3.51x | +| 32 KiB | 4.133 us | 1.031 us | 4.01x | +| 64 KiB | 3.453 us | 1.012 us | 3.41x | +| 1 MiB | 4.314 us | 1.962 us | 2.20x | +| 4 MiB | 8.752 us | 6.313 us | 1.39x | + +For 4-64 KiB, callers should batch naturally available copies because fixed +submission and completion costs dominate. From 64 KiB through 1 MiB, batching +still materially improves throughput but callers must not delay an otherwise +ready copy merely to fill a batch. Above 1 MiB, batching is optional: use it +for already-available parallel work, while latency-sensitive isolated copies +should remain single submissions. For example, 16 copies at an amortized +0.954 us still expose one batch completion after roughly 15.3 us. ## Failure and Cleanup @@ -169,8 +249,8 @@ initialization. It never resets a device or destroys an application-owned stream or context. Device submission failures are contained to the SDMA call. They must not write -a doorbell after failed validation or expose a nonzero event before both SQEs -are ready. +a doorbell after failed validation or expose a nonzero event before every SQE +in the submission is ready. ## Driver Baseline Migration @@ -198,6 +278,10 @@ Implementation acceptance requires: - no TileXR OPP artifact, vendor install directory, or custom OPP environment; - 48 distinct, simultaneously live, validated channels on the baseline host; - direct copies of 64 B, 4 KiB, and 1 MiB on channel 0 and channel 47; +- strided batches with counts 1, 2, 4, 8, 16, and 32, including SQ/tail/task-ID + wraparound and full comparison of every destination slice; +- representative 4 KiB through 4 MiB batch performance using equal backing + allocation and a rotating working set; - a concurrent multi-block run using distinct channels, followed by full byte comparisons and matching completion generations; - repeated init/use/shutdown loops with a healthy communicator context; diff --git a/src/include/tilexr_sdma.h b/src/include/tilexr_sdma.h index fd5b6035..ac12be91 100644 --- a/src/include/tilexr_sdma.h +++ b/src/include/tilexr_sdma.h @@ -89,6 +89,41 @@ __aicore__ inline uint64_t SDMACopyNbi( #endif } +__aicore__ inline uint64_t SDMACopyStridedNbi( + const __gm__ CommArgs* args, + __gm__ uint8_t* dst, + __gm__ uint8_t* src, + uint64_t bytes, + uint32_t copyCount, + uint64_t dstStrideBytes, + uint64_t srcStrideBytes, + uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP) +{ + if (copyCount == 1U) { + return SDMACopyNbi(args, dst, src, bytes, channelGroupIdx); + } +#if TILEXR_SDMA_A5_AICORE_COMPILE + if (!SDMAEnabled(args) || dst == nullptr || src == nullptr || bytes == 0U || + copyCount == 0U) { + return 0ULL; + } + const uint32_t resolvedGroup = SDMAResolveChannelGroup(channelGroupIdx); + return detail::A5SdmaCopyStridedNbi( + reinterpret_cast<__gm__ uint8_t*>(args->sdmaWorkspacePtr), + dst, src, bytes, copyCount, dstStrideBytes, srcStrideBytes, resolvedGroup); +#else + (void)args; + (void)dst; + (void)src; + (void)bytes; + (void)copyCount; + (void)dstStrideBytes; + (void)srcStrideBytes; + (void)channelGroupIdx; + return 0ULL; +#endif +} + __aicore__ inline bool SDMAWait( const __gm__ CommArgs* args, uint64_t eventHandle, diff --git a/src/include/tilexr_sdma_a5.h b/src/include/tilexr_sdma_a5.h index 63c564eb..0592593d 100644 --- a/src/include/tilexr_sdma_a5.h +++ b/src/include/tilexr_sdma_a5.h @@ -65,13 +65,7 @@ __aicore__ inline void A5SdmaRingDoorbell(uint64_t address, uint32_t tail) __aicore__ inline uint32_t A5SdmaReadCompletion(__gm__ A5SdmaCompletionLine* completion) { - __ubuf__ uint32_t line[TILEXR_SDMA_A5_COMPLETION_BYTES / sizeof(uint32_t)]; - copy_gm_to_ubuf_align_v2( - line, reinterpret_cast<__gm__ uint32_t*>(completion), - 0, 1, TILEXR_SDMA_A5_COMPLETION_BYTES, 0, 0, false, 0, 0, 0); - set_flag(PIPE_MTE2, PIPE_S, static_cast(0)); - wait_flag(PIPE_MTE2, PIPE_S, static_cast(0)); - return line[0]; + return AscendC::ReadGmByPassDCache(&completion->generation); } __aicore__ inline bool A5SdmaWorkspaceValid(const __gm__ A5SdmaWorkspace* workspace) @@ -86,24 +80,40 @@ __aicore__ inline bool A5SdmaWorkspaceValid(const __gm__ A5SdmaWorkspace* worksp workspace->header.workspaceSize == sizeof(A5SdmaWorkspace); } -__aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, - __gm__ uint8_t* destination, - __gm__ uint8_t* source, - uint64_t bytes, - uint32_t channelIndex) +__aicore__ inline uint64_t A5SdmaCopyStridedNbi( + __gm__ uint8_t* workspaceAddress, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint64_t bytes, + uint32_t copyCount, + uint64_t destinationStrideBytes, + uint64_t sourceStrideBytes, + uint32_t channelIndex) { __gm__ A5SdmaWorkspace* workspace = reinterpret_cast<__gm__ A5SdmaWorkspace*>(workspaceAddress); if (!A5SdmaWorkspaceValid(workspace) || destination == nullptr || source == nullptr || - !A5SdmaTransferLengthValid(bytes) || !A5SdmaChannelValid(channelIndex)) { + !A5SdmaTransferLengthValid(bytes) || !A5SdmaChannelValid(channelIndex) || + copyCount == 0U || copyCount == 0xFFFFFFFFU) { + return 0ULL; + } + const uint64_t destinationAddress = reinterpret_cast(destination); + const uint64_t sourceAddress = reinterpret_cast(source); + if (copyCount > 1U && + (!A5SdmaStridedRangeValid( + destinationAddress, bytes, copyCount, destinationStrideBytes) || + !A5SdmaStridedRangeValid( + sourceAddress, bytes, copyCount, sourceStrideBytes))) { return 0ULL; } __gm__ A5SdmaChannel* channel = &workspace->channels[channelIndex]; + const uint32_t requiredEntries = copyCount + 1U; if (channel->sqBase == 0U || channel->rtsqAddress == 0U || channel->completionPayloadAddress == 0U || channel->completionRecordAddress == 0U || channel->rtsqLength < sizeof(uint32_t) || - !A5SdmaQueueHasCapacity(channel->head, channel->tail, channel->depth) || + !A5SdmaQueueHasEntriesCapacity( + channel->head, channel->tail, channel->depth, requiredEntries) || channel->streamId > 0xFFFFU) { return 0ULL; } @@ -112,39 +122,43 @@ __aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, } const uint32_t generation = A5SdmaNextGeneration(channel->generation); - const uint32_t dataIndex = channel->tail; - const uint32_t completionIndex = (dataIndex + 1U) % channel->depth; - const uint32_t newTail = A5SdmaAdvanceTail(dataIndex, channel->depth); + const uint32_t firstIndex = channel->tail; __gm__ A5SdmaCompletionLine* payload = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( channel->completionPayloadAddress); - __gm__ A5SdmaCompletionLine* completion = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( - channel->completionRecordAddress); - __gm__ uint32_t* payloadWords = reinterpret_cast<__gm__ uint32_t*>(payload); - __gm__ uint32_t* completionWords = reinterpret_cast<__gm__ uint32_t*>(completion); - for (uint32_t index = 0U; index < TILEXR_SDMA_A5_COMPLETION_BYTES / sizeof(uint32_t); ++index) { - payloadWords[index] = 0U; - completionWords[index] = 0U; - } payload->generation = generation; __gm__ A5SdmaSqe* sqBase = reinterpret_cast<__gm__ A5SdmaSqe*>(channel->sqBase); - A5SdmaBuildSqe(sqBase + dataIndex, channel->streamId, channel->taskId, - reinterpret_cast(source), - reinterpret_cast(destination), - static_cast(bytes)); - A5SdmaBuildSqe(sqBase + completionIndex, channel->streamId, channel->taskId + 1U, - channel->completionPayloadAddress, - channel->completionRecordAddress, - TILEXR_SDMA_A5_COMPLETION_BYTES); + for (uint32_t copy = 0U; copy < copyCount; ++copy) { + const uint32_t sqeIndex = A5SdmaAdvanceTailBy(firstIndex, copy, channel->depth); + const uint64_t sourceOffset = static_cast(copy) * sourceStrideBytes; + const uint64_t destinationOffset = + static_cast(copy) * destinationStrideBytes; + A5SdmaBuildSqe( + sqBase + sqeIndex, channel->streamId, + A5SdmaAdvanceTaskIdBy(channel->taskId, copy), + sourceAddress + sourceOffset, destinationAddress + destinationOffset, + static_cast(bytes)); + } + const uint32_t completionIndex = + A5SdmaAdvanceTailBy(firstIndex, copyCount, channel->depth); + A5SdmaBuildSqe( + sqBase + completionIndex, channel->streamId, + A5SdmaAdvanceTaskIdBy(channel->taskId, copyCount), + channel->completionPayloadAddress, + channel->completionRecordAddress, + TILEXR_SDMA_A5_COMPLETION_BYTES); + const uint32_t newTail = + A5SdmaAdvanceTailBy(firstIndex, requiredEntries, channel->depth); channel->generation = generation; channel->tail = newTail; - channel->taskId = A5SdmaAdvanceTaskId(channel->taskId); + channel->taskId = A5SdmaAdvanceTaskIdBy(channel->taskId, requiredEntries); pipe_barrier(PIPE_ALL); A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(payload)); - A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(completion)); - A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + dataIndex)); - A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + completionIndex)); + for (uint32_t entry = 0U; entry < requiredEntries; ++entry) { + const uint32_t sqeIndex = A5SdmaAdvanceTailBy(firstIndex, entry, channel->depth); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + sqeIndex)); + } A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel)); A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel) + 64U); pipe_barrier(PIPE_ALL); @@ -153,6 +167,16 @@ __aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, return A5SdmaEncodeEvent(channelIndex, generation); } +__aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint64_t bytes, + uint32_t channelIndex) +{ + return A5SdmaCopyStridedNbi( + workspaceAddress, destination, source, bytes, 1U, 0U, 0U, channelIndex); +} + __aicore__ inline bool A5SdmaWaitEvent(__gm__ uint8_t* workspaceAddress, uint64_t event, uint32_t expectedChannel) diff --git a/src/include/tilexr_sdma_a5_types.h b/src/include/tilexr_sdma_a5_types.h index 507f082d..59d091be 100644 --- a/src/include/tilexr_sdma_a5_types.h +++ b/src/include/tilexr_sdma_a5_types.h @@ -126,7 +126,14 @@ TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueStateValid(uint32_t tail, uint TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTail(uint32_t tail, uint32_t depth) { - return (tail + 2U) % depth; + return static_cast((static_cast(tail) + 2U) % depth); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTailBy( + uint32_t tail, uint32_t entries, uint32_t depth) +{ + return static_cast( + (static_cast(tail) + static_cast(entries)) % depth); } TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaQueueDistance( @@ -135,13 +142,23 @@ TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaQueueDistance( return (tail + depth - head) % depth; } -TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasCapacity( - uint32_t head, uint32_t tail, uint32_t depth) +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasEntriesCapacity( + uint32_t head, uint32_t tail, uint32_t depth, uint32_t requiredEntries) { if (depth < 3U || head >= depth || tail >= depth) { return false; } - return A5SdmaQueueDistance(head, tail, depth) <= depth - 3U; + if (requiredEntries == 0U || requiredEntries >= depth) { + return false; + } + const uint32_t used = A5SdmaQueueDistance(head, tail, depth); + return requiredEntries <= depth - used - 1U; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasCapacity( + uint32_t head, uint32_t tail, uint32_t depth) +{ + return A5SdmaQueueHasEntriesCapacity(head, tail, depth, 2U); } TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskId(uint32_t taskId) @@ -149,6 +166,37 @@ TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskId(uint32_t taskId) return (taskId + 2U) & 0xFFFFU; } +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskIdBy( + uint32_t taskId, uint32_t entries) +{ + return static_cast( + (static_cast(taskId) + static_cast(entries)) & 0xFFFFULL); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaStridedRangeValid( + uint64_t address, uint64_t bytes, uint32_t copyCount, uint64_t strideBytes) +{ + if (address == 0U || bytes == 0U || copyCount == 0U) { + return false; + } + if (copyCount == 1U) { + return bytes - 1U <= 0xFFFFFFFFFFFFFFFFULL - address; + } + if (strideBytes < bytes) { + return false; + } + const uint64_t lastIndex = static_cast(copyCount - 1U); + if (strideBytes > 0xFFFFFFFFFFFFFFFFULL / lastIndex) { + return false; + } + const uint64_t lastOffset = strideBytes * lastIndex; + if (lastOffset > 0xFFFFFFFFFFFFFFFFULL - address) { + return false; + } + const uint64_t lastAddress = address + lastOffset; + return bytes - 1U <= 0xFFFFFFFFFFFFFFFFULL - lastAddress; +} + TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaNextGeneration(uint32_t generation) { return generation == 0xFFFFFFFFU ? 1U : generation + 1U; diff --git a/tests/sdma/CMakeLists.txt b/tests/sdma/CMakeLists.txt index b0c332bb..d335175b 100644 --- a/tests/sdma/CMakeLists.txt +++ b/tests/sdma/CMakeLists.txt @@ -3,6 +3,7 @@ project(TileXR_SDMA_Tests) set(CMAKE_CXX_STANDARD 14) set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_SKIP_BUILD_RPATH TRUE) option(BUILD_TILEXR_SDMA_DEMO "Build TileXR SDMA data-plane demo with Ascend C kernel" ON) set(TILEXR_SDMA_DEMO_SOC_TYPE "Ascend910B" CACHE STRING "SOC type used for the TileXR SDMA demo kernel") set_property(CACHE TILEXR_SDMA_DEMO_SOC_TYPE PROPERTY STRINGS Ascend910B Ascend950) @@ -56,8 +57,8 @@ add_executable(test_tilexr_sdma_metadata unit/test_tilexr_sdma_metadata.cpp ) target_include_directories(test_tilexr_sdma_metadata PRIVATE - ${TILEXR_ROOT}/install/include ${TILEXR_ROOT}/src/include + ${TILEXR_ROOT}/install/include ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/ ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/ ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ @@ -264,6 +265,7 @@ elseif(BUILD_TILEXR_SDMA_DEMO) -o "${TILEXR_SDMA_DEMO_KERNEL_SO}" DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/demo/tilexr_sdma_demo_kernel.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/demo/tilexr_sdma_benchmark_types.h" "${TILEXR_ROOT}/src/include/tilexr_sdma.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_a5.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_a5_types.h" @@ -300,6 +302,31 @@ elseif(BUILD_TILEXR_SDMA_DEMO) ascend_hal ) list(APPEND INSTALL_TARGETS tilexr_sdma_demo) + + if(TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend950") + add_executable(tilexr_sdma_benchmark + demo/tilexr_sdma_benchmark.cpp + ) + add_dependencies(tilexr_sdma_benchmark tilexr_sdma_demo_kernel) + target_include_directories(tilexr_sdma_benchmark PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/demo + ${TILEXR_ROOT}/install/include + ${TILEXR_ROOT}/src/include + ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/ + ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/ + ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ + ${ASCEND_DRIVER_PATH}/kernel/inc + ) + target_link_directories(tilexr_sdma_benchmark PRIVATE ${CMAKE_CURRENT_BINARY_DIR}) + target_link_libraries(tilexr_sdma_benchmark + "${TILEXR_SDMA_DEMO_KERNEL_SO}" + ${TILEXR_LIB} + ascendcl + runtime + ascend_hal + ) + list(APPEND INSTALL_TARGETS tilexr_sdma_benchmark) + endif() install(FILES "${TILEXR_SDMA_DEMO_KERNEL_SO}" DESTINATION ${CMAKE_INSTALL_PREFIX}/lib) message(STATUS "TileXR SDMA demo enabled with ${BISHENG_EXECUTABLE}, SOC=${TILEXR_SDMA_DEMO_SOC_TYPE}") endif() diff --git a/tests/sdma/demo/tilexr_sdma_benchmark.cpp b/tests/sdma/demo/tilexr_sdma_benchmark.cpp new file mode 100644 index 00000000..9eed7e9d --- /dev/null +++ b/tests/sdma/demo/tilexr_sdma_benchmark.cpp @@ -0,0 +1,362 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "acl/acl.h" +#include "tilexr_api.h" +#include "tilexr_sdma_a5_types.h" +#include "tilexr_sdma_benchmark_types.h" +#include "tilexr_types.h" + +extern "C" void launch_tilexr_sdma_benchmark( + void* stream, + GM_ADDR commArgs, + GM_ADDR dst, + GM_ADDR src, + GM_ADDR samples, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots); + +namespace { +constexpr uint32_t kDefaultBytes = 4096U; +constexpr uint32_t kDefaultWarmup = 20U; +constexpr uint32_t kDefaultIterations = 100U; +constexpr uint32_t kDefaultSamples = 10U; +constexpr uint32_t kAlignmentBytes = 64U; +constexpr uint32_t kMaxSamples = 100U; +constexpr uint32_t kMaxBatchCopies = 32U; +constexpr size_t kMinWorkingSetBytes = 64U * 1024U * 1024U; +constexpr int kDeviceId = 0; +constexpr double kAscend950CyclesPerUs = 1000.0; + +bool CheckAcl(const std::string& label, aclError ret) +{ + if (ret == ACL_SUCCESS) { + return true; + } + std::cerr << "ERROR: " << label << " failed with " << ret << std::endl; + return false; +} + +bool CheckTileXR(const std::string& label, int ret) +{ + if (ret == TileXR::TILEXR_SUCCESS) { + return true; + } + std::cerr << "ERROR: " << label << " failed with " << ret << std::endl; + return false; +} + +bool ParseUint32(const char* text, uint32_t* value) +{ + if (text == nullptr || value == nullptr || text[0] == '\0') { + return false; + } + errno = 0; + char* end = nullptr; + const unsigned long long parsed = std::strtoull(text, &end, 10); + if (errno != 0 || end == text || *end != '\0' || + parsed > std::numeric_limits::max()) { + return false; + } + *value = static_cast(parsed); + return true; +} + +void FillPattern(std::vector& data) +{ + for (size_t index = 0U; index < data.size(); ++index) { + data[index] = static_cast((index * 37U + 11U) & 0xFFU); + } +} + +void Cleanup(uint8_t* src, + uint8_t* dst, + TileXR::test::SdmaBenchmarkSample* samples, + TileXRCommPtr comm, + aclrtStream stream, + bool deviceSet, + bool aclInitialized) +{ + if (src != nullptr) { + (void)aclrtFree(src); + } + if (dst != nullptr) { + (void)aclrtFree(dst); + } + if (samples != nullptr) { + (void)aclrtFree(samples); + } + if (comm != nullptr) { + (void)TileXRCommDestroy(comm); + } + if (stream != nullptr) { + (void)aclrtDestroyStream(stream); + } + if (deviceSet) { + (void)aclrtResetDevice(kDeviceId); + } + if (aclInitialized) { + (void)aclFinalize(); + } +} + +const char* StatusName(uint32_t status) +{ + switch (status) { + case TileXR::test::TILEXR_SDMA_BENCHMARK_OK: + return "ok"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_DISABLED: + return "sdma_disabled"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED: + return "warmup_submit_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED: + return "warmup_wait_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED: + return "submit_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED: + return "wait_failed"; + default: + return "unknown"; + } +} +} // namespace + +int main(int argc, char** argv) +{ + uint32_t bytes = kDefaultBytes; + uint32_t warmupIterations = kDefaultWarmup; + uint32_t measuredIterations = kDefaultIterations; + uint32_t sampleCount = kDefaultSamples; + uint32_t channel = 0U; + uint32_t phaseProfile = 0U; + uint32_t batchCopies = 1U; + if (argc > 8 || + (argc >= 2 && !ParseUint32(argv[1], &bytes)) || + (argc >= 3 && !ParseUint32(argv[2], &warmupIterations)) || + (argc >= 4 && !ParseUint32(argv[3], &measuredIterations)) || + (argc >= 5 && !ParseUint32(argv[4], &sampleCount)) || + (argc >= 6 && !ParseUint32(argv[5], &channel)) || + (argc >= 7 && !ParseUint32(argv[6], &phaseProfile)) || + (argc == 8 && !ParseUint32(argv[7], &batchCopies)) || + bytes == 0U || (bytes % kAlignmentBytes) != 0U || measuredIterations == 0U || + sampleCount == 0U || sampleCount > kMaxSamples || + channel >= TileXR::detail::TILEXR_SDMA_A5_CHANNEL_COUNT || phaseProfile > 1U || + batchCopies == 0U || batchCopies > kMaxBatchCopies || + measuredIterations > std::numeric_limits::max() / batchCopies || + static_cast(bytes) * batchCopies > std::numeric_limits::max()) { + std::cerr << "ERROR: usage: tilexr_sdma_benchmark " + "[aligned-bytes [warmup [iterations [samples [channel " + "[phase-profile [batch-copies]]]]]]]" + << std::endl; + return 1; + } + + (void)setenv("TILEXR_ENABLE_SDMA", "1", 1); + bool aclInitialized = false; + bool deviceSet = false; + TileXRCommPtr comm = nullptr; + aclrtStream stream = nullptr; + uint8_t* src = nullptr; + uint8_t* dst = nullptr; + TileXR::test::SdmaBenchmarkSample* deviceSamples = nullptr; + + if (!CheckAcl("aclInit", aclInit(nullptr))) { + return 1; + } + aclInitialized = true; + if (!CheckAcl("aclrtSetDevice", aclrtSetDevice(kDeviceId))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + deviceSet = true; + const char* socName = aclrtGetSocName(); + if (socName == nullptr || std::strstr(socName, "Ascend950") == nullptr) { + std::cerr << "ERROR: device-cycle conversion is only defined here for Ascend950, got " + << (socName == nullptr ? "" : socName) << std::endl; + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + if (!CheckAcl("aclrtCreateStream", aclrtCreateStream(&stream)) || + !CheckTileXR("TileXRCommInitRankLocal", TileXRCommInitRankLocal(1, 0, &comm))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + bool sdmaAvailable = false; + GM_ADDR sdmaWorkspace = nullptr; + GM_ADDR commArgsDev = nullptr; + if (!CheckTileXR("TileXRSDMAAvailable", TileXRSDMAAvailable(comm, &sdmaAvailable)) || + !CheckTileXR("TileXRGetSDMAWorkspaceDev", TileXRGetSDMAWorkspaceDev(comm, &sdmaWorkspace)) || + !CheckTileXR("TileXRGetCommArgsDev", TileXRGetCommArgsDev(comm, commArgsDev)) || + !sdmaAvailable || sdmaWorkspace == nullptr || commArgsDev == nullptr) { + std::cerr << "ERROR: TileXR direct SDMA is unavailable" << std::endl; + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const size_t transferBytes = static_cast(bytes) * batchCopies; + const size_t workingSetSlots = std::max(1U, kMinWorkingSetBytes / transferBytes); + const size_t allocationBytes = transferBytes * workingSetSlots; + const size_t sampleBytes = static_cast(sampleCount) * + sizeof(TileXR::test::SdmaBenchmarkSample); + if (!CheckAcl("aclrtMalloc src", aclrtMalloc( + reinterpret_cast(&src), allocationBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + !CheckAcl("aclrtMalloc dst", aclrtMalloc( + reinterpret_cast(&dst), allocationBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + !CheckAcl("aclrtMalloc samples", aclrtMalloc( + reinterpret_cast(&deviceSamples), sampleBytes, ACL_MEM_MALLOC_HUGE_FIRST))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + std::vector hostSrc(allocationBytes); + std::vector hostDst(allocationBytes, 0U); + std::vector hostSamples(sampleCount); + FillPattern(hostSrc); + if (!CheckAcl("aclrtMemcpy H2D src", aclrtMemcpy( + src, allocationBytes, hostSrc.data(), hostSrc.size(), ACL_MEMCPY_HOST_TO_DEVICE)) || + !CheckAcl("aclrtMemcpy H2D dst", aclrtMemcpy( + dst, allocationBytes, hostDst.data(), hostDst.size(), ACL_MEMCPY_HOST_TO_DEVICE)) || + !CheckAcl("aclrtMemcpy H2D samples", aclrtMemcpy( + deviceSamples, sampleBytes, hostSamples.data(), sampleBytes, ACL_MEMCPY_HOST_TO_DEVICE))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + launch_tilexr_sdma_benchmark( + stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), + reinterpret_cast(deviceSamples), bytes, channel, + warmupIterations, measuredIterations, sampleCount, phaseProfile, batchCopies, + static_cast(workingSetSlots)); + if (!CheckAcl("aclrtSynchronizeStream", aclrtSynchronizeStream(stream)) || + !CheckAcl("aclrtMemcpy D2H dst", aclrtMemcpy( + hostDst.data(), hostDst.size(), dst, allocationBytes, ACL_MEMCPY_DEVICE_TO_HOST)) || + !CheckAcl("aclrtMemcpy D2H samples", aclrtMemcpy( + hostSamples.data(), sampleBytes, deviceSamples, sampleBytes, ACL_MEMCPY_DEVICE_TO_HOST))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const uint64_t submittedBatches = static_cast(warmupIterations) + + static_cast(measuredIterations) * sampleCount; + const size_t touchedSlots = static_cast(std::min( + static_cast(workingSetSlots), submittedBatches)); + const size_t touchedBytes = touchedSlots * transferBytes; + bool valid = std::equal(hostDst.begin(), hostDst.begin() + touchedBytes, hostSrc.begin()) && + std::all_of(hostDst.begin() + touchedBytes, hostDst.end(), + [](uint8_t value) { return value == 0U; }); + std::vector cyclesPerCopy; + cyclesPerCopy.reserve(sampleCount); + const uint32_t expectedCopies = measuredIterations * batchCopies; + for (uint32_t index = 0U; index < sampleCount; ++index) { + const auto& sample = hostSamples[index]; + if (sample.status != TileXR::test::TILEXR_SDMA_BENCHMARK_OK || + sample.completed != expectedCopies || sample.cycles == 0U) { + std::cerr << "ERROR: sample " << index << " status=" << StatusName(sample.status) + << " completed=" << sample.completed << "/" << expectedCopies + << " cycles=" << sample.cycles << std::endl; + valid = false; + continue; + } + cyclesPerCopy.push_back( + static_cast(sample.cycles) / static_cast(expectedCopies)); + } + if (!valid || cyclesPerCopy.size() != sampleCount) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const double meanCycles = std::accumulate(cyclesPerCopy.begin(), cyclesPerCopy.end(), 0.0) / + static_cast(cyclesPerCopy.size()); + double squaredDifferenceSum = 0.0; + for (double cycles : cyclesPerCopy) { + const double difference = cycles - meanCycles; + squaredDifferenceSum += difference * difference; + } + const double stddevCycles = std::sqrt( + squaredDifferenceSum / static_cast(cyclesPerCopy.size())); + const auto minmaxCycles = std::minmax_element(cyclesPerCopy.begin(), cyclesPerCopy.end()); + const double meanUs = meanCycles / kAscend950CyclesPerUs; + const double stddevUs = stddevCycles / kAscend950CyclesPerUs; + const double minUs = *minmaxCycles.first / kAscend950CyclesPerUs; + const double maxUs = *minmaxCycles.second / kAscend950CyclesPerUs; + const double bandwidthGBps = static_cast(bytes) / meanUs / 1000.0; + const double bandwidthGiBps = static_cast(bytes) / meanUs * 1000000.0 / + (1024.0 * 1024.0 * 1024.0); + double meanSubmitUs = 0.0; + double meanWaitUs = 0.0; + double meanCompletionUs = 0.0; + double meanReleaseUs = 0.0; + if (phaseProfile != 0U) { + for (const auto& sample : hostSamples) { + meanSubmitUs += static_cast(sample.submitCycles); + meanWaitUs += static_cast(sample.waitCycles); + meanCompletionUs += static_cast(sample.completionCycles); + meanReleaseUs += static_cast(sample.releaseCycles); + } + const double phaseDivisor = static_cast(sampleCount) * + static_cast(expectedCopies) * kAscend950CyclesPerUs; + meanSubmitUs /= phaseDivisor; + meanWaitUs /= phaseDivisor; + meanCompletionUs /= phaseDivisor; + meanReleaseUs /= phaseDivisor; + } + + std::cout << std::fixed << std::setprecision(6); + for (uint32_t index = 0U; index < sampleCount; ++index) { + std::cout << "SAMPLE index=" << index + << " total_cycles=" << hostSamples[index].cycles + << " cycles_per_copy=" << cyclesPerCopy[index] + << " latency_us=" << cyclesPerCopy[index] / kAscend950CyclesPerUs << std::endl; + } + std::cout << "RESULT soc=" << socName + << " bytes=" << bytes + << " channel=" << channel + << " warmup=" << warmupIterations + << " batches=" << measuredIterations + << " batch_copies=" << batchCopies + << " copies=" << expectedCopies + << " allocation_bytes=" << allocationBytes + << " working_set_slots=" << workingSetSlots + << " verified_slots=" << touchedSlots + << " samples=" << sampleCount + << " mean_cycles=" << meanCycles + << " mean_us=" << meanUs + << " stddev_us=" << stddevUs + << " min_us=" << minUs + << " max_us=" << maxUs + << " bandwidth_GBps=" << bandwidthGBps + << " bandwidth_GiBps=" << bandwidthGiBps << std::endl; + if (phaseProfile != 0U) { + std::cout << "PHASE submit_us=" << meanSubmitUs + << " wait_us=" << meanWaitUs + << " completion_us=" << meanCompletionUs + << " release_us=" << meanReleaseUs + << " combined_us=" << (meanSubmitUs + meanWaitUs) << std::endl; + } + std::cout << "PASS direct SDMA benchmark data verification" << std::endl; + + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 0; +} diff --git a/tests/sdma/demo/tilexr_sdma_benchmark_types.h b/tests/sdma/demo/tilexr_sdma_benchmark_types.h new file mode 100644 index 00000000..79c07f94 --- /dev/null +++ b/tests/sdma/demo/tilexr_sdma_benchmark_types.h @@ -0,0 +1,38 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_BENCHMARK_TYPES_H +#define TILEXR_SDMA_BENCHMARK_TYPES_H + +#include + +namespace TileXR { +namespace test { + +constexpr uint32_t TILEXR_SDMA_BENCHMARK_OK = 0U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_DISABLED = 1U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED = 2U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED = 3U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED = 4U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WAIT_FAILED = 5U; + +struct alignas(64) SdmaBenchmarkSample { + uint64_t cycles; + uint32_t completed; + uint32_t status; + uint64_t submitCycles; + uint64_t waitCycles; + uint64_t completionCycles; + uint64_t releaseCycles; + uint64_t reserved[2]; +}; + +static_assert(sizeof(SdmaBenchmarkSample) == 64U, + "SDMA benchmark sample must occupy one cache line"); + +} // namespace test +} // namespace TileXR + +#endif // TILEXR_SDMA_BENCHMARK_TYPES_H diff --git a/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp b/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp index 1ec0ac62..15e114d7 100644 --- a/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp +++ b/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp @@ -12,6 +12,20 @@ #include "kernel_operator.h" #include "tilexr_sdma.h" +#include "tilexr_sdma_benchmark_types.h" + +__aicore__ inline uint64_t BenchmarkSubmit( + __gm__ TileXR::CommArgs* args, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint32_t bytes, + uint32_t channel, + uint32_t batchCopies) +{ + return TileXR::SDMACopyStridedNbi( + args, destination, source, static_cast(bytes), batchCopies, + static_cast(bytes), static_cast(bytes), channel); +} extern "C" __global__ __aicore__ void tilexr_sdma_copy_kernel( GM_ADDR commArgsGM, @@ -112,3 +126,162 @@ extern "C" void launch_tilexr_sdma_copy( tilexr_sdma_copy_kernel<<>>( commArgs, dst, src, debug, bytes, firstChannel, iterations); } + +extern "C" __global__ __aicore__ void tilexr_sdma_benchmark_kernel( + GM_ADDR commArgsGM, + GM_ADDR dstGM, + GM_ADDR srcGM, + GM_ADDR samplesGM, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots) +{ + if ASCEND_IS_AIV { + auto args = reinterpret_cast<__gm__ TileXR::CommArgs*>(commArgsGM); + auto dst = reinterpret_cast<__gm__ uint8_t*>(dstGM); + auto src = reinterpret_cast<__gm__ uint8_t*>(srcGM); + auto samples = reinterpret_cast<__gm__ TileXR::test::SdmaBenchmarkSample*>(samplesGM); + const uint64_t batchBytes = static_cast(bytes) * batchCopies; + uint64_t operationIndex = 0U; + + uint32_t status = TileXR::SDMAEnabled(args) && workingSetSlots != 0U + ? TileXR::test::TILEXR_SDMA_BENCHMARK_OK + : TileXR::test::TILEXR_SDMA_BENCHMARK_DISABLED; + for (uint32_t iteration = 0U; + iteration < warmupIterations && status == TileXR::test::TILEXR_SDMA_BENCHMARK_OK; + ++iteration) { + const uint64_t offset = (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED; + } else if (!TileXR::SDMAWait(args, event, channel)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED; + } + } + + pipe_barrier(PIPE_ALL); + for (uint32_t sampleIndex = 0U; sampleIndex < sampleCount; ++sampleIndex) { + uint32_t completed = 0U; + uint64_t startCycle = 0U; + uint64_t endCycle = 0U; + uint64_t submitCycles = 0U; + uint64_t waitCycles = 0U; + uint64_t completionCycles = 0U; + uint64_t releaseCycles = 0U; + if (status == TileXR::test::TILEXR_SDMA_BENCHMARK_OK) { + startCycle = static_cast(AscendC::GetSystemCycle()); + if (phaseProfile != 0U) { + for (uint32_t iteration = 0U; iteration < measuredIterations; ++iteration) { + const uint64_t phaseStart = + static_cast(AscendC::GetSystemCycle()); + const uint64_t offset = + (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + const uint64_t submitEnd = + static_cast(AscendC::GetSystemCycle()); + submitCycles += submitEnd - phaseStart; + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED; + break; + } +#if TILEXR_SDMA_A5_AICORE_COMPILE + uint32_t eventChannel = 0U; + uint32_t eventGeneration = 0U; + if (!TileXR::detail::A5SdmaDecodeEvent( + event, eventChannel, eventGeneration)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + auto workspace = reinterpret_cast< + __gm__ TileXR::detail::A5SdmaWorkspace*>(args->sdmaWorkspacePtr); + auto completion = reinterpret_cast< + __gm__ TileXR::detail::A5SdmaCompletionLine*>( + workspace->channels[eventChannel].completionRecordAddress); + while (TileXR::detail::A5SdmaReadCompletion(completion) != + eventGeneration) { + } + const uint64_t completionEnd = + static_cast(AscendC::GetSystemCycle()); +#else + const uint64_t completionEnd = submitEnd; +#endif + const bool waitOk = TileXR::SDMAWait(args, event, channel); + const uint64_t waitEnd = + static_cast(AscendC::GetSystemCycle()); + waitCycles += waitEnd - submitEnd; + completionCycles += completionEnd - submitEnd; + releaseCycles += waitEnd - completionEnd; + if (!waitOk) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + completed += batchCopies; + } + } else { + for (uint32_t iteration = 0U; iteration < measuredIterations; ++iteration) { + const uint64_t offset = + (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED; + break; + } + if (!TileXR::SDMAWait(args, event, channel)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + completed += batchCopies; + } + } + pipe_barrier(PIPE_ALL); + endCycle = static_cast(AscendC::GetSystemCycle()); + } + samples[sampleIndex].cycles = endCycle - startCycle; + samples[sampleIndex].completed = completed; + samples[sampleIndex].status = status; + samples[sampleIndex].submitCycles = submitCycles; + samples[sampleIndex].waitCycles = waitCycles; + samples[sampleIndex].completionCycles = completionCycles; + samples[sampleIndex].releaseCycles = releaseCycles; + } + +#if TILEXR_SDMA_A5_AICORE_COMPILE + pipe_barrier(PIPE_ALL); + for (uint32_t sampleIndex = 0U; sampleIndex < sampleCount; ++sampleIndex) { + TileXR::detail::A5SdmaCleanCacheLine( + reinterpret_cast<__gm__ uint8_t*>(samples + sampleIndex)); + } + pipe_barrier(PIPE_ALL); + dsb(DSB_DDR); +#endif + } +} + +extern "C" void launch_tilexr_sdma_benchmark( + void* stream, + GM_ADDR commArgs, + GM_ADDR dst, + GM_ADDR src, + GM_ADDR samples, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots) +{ + tilexr_sdma_benchmark_kernel<<<1, nullptr, stream>>>( + commArgs, dst, src, samples, bytes, channel, + warmupIterations, measuredIterations, sampleCount, phaseProfile, + batchCopies, workingSetSlots); +} diff --git a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp index c4b7db1f..dec5ee0d 100644 --- a/tests/sdma/unit/test_tilexr_sdma_metadata.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_metadata.cpp @@ -76,12 +76,20 @@ void TestA5QueueAndTransferHelpers() CHECK_TRUE(!A5SdmaQueueStateValid(3U, 3U)); CHECK_EQ(A5SdmaAdvanceTail(0U, 3U), 2U); CHECK_EQ(A5SdmaAdvanceTail(2U, 3U), 1U); + CHECK_EQ(A5SdmaAdvanceTailBy(7U, 3U, 8U), 2U); + CHECK_EQ(A5SdmaAdvanceTailBy(0xFFFFFFFEU, 3U, 0xFFFFFFFFU), 2U); CHECK_EQ(A5SdmaQueueDistance(2U, 1U, 3U), 2U); CHECK_TRUE(A5SdmaQueueHasCapacity(0U, 0U, 3U)); CHECK_TRUE(A5SdmaQueueHasCapacity(1U, 2U, 8U)); CHECK_TRUE(!A5SdmaQueueHasCapacity(0U, 2U, 3U)); CHECK_TRUE(!A5SdmaQueueHasCapacity(3U, 0U, 3U)); + CHECK_TRUE(A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 7U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 8U)); + CHECK_TRUE(A5SdmaQueueHasEntriesCapacity(2U, 6U, 8U, 3U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(2U, 6U, 8U, 4U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 0U)); CHECK_EQ(A5SdmaAdvanceTaskId(0xFFFFU), 1U); + CHECK_EQ(A5SdmaAdvanceTaskIdBy(0xFFFFU, 17U), 16U); CHECK_TRUE(!A5SdmaTransferLengthValid(0U)); CHECK_TRUE(A5SdmaTransferLengthValid(1U)); CHECK_TRUE(A5SdmaTransferLengthValid(TILEXR_SDMA_A5_MAX_TRANSFER_BYTES)); @@ -89,6 +97,16 @@ void TestA5QueueAndTransferHelpers() CHECK_EQ(A5SdmaNextGeneration(0U), 1U); CHECK_EQ(A5SdmaNextGeneration(1U), 2U); CHECK_EQ(A5SdmaNextGeneration(0xFFFFFFFFU), 1U); + + constexpr uint64_t maxAddress = 0xFFFFFFFFFFFFFFFFULL; + CHECK_TRUE(A5SdmaStridedRangeValid(0x1000U, 64U, 16U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0x1000U, 64U, 16U, 63U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0x1000U, 64U, 0U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0U, 64U, 1U, 0U)); + CHECK_TRUE(A5SdmaStridedRangeValid(maxAddress - 63U, 64U, 1U, 0U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(maxAddress - 62U, 64U, 1U, 0U)); + CHECK_TRUE(A5SdmaStridedRangeValid(maxAddress - 127U, 64U, 2U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(maxAddress - 126U, 64U, 2U, 64U)); } void TestA5EventHelpers() diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index c3a0b03b..afb5fc64 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -127,6 +127,7 @@ void TestBuildSelectsExplicitSoc() "TILEXR_SDMA_AICORE_ARCH \"--cce-aicore-arch=dav-c310-vec\""); CheckNeedle(cmakePath, cmakeText, "Unsupported TILEXR_SDMA_DEMO_SOC_TYPE="); + CheckNeedle(cmakePath, cmakeText, "set(CMAKE_SKIP_BUILD_RPATH TRUE)"); } void TestAscend950UsesOwnedDirectBackend() @@ -150,8 +151,18 @@ void TestAscend950UsesOwnedDirectBackend() const std::string devicePath = "src/include/tilexr_sdma_a5.h"; const auto device = ReadFile(devicePath); + CheckNeedle(devicePath, device, "A5SdmaCopyStridedNbi"); CheckNeedle(devicePath, device, "while (A5SdmaReadCompletion(completion) != generation)"); + const std::string publicPath = "src/include/tilexr_sdma.h"; + const auto publicHeader = ReadFile(publicPath); + CheckNeedle(publicPath, publicHeader, "SDMACopyStridedNbi"); + + const std::string benchmarkKernelPath = "tests/sdma/demo/tilexr_sdma_demo_kernel.cpp"; + const auto benchmarkKernel = ReadFile(benchmarkKernelPath); + CheckNeedle(benchmarkKernelPath, benchmarkKernel, "SDMACopyStridedNbi"); + CheckNoNeedle(benchmarkKernelPath, benchmarkKernel, "BenchmarkBatchCopyNbi"); + const std::string demoPath = "tests/sdma/demo/tilexr_sdma_demo.cpp"; const auto demo = ReadFile(demoPath); CheckNeedle(demoPath, demo, "aclrtSynchronizeStream(stream)"); From ce12aa912e564fe7cc0ca2104f7248d348271ed2 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Thu, 30 Jul 2026 16:00:41 +0800 Subject: [PATCH 5/7] refactor: simplify SDMA initialization --- src/comm/sdma/tilexr_sdma_a5_backend.cpp | 4 ++-- src/comm/sdma/tilexr_sdma_transport.cpp | 5 +++-- src/comm/tilexr_comm.cpp | 15 ++++++++------- tests/sdma/unit/test_tilexr_sdma_source_guard.cpp | 2 ++ 4 files changed, 15 insertions(+), 11 deletions(-) diff --git a/src/comm/sdma/tilexr_sdma_a5_backend.cpp b/src/comm/sdma/tilexr_sdma_a5_backend.cpp index 9873d8f1..aa5320fa 100644 --- a/src/comm/sdma/tilexr_sdma_a5_backend.cpp +++ b/src/comm/sdma/tilexr_sdma_a5_backend.cpp @@ -464,8 +464,8 @@ TileXRA5SDMABackend::~TileXRA5SDMABackend() bool TileXRA5SDMABackend::Init(int32_t deviceId) { - if (!Shutdown()) { - TILEXR_LOG(WARN) << "TileXR A5 SDMA previous resources could not be released"; + if (impl_ != nullptr) { + TILEXR_LOG(ERROR) << "TileXR A5 SDMA backend contains state before initialization"; return false; } std::unique_ptr state(new (std::nothrow) Impl()); diff --git a/src/comm/sdma/tilexr_sdma_transport.cpp b/src/comm/sdma/tilexr_sdma_transport.cpp index 16fdb486..9b770757 100644 --- a/src/comm/sdma/tilexr_sdma_transport.cpp +++ b/src/comm/sdma/tilexr_sdma_transport.cpp @@ -47,9 +47,10 @@ bool TileXRSDMATransport::EnvEnabled() int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) { - if (!Shutdown()) { + if (impl_ != nullptr) { + TILEXR_LOG(ERROR) << "TileXR SDMA transport contains state before initialization"; lastStatus_ = SDMAInitStatus::INIT_FAILED; - return TILEXR_SUCCESS; + return TILEXR_ERROR_INTERNAL; } options_ = options; available_ = false; diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp index d06696c4..ee58ffba 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -191,12 +191,9 @@ int TileXRComm::ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &st int TileXRComm::InitSDMA() { if (sdmaTransport_ != nullptr) { - if (!sdmaTransport_->Shutdown()) { - TILEXR_LOG(WARN) << "TileXR previous SDMA resources are still pending cleanup"; - sdmaInitStatus_ = SDMAInitStatus::INIT_FAILED; - return TILEXR_SUCCESS; - } - sdmaTransport_.reset(); + TILEXR_LOG(ERROR) << "TileXR SDMA transport exists before initialization"; + sdmaInitStatus_ = SDMAInitStatus::INIT_FAILED; + return TILEXR_ERROR_INTERNAL; } sdmaTransport_.reset(new (nothrow) TileXRSDMATransport()); if (sdmaTransport_ == nullptr) { @@ -209,7 +206,11 @@ int TileXRComm::InitSDMA() options.devId = devId_; int ret = sdmaTransport_->Init(options); sdmaInitStatus_ = sdmaTransport_->GetLastStatus(); - if (ret != TILEXR_SUCCESS || !sdmaTransport_->IsAvailable()) { + if (ret != TILEXR_SUCCESS) { + TILEXR_LOG(ERROR) << "TileXR SDMA transport initialization failed"; + return ret; + } + if (!sdmaTransport_->IsAvailable()) { if (sdmaInitStatus_ != SDMAInitStatus::DISABLED_BY_ENV) { TILEXR_LOG(WARN) << "TileXR SDMA init unavailable, status " << static_cast(sdmaInitStatus_); } diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index afb5fc64..fcb926a8 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -174,6 +174,8 @@ void TestSdmaFailureIsScopedToOneCommunicator() const auto comm = ReadFile(commPath); CheckNoNeedle(commPath, comm, "g_sdmaUnavailable"); CheckNeedle(commPath, comm, "sdmaTransport_->Init(options)"); + CheckNeedle(commPath, comm, "TileXR SDMA transport exists before initialization"); + CheckNoNeedle(commPath, comm, "TileXR previous SDMA resources are still pending cleanup"); CheckNeedle(commPath, comm, "bool TileXRComm::PrepareDestroy()"); CheckNeedle(commPath, comm, "if (UpdateCommArgsDev() != TILEXR_SUCCESS)"); From 97c0bed7670fce24213277b7b2fa838c9296e073 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Thu, 30 Jul 2026 17:35:42 +0800 Subject: [PATCH 6/7] fix: skip inactive SDMA teardown upload --- src/comm/tilexr_comm.cpp | 3 +++ tests/sdma/unit/test_tilexr_sdma_source_guard.cpp | 1 + 2 files changed, 4 insertions(+) diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp index ee58ffba..05f6e756 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -244,6 +244,9 @@ int TileXRComm::InitSDMA() bool TileXRComm::ResetSDMAState() { + if (sdmaTransport_ == nullptr) { + return true; + } commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; if (UpdateCommArgsDev() != TILEXR_SUCCESS) { diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp index fcb926a8..39b28fa0 100644 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp @@ -177,6 +177,7 @@ void TestSdmaFailureIsScopedToOneCommunicator() CheckNeedle(commPath, comm, "TileXR SDMA transport exists before initialization"); CheckNoNeedle(commPath, comm, "TileXR previous SDMA resources are still pending cleanup"); CheckNeedle(commPath, comm, "bool TileXRComm::PrepareDestroy()"); + CheckNeedle(commPath, comm, "bool TileXRComm::ResetSDMAState()\n{\n if (sdmaTransport_ == nullptr)"); CheckNeedle(commPath, comm, "if (UpdateCommArgsDev() != TILEXR_SUCCESS)"); const std::string wrapPath = "src/comm/comm_wrap.cpp"; From 6f7662bc016fcf99676cccd0f07425d8ed857e33 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Thu, 30 Jul 2026 19:05:47 +0800 Subject: [PATCH 7/7] test: remove SDMA source guard --- tests/sdma/CMakeLists.txt | 8 - tests/sdma/run_tests.sh | 1 - .../unit/test_tilexr_sdma_source_guard.cpp | 204 ------------------ 3 files changed, 213 deletions(-) delete mode 100644 tests/sdma/unit/test_tilexr_sdma_source_guard.cpp diff --git a/tests/sdma/CMakeLists.txt b/tests/sdma/CMakeLists.txt index d335175b..6efa6f99 100644 --- a/tests/sdma/CMakeLists.txt +++ b/tests/sdma/CMakeLists.txt @@ -121,13 +121,6 @@ target_compile_definitions(test_tilexr_sdma_comm_wiring PRIVATE TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" ) -add_executable(test_tilexr_sdma_source_guard - unit/test_tilexr_sdma_source_guard.cpp -) -target_compile_definitions(test_tilexr_sdma_source_guard PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - add_executable(test_tilexr_sdma_header_compile unit/test_tilexr_sdma_header_compile.cpp ) @@ -161,7 +154,6 @@ set(INSTALL_TARGETS test_tilexr_sdma_transport_disabled test_tilexr_sdma_a5_validation test_tilexr_sdma_comm_wiring - test_tilexr_sdma_source_guard test_tilexr_sdma_header_compile test_tilexr_sdma_disabled_comm ) diff --git a/tests/sdma/run_tests.sh b/tests/sdma/run_tests.sh index 3591e4ba..d179edd7 100755 --- a/tests/sdma/run_tests.sh +++ b/tests/sdma/run_tests.sh @@ -65,7 +65,6 @@ fi "${INSTALL_DIR}/bin/test_tilexr_sdma_transport_disabled" "${INSTALL_DIR}/bin/test_tilexr_sdma_a5_validation" "${INSTALL_DIR}/bin/test_tilexr_sdma_comm_wiring" -"${INSTALL_DIR}/bin/test_tilexr_sdma_source_guard" "${INSTALL_DIR}/bin/test_tilexr_sdma_header_compile" if [ "${HAL_AVAILABLE}" -eq 1 ]; then diff --git a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp b/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp deleted file mode 100644 index 39b28fa0..00000000 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ /dev/null @@ -1,204 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path, bool required = true) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - if (required) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - } - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckNoNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "unexpected dependency in " << path << ": " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void CheckNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -void TestCommSourcesDoNotUseShmem() -{ - const std::vector paths = { - "src/comm/CMakeLists.txt", - "src/comm/tilexr_comm.cpp", - "src/comm/comm_wrap.cpp", - "src/comm/tilexr_comm.h", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::vector forbidden = { - "shmem", - "shmem.h", - "libshmem", - "aclshmem", - "ACLSHMEM", - }; - for (const auto& path : paths) { - const auto text = ReadFile(path); - for (const auto& needle : forbidden) { - CheckNoNeedle(path, text, needle); - } - } -} - -void TestOnlyCompatIncludesSdmaIntrinsics() -{ - const std::vector paths = { - "src/include/tilexr_sdma.h", - "src/include/tilexr_sdma_compat.h", - "src/include/tilexr_sdma_types.h", - "src/include/comm_args.h", - "src/comm/tilexr_comm.cpp", - "src/comm/tilexr_comm.h", - "src/comm/comm_wrap.cpp", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::string ptoPrefix = "pto/npu/comm/async/sdma/"; - const std::string intrinHeader = ptoPrefix + "sdma_async_intrin.hpp"; - const std::string workspaceHeader = ptoPrefix + "sdma_workspace_manager.hpp"; - for (const auto& path : paths) { - const bool required = path != "src/include/tilexr_sdma.h"; - const auto text = ReadFile(path, required); - if (path == "src/include/tilexr_sdma_compat.h") { - CheckNeedle(path, text, intrinHeader); - CheckNoNeedle(path, text, workspaceHeader); - continue; - } - if (path == "src/comm/sdma/tilexr_sdma_transport.cpp") { - CheckNeedle(path, text, workspaceHeader); - CheckNoNeedle(path, text, intrinHeader); - continue; - } - CheckNoNeedle(path, text, ptoPrefix); - } -} - -void TestBuildSelectsExplicitSoc() -{ - const std::string buildPath = "tests/sdma/build.sh"; - const auto buildText = ReadFile(buildPath); - CheckNeedle(buildPath, buildText, - "SDMA_SOC_TYPE=\"${2:-${TILEXR_SDMA_DEMO_SOC_TYPE:-Ascend910B}}\""); - CheckNeedle(buildPath, buildText, - "-DTILEXR_SDMA_DEMO_SOC_TYPE=\"${SDMA_SOC_TYPE}\""); - - const std::string cmakePath = "tests/sdma/CMakeLists.txt"; - const auto cmakeText = ReadFile(cmakePath); - CheckNeedle(cmakePath, cmakeText, - "TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL \"Ascend950\""); - CheckNeedle(cmakePath, cmakeText, - "TILEXR_SDMA_NPU_ARCH \"dav-3510\""); - CheckNeedle(cmakePath, cmakeText, - "TILEXR_SDMA_AICORE_ARCH \"--cce-aicore-arch=dav-c310-vec\""); - CheckNeedle(cmakePath, cmakeText, - "Unsupported TILEXR_SDMA_DEMO_SOC_TYPE="); - CheckNeedle(cmakePath, cmakeText, "set(CMAKE_SKIP_BUILD_RPATH TRUE)"); -} - -void TestAscend950UsesOwnedDirectBackend() -{ - const std::string transportPath = "src/comm/sdma/tilexr_sdma_transport.cpp"; - const auto transport = ReadFile(transportPath); - CheckNeedle(transportPath, transport, "TileXRA5SDMABackend"); - CheckNeedle(transportPath, transport, "ClassifySDMABackend(socName)"); - - const std::string backendPath = "src/comm/sdma/tilexr_sdma_a5_backend.cpp"; - const auto backend = ReadFile(backendPath); - CheckNeedle(backendPath, backend, "aclnnShmemSdmaStarsQuery"); - CheckNeedle(backendPath, backend, "RES_ADDR_TYPE_STARS_RTSQ"); - CheckNeedle(backendPath, backend, "TILEXR_SDMA_A5_CHANNEL_COUNT"); - CheckNeedle(backendPath, backend, "kExpectedAicpuQueryFailure"); - CheckNeedle(backendPath, backend, - "LoadSymbol(RTLD_DEFAULT, \"halResAddrMap\", mapResource)"); - CheckNeedle(backendPath, backend, "pendingQueries"); - CheckNoNeedle(backendPath, backend, "halResAddrMap("); - CheckNoNeedle(backendPath, backend, "halResAddrUnmap("); - - const std::string devicePath = "src/include/tilexr_sdma_a5.h"; - const auto device = ReadFile(devicePath); - CheckNeedle(devicePath, device, "A5SdmaCopyStridedNbi"); - CheckNeedle(devicePath, device, "while (A5SdmaReadCompletion(completion) != generation)"); - - const std::string publicPath = "src/include/tilexr_sdma.h"; - const auto publicHeader = ReadFile(publicPath); - CheckNeedle(publicPath, publicHeader, "SDMACopyStridedNbi"); - - const std::string benchmarkKernelPath = "tests/sdma/demo/tilexr_sdma_demo_kernel.cpp"; - const auto benchmarkKernel = ReadFile(benchmarkKernelPath); - CheckNeedle(benchmarkKernelPath, benchmarkKernel, "SDMACopyStridedNbi"); - CheckNoNeedle(benchmarkKernelPath, benchmarkKernel, "BenchmarkBatchCopyNbi"); - - const std::string demoPath = "tests/sdma/demo/tilexr_sdma_demo.cpp"; - const auto demo = ReadFile(demoPath); - CheckNeedle(demoPath, demo, "aclrtSynchronizeStream(stream)"); -} - -void TestSdmaFailureIsScopedToOneCommunicator() -{ - const std::string commPath = "src/comm/tilexr_comm.cpp"; - const auto comm = ReadFile(commPath); - CheckNoNeedle(commPath, comm, "g_sdmaUnavailable"); - CheckNeedle(commPath, comm, "sdmaTransport_->Init(options)"); - CheckNeedle(commPath, comm, "TileXR SDMA transport exists before initialization"); - CheckNoNeedle(commPath, comm, "TileXR previous SDMA resources are still pending cleanup"); - CheckNeedle(commPath, comm, "bool TileXRComm::PrepareDestroy()"); - CheckNeedle(commPath, comm, "bool TileXRComm::ResetSDMAState()\n{\n if (sdmaTransport_ == nullptr)"); - CheckNeedle(commPath, comm, "if (UpdateCommArgsDev() != TILEXR_SUCCESS)"); - - const std::string wrapPath = "src/comm/comm_wrap.cpp"; - const auto wrap = ReadFile(wrapPath); - CheckNeedle(wrapPath, wrap, "if (!c->PrepareDestroy())"); - CheckNeedle(wrapPath, wrap, "return TILEXR_ERROR_INTERNAL;"); -} - -} // namespace - -int main() -{ - TestCommSourcesDoNotUseShmem(); - TestOnlyCompatIncludesSdmaIntrinsics(); - TestBuildSelectsExplicitSoc(); - TestAscend950UsesOwnedDirectBackend(); - TestSdmaFailureIsScopedToOneCommunicator(); - if (g_failures != 0) { - std::cerr << g_failures << " SDMA source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR SDMA source guard checks passed" << std::endl; - return 0; -}