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 5c64771e..15578b31 100644 --- a/src/comm/CMakeLists.txt +++ b/src/comm/CMakeLists.txt @@ -170,6 +170,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}) @@ -237,6 +239,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 cccb7c5d..277468b9 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -52,8 +52,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; constexpr const char* TILEXR_ENABLE_CCU_BACKEND_ENV = "TILEXR_ENABLE_CCU_BACKEND"; @@ -258,15 +256,14 @@ int TileXRComm::EnableCcuBackendForTest() 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"; @@ -281,10 +278,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; @@ -297,7 +294,9 @@ int TileXRComm::InitSDMA() commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; sdmaWorkspaceDev_ = nullptr; - sdmaTransport_.reset(); + if (sdmaTransport_->Shutdown()) { + sdmaTransport_.reset(); + } return TILEXR_SUCCESS; } @@ -308,16 +307,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 @@ -973,7 +980,7 @@ TileXRComm::~TileXRComm() ccuBackend_->Shutdown(); ccuBackend_.reset(); } - 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 48dda0c3..532fd196 100644 --- a/src/comm/tilexr_comm.h +++ b/src/comm/tilexr_comm.h @@ -58,6 +58,7 @@ class TileXRComm { bool IsSDMAAvailable() const; GM_ADDR GetSDMAWorkspacePtr() const; SDMAInitStatus GetSDMAInitStatus() const; + bool PrepareDestroy(); std::string PrintDFX(); friend class Lccl; friend class Lcoc; @@ -86,7 +87,7 @@ class TileXRComm { int InitCcuBackendIfEnabled(); 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;