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..7d3b9069 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,219 @@ 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); + +uint64_t batchEvent = TileXR::SDMACopyStridedNbi( + args, dst, src, bytes, copyCount, dstStrideBytes, srcStrideBytes, channel); +bool batchOk = TileXR::SDMAWait(args, batchEvent, 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 +`SDMACopyStridedNbi` submits equal-sized, non-overlapping strided slices. On A5, +`copyCount > 1` writes one data SQE per slice, appends one completion SQE, and +rings one doorbell. The returned event completes the whole batch. A2/A3 retains +the existing PTO single-copy path; multi-copy strided submission returns event +0 there. `copyCount == 1` delegates to `SDMACopyNbi` on every backend. -`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`. +## A2/A3 PTO Backend -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`. +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`. -## CANN Compatibility +## A5 Direct Backend -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`. +On Ascend950/A5, TileXR owns the Host initializer and AIV submission path. It +does not ship a TileXR AICPU binary or operator package. -Runtime must load `libascend_hal.so` from the driver path, typically: +During communicator initialization, the Host backend: -```text -/usr/local/Ascend/driver/lib64/driver/libascend_hal.so -``` +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`. -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`. +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, +waits for the matching completion record, and releases the channel only after +completion. -## Build And Test +For a strided batch, the AIV backend reserves `copyCount + 1` SQ entries, +advances consecutive task IDs, and reuses the same completion/wait path. Queue +capacity, count, length, stride, and 64-bit address arithmetic are validated +before the channel is claimed. Multi-copy source and destination strides must +each be at least the copy length. -Build tests against a selected CANN install: +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. -```bash -bash tests/sdma/build.sh /path/to/cann -bash tests/sdma/run_tests.sh /path/to/cann -``` +## Build -Run data-plane demo: +Build the core and SDMA tests against CANN 9.1.0: ```bash -bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 0 64 4096 1048576 +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 ``` -Expected demo success line: +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. -```text -PASS TileXR SDMA copied bytes correctly -``` +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. -## Acceptance +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. -For release validation, run the unit tests and demo against both CANN versions: +## Hardware Validation -```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 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 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 +```bash +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 ``` -## 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: +Every run must report a nonzero event, a successful wait, the expected +generation, and a full byte-for-byte destination comparison. A skipped or +unavailable backend is not an A5 data-plane pass. + +### Batch performance and selection + +The A5-only benchmark accepts: ```text -PASS TileXR SDMA copied 64 bytes correctly -PASS TileXR SDMA copied 4096 bytes correctly -PASS TileXR SDMA copied 1048576 bytes correctly +tilexr_sdma_benchmark \ + [aligned-bytes [warmup-batches [measured-batches [samples \ + [channel [phase-profile [batch-copies]]]]]]] ``` -Remaining validation: +The following device-cycle measurements use a rotating 64 MiB working set and +the same `ACL_MEM_MALLOC_HUGE_FIRST` backing size for single and batch runs. +Each result includes SQE construction, cache maintenance, doorbell, device +completion, and channel release, while excluding Host launch and allocation. + +| Bytes per copy | Single completion | Batch 16 amortized | Speedup | +| ---: | ---: | ---: | ---: | +| 4 KiB | 3.351 us | 0.954 us | 3.51x | +| 8 KiB | 3.372 us | 0.975 us | 3.46x | +| 16 KiB | 3.415 us | 0.972 us | 3.51x | +| 32 KiB | 4.133 us | 1.031 us | 4.01x | +| 64 KiB | 3.453 us | 1.012 us | 3.41x | +| 1 MiB | 4.314 us | 1.962 us | 2.20x | +| 4 MiB | 8.752 us | 6.313 us | 1.39x | + +Use batching for naturally available 4-64 KiB work. From 64 KiB through +1 MiB it remains throughput-beneficial, but do not wait to fill a batch. Above +1 MiB, batch only already-available parallel copies; send latency-sensitive +isolated copies immediately. Batch values are amortized: all copies share one +event that becomes visible only after the complete batch. + +## Validation Record + +Production A5 acceptance was exercised on: + +- 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; +- formal `SDMACopyStridedNbi` batches: counts 1, 2, 4, 8, 16, and 32 at + 4 KiB on channel 0, with 20,000 warmup batches and repeated SQ-tail/task-ID + wraparound; +- channel 47: batch 16 at 4 KiB completed at 0.907 us/copy amortized; +- representative batch 16 results through the production API: 4 KiB at + 0.943 us/copy, 64 KiB at 1.044 us/copy, 1 MiB at 1.909 us/copy, and 4 MiB + at 6.315 us/copy; +- validation: every successful run compared the full source and destination. + +The A2/A3 PTO path was previously exercised on Ascend 910B3 with CANN 9.0.0 and +9.1.0 and driver `25.5.0`. That driver value records the tested machine; it is +not TileXR's minimum supported driver. + +## Runtime Dependency Checks + +Runtime must resolve the real driver HAL, typically: -- Broader parameter/performance matrix below; current demo covers one device, - one stream, channel group `0`, and sizes 64 B, 4 KiB, and 1 MiB. +```text +/usr/local/Ascend/driver/lib64/driver/libascend_hal.so +``` -## 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..7786f2d9 --- /dev/null +++ b/docs/plans/2026-07-28-a5-direct-sdma-implementation.md @@ -0,0 +1,418 @@ +# A5 Direct SDMA Implementation Plan + +Date: 2026-07-28 +Status: Approved for execution +Design: `docs/specs/2026-07-28-a5-direct-sdma-design.md` +Updated: 2026-07-29 (Task 9 adds approved strided batch submission) + +## Goal + +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 +single-copy two-SQE or strided-batch `N + 1` sequence directly. + +The project-wide supported baseline becomes CANN 9.1.0 with driver +`25.1.rc1` or later. The validated hardware baseline is +`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`, `SDMACopyStridedNbi`, and `SDMAWait` to A5 direct + submission for dav-3510 kernels while keeping PTO single-copy behavior on + A2/A3. +- Extend SDMA unit, integration, build, demo, dependency, and hardware checks. +- Change requirements and CI provisioning from driver `25.5.0+` to + `25.1.rc1+`, including Huawei RC/build-suffix comparison. +- 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, workspace ABI, or event-zero semantics. +- Do not replace PTO SDMA on A2/A3, make SDMA mandatory for communicator init, + add more than one outstanding event per channel, or claim Host/simulator + checks prove the A5 data plane. +- 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 original backend +acceptance gate. Task 9 depends on Tasks 1, 4, 5, and 8 and is the final batch +extension gate; PoC removal is not considered complete until its evidence is +covered there. + +## Task 1: Define The A5 ABI And Test Seams + +**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, update the completion payload, build data and 64-byte completion-copy SQEs with +type 11, `wrCqe=0`, credit 254, flush all touched cache lines, issue ordering, +and ring RTSQ only after both SQEs are ready. `SDMAWait(0)` succeeds; other +waits reject bad channel/generation/stale events, poll matching completion, and +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. + +## Task 9: Promote Strided Batch Submission + +**Objective and role:** Move the hardware-proven batch prototype into the +installed TileXR device API so multiple independent same-sized copies can share +one A5 doorbell and one completion event. + +**Background and prerequisites:** Representative A5 measurements show +3.4-4.0x amortized improvement for 4-64 KiB, 2.2x at 1 MiB, and 1.39x at +4 MiB. The accepted design adds `SDMACopyStridedNbi` without changing Host or +workspace ABI. One batch uses `N` data SQEs plus one completion-copy SQE and +retains one outstanding event per channel. + +**Modification scope:** Update `src/include/tilexr_sdma_a5_types.h` with pure +multi-entry queue helpers, implement the A5 strided batch in +`src/include/tilexr_sdma_a5.h`, expose the architecture-routed wrapper in +`src/include/tilexr_sdma.h`, and replace the private implementation in +`tests/sdma/demo/tilexr_sdma_demo_kernel.cpp`. Extend metadata, header/source, +benchmark, and transport documentation tests as appropriate. + +**Constraints and non-goals:** Preserve C++14 and CANN 9.1, the A5 workspace +layout/version, Host APIs, event format, existing single-copy behavior, A2/A3 +PTO behavior, and one outstanding event per channel. Validate all inputs before +claiming a channel. Require non-overlapping strides for multi-copy batches, +reject address arithmetic overflow and insufficient queue capacity, leave one +SQ entry unused, and ring exactly one doorbell. Do not add descriptor arrays, +begin/append/commit state, CQE polling, broadcast/overlap semantics, or a +multi-copy PTO implementation. + +**Acceptance and verification:** Unit-test arbitrary-entry capacity, tail and +task-ID wrap, invalid count/stride/overflow inputs, and unchanged two-entry +single-copy helpers. Compile the installed header and benchmark kernel for +dav-3510. On A5 run counts 1, 2, 4, 8, 16, and 32 across queue/task-ID wraps, +verify every destination slice, repeat init/use/destroy, and rerun 4 KiB, +64 KiB, 1 MiB, and 4 MiB representative timings with equal 64 MiB backing and +a rotating working set. Run the eight focused SDMA test executables and inspect +the final dependency/RPATH state. + +**Artifacts and interfaces:** Installed `SDMACopyStridedNbi`, reusable A5 +multi-entry queue helpers, benchmark coverage using only the production API, +and recorded hardware semantics/performance evidence. + +## Key Risks And Controls + +- **Expected AICPU failure corrupts the active context:** every query runs in a + 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 `N + 1` slots while leaving + one unused, permit one outstanding event per channel, atomically claim it, + and assign distinct channels to concurrent blocks. +- **Cache/order mismatch:** flush the payload and every written SQE; issue the + validated barrier sequence before the RTSQ tail write. +- **ABI drift between Host and AIV:** one installed fixed-width header and + static assertions are compiled in both modes. +- **Optional runtime becomes a hard dependency:** dynamically resolve the + 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 nine tasks pass, A2/A3 PTO selection is +preserved, A5 device 5 passes the full hardware matrix with 48 live channels, +driver `25.1.rc1.b188` is accepted by the shared comparator, dependency checks +are clean, and the repository contains neither the PoC nor a TileXR OPP path. diff --git a/docs/specs/2026-07-28-a5-direct-sdma-design.md b/docs/specs/2026-07-28-a5-direct-sdma-design.md new file mode 100644 index 00000000..907c3fe0 --- /dev/null +++ b/docs/specs/2026-07-28-a5-direct-sdma-design.md @@ -0,0 +1,306 @@ +# A5 Direct SDMA Design + +Date: 2026-07-28 +Status: Approved on 2026-07-28 +Updated: 2026-07-29 (strided batch submission approved for implementation) + +## 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 existing API behavior; +- add an A5 strided batch device API without changing Host or workspace ABI; +- remain best-effort when any A5 capability is unavailable; +- ship no TileXR OPP package and require no custom OPP environment variable; +- keep `reference/` comparison-only. + +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. + +The installed device header additionally exposes: + +```cpp +uint64_t SDMACopyStridedNbi( + const __gm__ CommArgs* args, + __gm__ uint8_t* dst, + __gm__ uint8_t* src, + uint64_t bytes, + uint32_t copyCount, + uint64_t dstStrideBytes, + uint64_t srcStrideBytes, + uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP); +``` + +For copy `i`, the source and destination are `src + i * srcStrideBytes` and +`dst + i * dstStrideBytes`. `copyCount == 1` delegates to `SDMACopyNbi` and +preserves its behavior. On A5, `copyCount > 1` submits one ordered batch. On +A2/A3, multi-copy batch submission is unavailable and returns event 0; the +existing PTO single-copy path is unchanged. + +For `copyCount > 1`, both strides must be at least `bytes`, every individual +length must fit the A5 SQE, address arithmetic must not overflow, and the +channel must have room for `copyCount + 1` SQEs. These conservative rules keep +all source and destination slices distinct. The API does not promise broadcast +or overlapping-copy semantics. + +Host selects the backend by runtime SoC: + +- A2/A3 continues to use PTO `SdmaWorkspaceManager` and PTO device intrinsics; +- 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. Write the generation into the channel's completion payload. The completion + record is not cleared because generations are monotonic and the completion + SQE overwrites the full record before it can match the new generation. +4. Build two consecutive 64-byte A5 SQEs at the current tail: one data copy and + one 64-byte copy from the payload to the completion record. +5. Set SQE type 11, `wrCqe=0`, kernel credit 254, and the validated A5 address, + substream, length, and task fields. +6. Clean the written payload 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 +generation through `ReadGmByPassDCache`, then releases the outstanding state. +A zero event remains an immediate successful no-op. + +### Strided batch submission + +`SDMACopyStridedNbi` uses the same channel claim, event encoding, completion +record, and wait path. After all validation succeeds, it: + +1. reserves `copyCount + 1` entries while leaving one SQ slot unused; +2. writes `copyCount` ordered data SQEs with consecutive task IDs; +3. appends one 64-byte completion-copy SQE; +4. cleans the payload and every written SQE cache line; +5. advances the tail and task ID by `copyCount + 1` modulo their hardware + widths; +6. rings one doorbell and returns one event for the complete batch. + +No validation failure after the channel claim is allowed. All pointer, stride, +length, count, queue-capacity, and workspace checks therefore occur before the +atomic claim. `SDMAWait` releases the channel only after the final completion +SQE, so every data SQE in the batch is complete when the event becomes visible. + +The returned event represents the whole batch. Reported per-copy latency is +total submit-through-wait time divided by `copyCount`; callers cannot observe +or wait for an individual copy through this event. + +## Batch Performance Evidence + +The batch decision was validated on `Ascend950PR_9589`, CANN 9.1.0, driver +`25.1.rc1.b188`, physical device 5. The benchmark used AIV +`GetSystemCycle()` at 1000 cycles/us, a rotating 64 MiB source/destination +working set with the same allocation policy for single and batch runs, device +warmup before sampling, ten samples, and full working-set byte comparison. +Host launch, initialization, allocation, H2D/D2H, and stream synchronization +were outside the timed region. SQE construction, cache maintenance, doorbell, +device completion, and channel release were included. + +| Bytes per copy | Single completion | Batch 16 amortized | Speedup | +| ---: | ---: | ---: | ---: | +| 4 KiB | 3.351 us | 0.954 us | 3.51x | +| 8 KiB | 3.372 us | 0.975 us | 3.46x | +| 16 KiB | 3.415 us | 0.972 us | 3.51x | +| 32 KiB | 4.133 us | 1.031 us | 4.01x | +| 64 KiB | 3.453 us | 1.012 us | 3.41x | +| 1 MiB | 4.314 us | 1.962 us | 2.20x | +| 4 MiB | 8.752 us | 6.313 us | 1.39x | + +For 4-64 KiB, callers should batch naturally available copies because fixed +submission and completion costs dominate. From 64 KiB through 1 MiB, batching +still materially improves throughput but callers must not delay an otherwise +ready copy merely to fill a batch. Above 1 MiB, batching is optional: use it +for already-available parallel work, while latency-sensitive isolated copies +should remain single submissions. For example, 16 copies at an amortized +0.954 us still expose one batch completion after roughly 15.3 us. + +## Failure and Cleanup + +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 every SQE +in the submission is 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; +- strided batches with counts 1, 2, 4, 8, 16, and 32, including SQ/tail/task-ID + wraparound and full comparison of every destination slice; +- representative 4 KiB through 4 MiB batch performance using equal backing + allocation and a rotating working set; +- a concurrent multi-block run using distinct channels, followed by full byte + comparisons and matching completion generations; +- repeated init/use/shutdown loops with a healthy communicator context; +- real driver `libascend_hal.so` resolution and no CANN `devlib` RPATH/RUNPATH. + +Host/source tests and simulator runs do not prove the A5 data plane. Hardware +claims must name the exercised SoC, driver, device, channel count, sizes, and +concurrency level. + +## Repository Cleanup + +The implementation change replaces the PoC-only material with this design and +the production backend. It removes: + +- `tests/sdma/a5_aicpu_probe/`; +- the A5 PoC and obsolete adaptation plans under `docs/plans/`; +- the temporary Ascend950 preflight that disables the PTO path; +- documentation that describes the production A5 path as pending or requires + driver 25.5.0. + +The existing SDMA demo becomes the single data-plane acceptance entry point for +both A2/A3 and A5. diff --git a/scripts/README.md b/scripts/README.md index c37f0f57..fe288997 100644 --- a/scripts/README.md +++ b/scripts/README.md @@ -75,6 +75,21 @@ bash scripts/download_open_source_deps.sh --check ## Testing +### SDMA + +The production SDMA build, unit-test, and hardware-demo entry points live under +`tests/sdma/`: + +```bash +bash tests/sdma/build.sh /path/to/cann Ascend950 +bash tests/sdma/run_tests.sh /path/to/cann +bash tests/sdma/demo/run_tilexr_sdma_demo.sh /path/to/cann 5 64 4096 1048576 +``` + +Omit `Ascend950` to compile the Ascend910B PTO demo. See +[`docs/SDMA_TRANSPORT.md`](../docs/SDMA_TRANSPORT.md) for channel, concurrency, +repeat, and dependency checks. + ### Pull-request CI The CI operator runbook is [docs/CI.md](../docs/CI.md). Primary entrypoints are: diff --git a/scripts/ci/control/build_blue.sh b/scripts/ci/control/build_blue.sh index 9880e917..59f17d5e 100755 --- a/scripts/ci/control/build_blue.sh +++ b/scripts/ci/control/build_blue.sh @@ -246,6 +246,8 @@ rm -rf \ "${SOURCE_DIR}/tests/udma/install" \ "${SOURCE_DIR}/tests/sdma/build" \ "${SOURCE_DIR}/tests/sdma/install" \ + "${SOURCE_DIR}/tests/sdma/build-a5" \ + "${SOURCE_DIR}/tests/sdma/install-a5" \ "${SOURCE_DIR}/tests/ep/build" \ "${SOURCE_DIR}/tests/ep/install" \ "${SOURCE_DIR}/tests/memory/build" \ @@ -278,12 +280,22 @@ run_case udma-source-guard "${SOURCE_DIR}/tests/udma/install/bin/test_tilexr_udm run_logged_step "sdma-build" bash \ "${SOURCE_DIR}/tests/sdma/build.sh" "${ASCEND_HOME_PATH}" run_case sdma-metadata "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_metadata" +run_case sdma-a5-validation "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_a5_validation" run_case sdma-api-invalid "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_api_invalid" run_case sdma-transport-disabled "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_transport_disabled" run_case sdma-comm-wiring "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_comm_wiring" run_case sdma-source-guard "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_source_guard" run_case sdma-header-compile "${SOURCE_DIR}/tests/sdma/install/bin/test_tilexr_sdma_header_compile" +run_logged_step "sdma-a5-configure" cmake \ + -S "${SOURCE_DIR}/tests/sdma" -B "${SOURCE_DIR}/tests/sdma/build-a5" \ + -DCMAKE_INSTALL_PREFIX="${SOURCE_DIR}/tests/sdma/install-a5" \ + -DBUILD_TILEXR_SDMA_DEMO=ON \ + -DTILEXR_SDMA_DEMO_SOC_TYPE=Ascend950 +run_logged_step "sdma-a5-build" cmake \ + --build "${SOURCE_DIR}/tests/sdma/build-a5" \ + --target tilexr_sdma_demo_kernel -j"${BUILD_JOBS}" + run_logged_step "ep-build" bash "${SOURCE_DIR}/tests/ep/build.sh" full run_case ep-layout "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_layout" run_case ep-api-sources "${SOURCE_DIR}/tests/ep/install/bin/test_tilexr_ep_api_sources" @@ -319,7 +331,7 @@ for required_header in \ tilexr_api.h tilexr_types.h comm_args.h tilexr_sync.h \ tilexr_data_as_flag.h tilexr_perf_trace.h \ tilexr_udma.h tilexr_udma_reg.h tilexr_udma_types.h \ - tilexr_sdma_config.h tilexr_sdma_types.h tilexr_sdma.h tilexr_sdma_compat.h \ + tilexr_sdma_config.h tilexr_sdma_types.h tilexr_sdma_a5_types.h tilexr_sdma_a5.h tilexr_sdma.h tilexr_sdma_compat.h \ tilexr_ep.h tilexr_collectives.h tilexr_collectives_perf.h do require_regular_file "${SOURCE_DIR}/install/include/${required_header}" diff --git a/scripts/ci/provision/cann.sh b/scripts/ci/provision/cann.sh index af3947f4..7c453f22 100755 --- a/scripts/ci/provision/cann.sh +++ b/scripts/ci/provision/cann.sh @@ -143,8 +143,8 @@ check_driver_version() { local driver_version driver_version="$(awk -F= '$1 == "Version" {print $2; exit}' \ /usr/local/Ascend/driver/version.info 2>/dev/null)" || driver_version="" - if ! version_at_least "${driver_version}" 25.5.0; then - echo "ERROR: driver >= 25.5.0 is required, found ${driver_version:-unknown}" >&2 + if ! version_at_least "${driver_version}" 25.1.rc1; then + echo "ERROR: driver >= 25.1.rc1 is required, found ${driver_version:-unknown}" >&2 return 1 fi } diff --git a/scripts/ci/provision/common.sh b/scripts/ci/provision/common.sh index c53b2466..b0e49cc0 100755 --- a/scripts/ci/provision/common.sh +++ b/scripts/ci/provision/common.sh @@ -46,22 +46,45 @@ require_root() { version_at_least() { local actual="$1" local minimum="$2" - local actual_major actual_minor actual_patch - local minimum_major minimum_minor minimum_patch - - [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)$ ]] || return 1 - actual_major=$((10#${BASH_REMATCH[1]})) - actual_minor=$((10#${BASH_REMATCH[2]})) - actual_patch=$((10#${BASH_REMATCH[3]})) - [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)$ ]] || return 1 - minimum_major=$((10#${BASH_REMATCH[1]})) - minimum_minor=$((10#${BASH_REMATCH[2]})) - minimum_patch=$((10#${BASH_REMATCH[3]})) + local actual_major actual_minor actual_release actual_stage + local minimum_major minimum_minor minimum_release minimum_stage + + if [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.[rR][cC]([0-9]+)(\.b([0-9]+))?$ ]]; then + [[ "${BASH_REMATCH[3]}" != 0 ]] || return 1 + actual_major=$((10#${BASH_REMATCH[1]})) + actual_minor=$((10#${BASH_REMATCH[2]})) + actual_release=$((10#${BASH_REMATCH[3]})) + actual_stage=0 + elif [[ "${actual}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)(\.b([0-9]+))?$ ]]; then + actual_major=$((10#${BASH_REMATCH[1]})) + actual_minor=$((10#${BASH_REMATCH[2]})) + actual_release=$((10#${BASH_REMATCH[3]})) + actual_stage=1 + else + return 1 + fi + + if [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.[rR][cC]([0-9]+)(\.b([0-9]+))?$ ]]; then + [[ "${BASH_REMATCH[3]}" != 0 ]] || return 1 + minimum_major=$((10#${BASH_REMATCH[1]})) + minimum_minor=$((10#${BASH_REMATCH[2]})) + minimum_release=$((10#${BASH_REMATCH[3]})) + minimum_stage=0 + elif [[ "${minimum}" =~ ^([0-9]+)\.([0-9]+)\.([0-9]+)(\.b([0-9]+))?$ ]]; then + minimum_major=$((10#${BASH_REMATCH[1]})) + minimum_minor=$((10#${BASH_REMATCH[2]})) + minimum_release=$((10#${BASH_REMATCH[3]})) + minimum_stage=1 + else + return 1 + fi (( actual_major > minimum_major || (actual_major == minimum_major && actual_minor > minimum_minor) || (actual_major == minimum_major && actual_minor == minimum_minor && - actual_patch >= minimum_patch) )) + actual_stage > minimum_stage) || + (actual_major == minimum_major && actual_minor == minimum_minor && + actual_stage == minimum_stage && actual_release >= minimum_release) )) } ci_primary_group_has_non_root_gid() { diff --git a/scripts/ci/provision/verify.sh b/scripts/ci/provision/verify.sh index 6bc504d2..712f2f0d 100755 --- a/scripts/ci/provision/verify.sh +++ b/scripts/ci/provision/verify.sh @@ -33,8 +33,8 @@ check_driver_version() { local driver_version driver_version="$(awk -F= '$1 == "Version" {print $2; exit}' \ /usr/local/Ascend/driver/version.info 2>/dev/null)" || driver_version="" - if ! version_at_least "${driver_version}" 25.5.0; then - echo "ERROR: driver >= 25.5.0 is required, found ${driver_version:-unknown}" >&2 + if ! version_at_least "${driver_version}" 25.1.rc1; then + echo "ERROR: driver >= 25.1.rc1 is required, found ${driver_version:-unknown}" >&2 return 1 fi } diff --git a/src/comm/CMakeLists.txt b/src/comm/CMakeLists.txt index e0fd5253..91e8e0ab 100644 --- a/src/comm/CMakeLists.txt +++ b/src/comm/CMakeLists.txt @@ -118,6 +118,9 @@ set(TILEXR_SOURCE_FILE tilexr_comm.cpp udma/tilexr_udma_transport.cpp sdma/tilexr_sdma_transport.h sdma/tilexr_sdma_transport.cpp + sdma/tilexr_sdma_a5_cleanup.h + sdma/tilexr_sdma_a5_backend.h + sdma/tilexr_sdma_a5_backend.cpp ) add_library(tile-comm SHARED ${TILEXR_SOURCE_FILE}) @@ -178,6 +181,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..aa5320fa --- /dev/null +++ b/src/comm/sdma/tilexr_sdma_a5_backend.cpp @@ -0,0 +1,766 @@ +/* + * 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 + +#include "acl/acl.h" +#include "aclnn/aclnn_base.h" +#include "driver/ascend_hal.h" +#include "sdma/tilexr_sdma_a5_cleanup.h" +#include "tilexr_log.h" + +#ifndef ACL_STREAM_DEVICE_USE_ONLY +#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 HalResAddrMapFn = drvError_t (*)(unsigned int, res_addr_info*, + unsigned long*, unsigned int*); +using HalResAddrUnmapFn = drvError_t (*)(unsigned int, res_addr_info*); +using AclCreateTensorFn = aclTensor* (*)(const int64_t*, uint64_t, aclDataType, + const int64_t*, int64_t, aclFormat, + const int64_t*, uint64_t, void*); +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; + } + + if (!LoadSymbol(RTLD_DEFAULT, "halResAddrMap", mapResource) || + !LoadSymbol(RTLD_DEFAULT, "halResAddrUnmap", unmapResource)) { + Close(); + return false; + } + return true; + } + + void Close() + { + getPhysicalDevice = nullptr; + getSqId = nullptr; + getCqId = nullptr; + getDeviceInfo = nullptr; + createTensor = nullptr; + destroyTensor = nullptr; + prepareQuery = nullptr; + executeQuery = nullptr; + mapResource = nullptr; + unmapResource = 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; + HalResAddrMapFn mapResource = nullptr; + HalResAddrUnmapFn unmapResource = nullptr; + +private: + void* runtimeHandle_ = nullptr; + void* opapiHandle_ = nullptr; +}; + +struct QuerySnapshot { + uint32_t flag = 0U; + uint32_t totalQueueCount = 0U; + aclError syncStatus = ACL_SUCCESS; + std::vector channels; +}; + +int CleanupSetCurrentContext(void*, void* context) +{ + return static_cast( + aclrtSetCurrentContext(static_cast(context))); +} + +int CleanupDestroyStream(void*, void* stream) +{ + return static_cast(aclrtDestroyStream(static_cast(stream))); +} + +int CleanupDestroyContext(void*, void* context) +{ + return static_cast(aclrtDestroyContext(static_cast(context))); +} + +int CleanupFreeDevice(void*, void* address) +{ + return static_cast(aclrtFree(address)); +} + +int CleanupDestroyTensor(void* opaque, const void* tensor) +{ + A5RuntimeApi* api = static_cast(opaque); + return api == nullptr || api->destroyTensor == nullptr + ? -1 + : api->destroyTensor(static_cast(tensor)); +} + +detail::A5QueryCleanupOps MakeQueryCleanupOps(A5RuntimeApi& api) +{ + detail::A5QueryCleanupOps ops; + ops.opaque = &api; + ops.setCurrentContext = CleanupSetCurrentContext; + ops.destroyStream = CleanupDestroyStream; + ops.destroyContext = CleanupDestroyContext; + ops.freeDevice = CleanupFreeDevice; + ops.destroyTensor = CleanupDestroyTensor; + return ops; +} + +bool AllocateTrackedBuffer(std::vector& buffers, size_t bytes, + bool zero, void*& address) +{ + address = nullptr; + if (aclrtMalloc(&address, bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) { + return false; + } + buffers.push_back(address); + return !zero || aclrtMemset(address, bytes, 0, bytes) == ACL_SUCCESS; +} + +bool CreateTrackedUint64Tensor(A5RuntimeApi& api, + detail::A5PendingQueryCleanup& cleanup, + void* address, int64_t elements, + aclTensor*& tensor) +{ + const int64_t shape[] = {elements}; + const int64_t strides[] = {1}; + tensor = api.createTensor(shape, 1U, ACL_UINT64, strides, 0, + ACL_FORMAT_ND, shape, 1U, address); + if (tensor == nullptr) { + return false; + } + cleanup.tensors.push_back(tensor); + return true; +} + +bool CheckRuntimeHealth(detail::A5PendingQueryCleanup& cleanup, void* scratch) +{ + aclrtStream stream = nullptr; + const aclError createStatus = aclrtCreateStream(&stream); + cleanup.healthStream = stream; + if (createStatus != ACL_SUCCESS || stream == nullptr) { + return false; + } + const aclError memsetStatus = aclrtMemsetAsync( + scratch, sizeof(uint64_t), 0xA5, sizeof(uint64_t), stream); + const aclError syncStatus = memsetStatus == ACL_SUCCESS + ? aclrtSynchronizeStream(stream) + : memsetStatus; + uint64_t value = 0U; + const aclError copyStatus = syncStatus == ACL_SUCCESS + ? aclrtMemcpy(&value, sizeof(value), scratch, sizeof(value), + ACL_MEMCPY_DEVICE_TO_HOST) + : syncStatus; + return memsetStatus == ACL_SUCCESS && syncStatus == ACL_SUCCESS && + copyStatus == ACL_SUCCESS && value == 0xA5A5A5A5A5A5A5A5ULL; +} + +bool FinishQuery(A5RuntimeApi& api, + detail::A5PendingQueryCleanup& cleanup, + std::vector& pending, + bool result) +{ + const bool released = detail::CleanupA5QueryResources( + cleanup, MakeQueryCleanupOps(api), cleanup.ownerContext); + if (!cleanup.Empty()) { + pending.push_back(std::move(cleanup)); + } + return result && released; +} + +bool RunBuiltinQuery(A5RuntimeApi& api, + int32_t logicalDevice, + const std::vector& streams, + std::vector& pending, + QuerySnapshot& snapshot) +{ + if (streams.empty() || streams.size() > detail::TILEXR_SDMA_A5_CHANNEL_COUNT) { + return false; + } + + detail::A5PendingQueryCleanup cleanup; + aclrtContext ownerContext = nullptr; + if (aclrtGetCurrentContext(&ownerContext) != ACL_SUCCESS || + ownerContext == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query requires an active owner context"; + return false; + } + cleanup.ownerContext = ownerContext; + + const size_t streamsBytes = streams.size() * sizeof(detail::A5BuiltinStreamInfo); + void* streamsDev = nullptr; + void* resourceDev = nullptr; + void* builtinWorkspaceDev = nullptr; + void* inputDev = nullptr; + void* outputDev = nullptr; + void* opWorkspaceDev = nullptr; + if (!AllocateTrackedBuffer(cleanup.ownerBuffers, streamsBytes, false, streamsDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + sizeof(detail::A5BuiltinOpResource), true, resourceDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + kBuiltinWorkspaceBytes, true, builtinWorkspaceDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + 2U * sizeof(uint64_t), false, inputDev) || + !AllocateTrackedBuffer(cleanup.ownerBuffers, + sizeof(uint64_t), true, outputDev) || + aclrtMemcpy(streamsDev, streamsBytes, streams.data(), streamsBytes, + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query buffer setup failed"; + return FinishQuery(api, cleanup, pending, false); + } + + detail::A5BuiltinOpResource resource {}; + resource.size = streams.size(); + resource.streamsAddress = reinterpret_cast(streamsDev); + resource.workspaceAddress = reinterpret_cast(builtinWorkspaceDev); + const uint64_t inputs[] = { + reinterpret_cast(resourceDev), + reinterpret_cast(builtinWorkspaceDev), + }; + if (aclrtMemcpy(resourceDev, sizeof(resource), &resource, sizeof(resource), + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS || + aclrtMemcpy(inputDev, sizeof(inputs), inputs, sizeof(inputs), + ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query resource upload failed"; + return FinishQuery(api, cleanup, pending, false); + } + + aclrtContext isolated = nullptr; + const aclError createContextStatus = aclrtCreateContext(&isolated, logicalDevice); + cleanup.isolatedContext = isolated; + if (createContextStatus != ACL_SUCCESS || isolated == nullptr) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA isolated query context creation failed"; + return FinishQuery(api, cleanup, pending, false); + } + + aclTensor* inputTensor = nullptr; + aclTensor* outputTensor = nullptr; + aclrtStream queryStream = nullptr; + bool queryLaunched = false; + aclError syncStatus = ACL_SUCCESS; + uint64_t opWorkspaceBytes = 0U; + aclOpExecutor* executor = nullptr; + bool ok = CreateTrackedUint64Tensor( + api, cleanup, inputDev, 2, inputTensor) && + CreateTrackedUint64Tensor(api, cleanup, outputDev, 1, outputTensor) && + api.prepareQuery(inputTensor, outputTensor, + &opWorkspaceBytes, &executor) == ACL_SUCCESS; + if (ok && opWorkspaceBytes != 0U) { + ok = AllocateTrackedBuffer(cleanup.isolatedBuffers, + static_cast(opWorkspaceBytes), false, + opWorkspaceDev); + } + if (ok) { + const aclError createStreamStatus = aclrtCreateStreamWithConfig( + &queryStream, 0, ACL_STREAM_FAST_LAUNCH | ACL_STREAM_FAST_SYNC); + cleanup.queryStream = queryStream; + ok = createStreamStatus == ACL_SUCCESS && queryStream != nullptr; + } + if (ok) { + aclrtStreamAttrValue failureMode {}; + failureMode.failureMode = 0; + ok = aclrtSetStreamAttribute( + queryStream, ACL_STREAM_ATTR_FAILURE_MODE, &failureMode) == ACL_SUCCESS; + } + if (ok) { + const aclnnStatus launchStatus = api.executeQuery( + opWorkspaceDev, opWorkspaceBytes, executor, queryStream); + queryLaunched = launchStatus == ACL_SUCCESS; + ok = queryLaunched; + if (queryLaunched) { + syncStatus = aclrtSynchronizeStream(queryStream); + } + } + + if (aclrtSetCurrentContext(ownerContext) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not restore query owner context"; + return FinishQuery(api, cleanup, pending, false); + } + if (!ok || !queryLaunched) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA built-in query launch failed"; + return FinishQuery(api, cleanup, pending, false); + } + + const size_t snapshotBytes = sizeof(detail::A5BuiltinWorkspaceHeader) + + streams.size() * sizeof(detail::A5BuiltinChannelInfo); + std::vector bytes(snapshotBytes, 0U); + if (aclrtMemcpy(bytes.data(), bytes.size(), builtinWorkspaceDev, bytes.size(), + ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query workspace download failed"; + return FinishQuery(api, cleanup, pending, 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(cleanup, outputDev)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA context health check failed after expected AICPU error"; + return FinishQuery(api, cleanup, pending, false); + } + return FinishQuery(api, cleanup, pending, 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; + A5RuntimeApi api; + aclrtContext ownerContext = nullptr; + aclrtContext restoreContext = nullptr; + bool restorePending = false; + void* workspaceDev = nullptr; + std::vector pendingQueries; + std::array channels {}; + + bool HasOwnedResources() const + { + if (workspaceDev != nullptr || !pendingQueries.empty()) { + return true; + } + for (const OwnedChannel& channel : channels) { + if (channel.mapped || channel.stream != nullptr) { + return true; + } + } + return false; + } + + void EraseCompletedQueries() + { + auto query = pendingQueries.begin(); + while (query != pendingQueries.end()) { + if (query->Empty()) { + query = pendingQueries.erase(query); + } else { + ++query; + } + } + } +}; + +TileXRA5SDMABackend::TileXRA5SDMABackend() = default; + +TileXRA5SDMABackend::~TileXRA5SDMABackend() +{ + (void)Shutdown(); +} + +bool TileXRA5SDMABackend::Init(int32_t deviceId) +{ + if (impl_ != nullptr) { + TILEXR_LOG(ERROR) << "TileXR A5 SDMA backend contains state before initialization"; + 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; + } + + if (!state->api.Load() || + state->api.getPhysicalDevice( + static_cast(deviceId), &state->physicalDevice) != 0) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA runtime discovery failed"; + return false; + } + int64_t physicalDie = -1; + if (state->api.getDeviceInfo( + static_cast(deviceId), kDeviceInfoModuleType, + kPhysicalDieInfoType, &physicalDie) != 0 || + physicalDie < 0 || static_cast(physicalDie) > + static_cast(std::numeric_limits::max())) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA physical die discovery failed"; + return false; + } + state->physicalDieId = static_cast(physicalDie); + impl_ = std::move(state); + A5RuntimeApi& api = impl_->api; + + 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 = api.mapResource( + 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, impl_->pendingQueries, 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, impl_->pendingQueries, 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; + } + } + + const detail::A5QueryCleanupOps queryCleanupOps = + MakeQueryCleanupOps(impl_->api); + for (size_t reverse = impl_->pendingQueries.size(); reverse > 0U; --reverse) { + detail::A5PendingQueryCleanup& query = impl_->pendingQueries[reverse - 1U]; + if (query.restorePending && + !detail::CleanupA5QueryResources( + query, queryCleanupOps, query.restoreContext)) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA query context restore remains pending"; + return false; + } + } + impl_->EraseCompletedQueries(); + + aclrtContext previous = nullptr; + if (aclrtGetCurrentContext(&previous) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not capture current context for cleanup"; + return false; + } + for (size_t reverse = impl_->pendingQueries.size(); reverse > 0U; --reverse) { + detail::A5PendingQueryCleanup& query = impl_->pendingQueries[reverse - 1U]; + if (!detail::CleanupA5QueryResources(query, queryCleanupOps, previous)) { + impl_->EraseCompletedQueries(); + TILEXR_LOG(WARN) << "TileXR A5 SDMA query cleanup incomplete; retained for retry"; + return false; + } + } + impl_->EraseCompletedQueries(); + + if (impl_->ownerContext == nullptr || + aclrtSetCurrentContext(impl_->ownerContext) != ACL_SUCCESS) { + TILEXR_LOG(WARN) << "TileXR A5 SDMA could not switch to owner context for cleanup"; + 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 (impl_->api.unmapResource( + 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_a5_cleanup.h b/src/comm/sdma/tilexr_sdma_a5_cleanup.h new file mode 100644 index 00000000..b5098021 --- /dev/null +++ b/src/comm/sdma/tilexr_sdma_a5_cleanup.h @@ -0,0 +1,173 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_A5_CLEANUP_H +#define TILEXR_SDMA_A5_CLEANUP_H + +#include +#include + +namespace TileXR { +namespace detail { + +using A5CleanupFn = int (*)(void*, void*); +using A5ConstCleanupFn = int (*)(void*, const void*); + +struct A5QueryCleanupOps { + void* opaque = nullptr; + A5CleanupFn setCurrentContext = nullptr; + A5CleanupFn destroyStream = nullptr; + A5CleanupFn destroyContext = nullptr; + A5CleanupFn freeDevice = nullptr; + A5ConstCleanupFn destroyTensor = nullptr; +}; + +struct A5PendingQueryCleanup { + void* ownerContext = nullptr; + void* isolatedContext = nullptr; + void* queryStream = nullptr; + void* healthStream = nullptr; + std::vector ownerBuffers; + std::vector isolatedBuffers; + std::vector tensors; + void* restoreContext = nullptr; + bool restorePending = false; + + bool Empty() const + { + return isolatedContext == nullptr && queryStream == nullptr && + healthStream == nullptr && ownerBuffers.empty() && + isolatedBuffers.empty() && tensors.empty() && !restorePending; + } +}; + +inline bool A5SetCleanupContext(const A5QueryCleanupOps& ops, void* context) +{ + return ops.setCurrentContext != nullptr && + ops.setCurrentContext(ops.opaque, context) == 0; +} + +inline bool A5ReleaseCleanupHandle(void*& handle, A5CleanupFn cleanup, + const A5QueryCleanupOps& ops) +{ + if (handle == nullptr) { + return true; + } + if (cleanup != nullptr && cleanup(ops.opaque, handle) == 0) { + handle = nullptr; + return true; + } + return false; +} + +inline bool A5ReleaseCleanupHandles(std::vector& handles, + A5CleanupFn cleanup, + const A5QueryCleanupOps& ops) +{ + bool released = true; + size_t reverse = handles.size(); + while (reverse > 0U) { + --reverse; + if (handles[reverse] != nullptr && cleanup != nullptr && + cleanup(ops.opaque, handles[reverse]) == 0) { + handles.erase(handles.begin() + static_cast(reverse)); + } else { + released = false; + } + } + return released; +} + +inline bool A5ReleaseCleanupTensors(std::vector& tensors, + const A5QueryCleanupOps& ops) +{ + bool released = true; + size_t reverse = tensors.size(); + while (reverse > 0U) { + --reverse; + if (tensors[reverse] != nullptr && ops.destroyTensor != nullptr && + ops.destroyTensor(ops.opaque, tensors[reverse]) == 0) { + tensors.erase(tensors.begin() + static_cast(reverse)); + } else { + released = false; + } + } + return released; +} + +inline bool CleanupA5QueryResources(A5PendingQueryCleanup& state, + const A5QueryCleanupOps& ops, + void* callerContext) +{ + bool released = true; + if (state.restorePending) { + if (!A5SetCleanupContext(ops, state.restoreContext)) { + return false; + } + state.restoreContext = nullptr; + state.restorePending = false; + } + + if (state.healthStream != nullptr) { + if (state.ownerContext == nullptr || + !A5SetCleanupContext(ops, state.ownerContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandle( + state.healthStream, ops.destroyStream, ops) && released; + } + } + + const bool hasIsolatedResources = state.queryStream != nullptr || + !state.isolatedBuffers.empty() || !state.tensors.empty(); + if (state.isolatedContext != nullptr) { + if (!A5SetCleanupContext(ops, state.isolatedContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandle( + state.queryStream, ops.destroyStream, ops) && released; + if (state.queryStream == nullptr) { + released = A5ReleaseCleanupHandles( + state.isolatedBuffers, ops.freeDevice, ops) && released; + released = A5ReleaseCleanupTensors(state.tensors, ops) && released; + } + if (state.queryStream == nullptr && state.tensors.empty() && + state.isolatedBuffers.empty()) { + released = A5ReleaseCleanupHandle( + state.isolatedContext, ops.destroyContext, ops) && released; + } + } + } else if (hasIsolatedResources) { + released = false; + } + + const bool isolatedReleased = state.isolatedContext == nullptr && + state.queryStream == nullptr && state.isolatedBuffers.empty() && + state.tensors.empty(); + if (!state.ownerBuffers.empty() && state.healthStream == nullptr && + isolatedReleased) { + if (state.ownerContext == nullptr || + !A5SetCleanupContext(ops, state.ownerContext)) { + released = false; + } else { + released = A5ReleaseCleanupHandles( + state.ownerBuffers, ops.freeDevice, ops) && released; + } + } else if (!state.ownerBuffers.empty()) { + released = false; + } + + if (!A5SetCleanupContext(ops, callerContext)) { + state.restoreContext = callerContext; + state.restorePending = true; + released = false; + } + return released && state.Empty(); +} + +} // namespace detail +} // namespace TileXR + +#endif // TILEXR_SDMA_A5_CLEANUP_H diff --git a/src/comm/sdma/tilexr_sdma_transport.cpp b/src/comm/sdma/tilexr_sdma_transport.cpp index 08ee0b6a..9b770757 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,11 @@ bool TileXRSDMATransport::EnvEnabled() int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) { - Shutdown(); + if (impl_ != nullptr) { + TILEXR_LOG(ERROR) << "TileXR SDMA transport contains state before initialization"; + lastStatus_ = SDMAInitStatus::INIT_FAILED; + return TILEXR_ERROR_INTERNAL; + } options_ = options; available_ = false; workspaceDev_ = nullptr; @@ -54,19 +62,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 +125,35 @@ int TileXRSDMATransport::Init(const TileXRSDMATransportOptions& options) #else lastStatus_ = SDMAInitStatus::PTO_UNAVAILABLE; TILEXR_LOG(WARN) << "TileXR SDMA PTO headers unavailable at build time"; + impl_.reset(); return TILEXR_SUCCESS; #endif } -void TileXRSDMATransport::Shutdown() +bool TileXRSDMATransport::Shutdown() { -#if TILEXR_HAVE_PTO_SDMA + bool cleanupComplete = true; if (impl_ != nullptr) { - impl_->workspaceManager.Finalize(); - impl_.reset(); - } + if (impl_->a5Backend != nullptr) { + if (impl_->a5Backend->Shutdown()) { + impl_->a5Backend.reset(); + } else { + cleanupComplete = false; + } + } +#if TILEXR_HAVE_PTO_SDMA + if (impl_->ptoInitialized) { + impl_->workspaceManager.Finalize(); + impl_->ptoInitialized = false; + } #endif + if (cleanupComplete) { + impl_.reset(); + } + } available_ = false; workspaceDev_ = nullptr; + return cleanupComplete; } bool TileXRSDMATransport::IsAvailable() const diff --git a/src/comm/sdma/tilexr_sdma_transport.h b/src/comm/sdma/tilexr_sdma_transport.h index 2f8ed9f9..93f778c3 100644 --- a/src/comm/sdma/tilexr_sdma_transport.h +++ b/src/comm/sdma/tilexr_sdma_transport.h @@ -6,6 +6,7 @@ #ifndef TILEXR_SDMA_TRANSPORT_H #define TILEXR_SDMA_TRANSPORT_H +#include #include #include "comm_args.h" @@ -13,6 +14,35 @@ namespace TileXR { +namespace detail { + +enum class SDMABackendKind : uint32_t { + UNSUPPORTED = 0U, + PTO = 1U, + A5_DIRECT = 2U, +}; + +inline bool SDMASocHasPrefix(const char* socName, const char* prefix) +{ + return socName != nullptr && + std::strncmp(socName, prefix, std::strlen(prefix)) == 0; +} + +inline SDMABackendKind ClassifySDMABackend(const char* socName) +{ + if (SDMASocHasPrefix(socName, "Ascend950")) { + return SDMABackendKind::A5_DIRECT; + } + if (SDMASocHasPrefix(socName, "Ascend910B") || + SDMASocHasPrefix(socName, "Ascend910A") || + SDMASocHasPrefix(socName, "Ascend910_93")) { + return SDMABackendKind::PTO; + } + return SDMABackendKind::UNSUPPORTED; +} + +} // namespace detail + struct TileXRSDMATransportOptions { int devId = 0; }; @@ -25,7 +55,7 @@ class TileXRSDMATransport { TileXRSDMATransport& operator=(const TileXRSDMATransport&) = delete; int Init(const TileXRSDMATransportOptions& options); - void Shutdown(); + bool Shutdown(); bool IsAvailable() const; GM_ADDR GetWorkspaceDev() const; diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp index 69736d0d..05f6e756 100644 --- a/src/comm/tilexr_comm.cpp +++ b/src/comm/tilexr_comm.cpp @@ -50,8 +50,6 @@ constexpr int TILEXR_INIT_TIMEOUT = 600; static map g_localPeerMemMap; static map g_devList; static std::mutex g_mtx; -static std::mutex g_sdmaMtx; -static bool g_sdmaUnavailable = false; // 如果是互联的链路，返回false； 对910B2C那些不互联的链路，返回true @@ -192,15 +190,11 @@ int TileXRComm::ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &st int TileXRComm::InitSDMA() { - { - lock_guard lock(g_sdmaMtx); - if (g_sdmaUnavailable) { - TILEXR_LOG(INFO) << "InitSDMA skipped after previous SDMA init failure"; - sdmaInitStatus_ = SDMAInitStatus::PTO_UNAVAILABLE; - return TILEXR_SUCCESS; - } + if (sdmaTransport_ != nullptr) { + TILEXR_LOG(ERROR) << "TileXR SDMA transport exists before initialization"; + sdmaInitStatus_ = SDMAInitStatus::INIT_FAILED; + return TILEXR_ERROR_INTERNAL; } - sdmaTransport_.reset(new (nothrow) TileXRSDMATransport()); if (sdmaTransport_ == nullptr) { TILEXR_LOG(WARN) << "TileXRSDMATransport allocation failed, SDMA disabled"; @@ -212,13 +206,17 @@ int TileXRComm::InitSDMA() options.devId = devId_; int ret = sdmaTransport_->Init(options); sdmaInitStatus_ = sdmaTransport_->GetLastStatus(); - if (ret != TILEXR_SUCCESS || !sdmaTransport_->IsAvailable()) { + if (ret != TILEXR_SUCCESS) { + TILEXR_LOG(ERROR) << "TileXR SDMA transport initialization failed"; + return ret; + } + if (!sdmaTransport_->IsAvailable()) { if (sdmaInitStatus_ != SDMAInitStatus::DISABLED_BY_ENV) { TILEXR_LOG(WARN) << "TileXR SDMA init unavailable, status " << static_cast(sdmaInitStatus_); - lock_guard lock(g_sdmaMtx); - g_sdmaUnavailable = true; } - sdmaTransport_.reset(); + if (sdmaTransport_->Shutdown()) { + sdmaTransport_.reset(); + } sdmaWorkspaceDev_ = nullptr; commArgs_.sdmaWorkspacePtr = nullptr; return TILEXR_SUCCESS; @@ -231,7 +229,9 @@ int TileXRComm::InitSDMA() commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; sdmaWorkspaceDev_ = nullptr; - sdmaTransport_.reset(); + if (sdmaTransport_->Shutdown()) { + sdmaTransport_.reset(); + } return TILEXR_SUCCESS; } @@ -242,16 +242,30 @@ int TileXRComm::InitSDMA() return TILEXR_SUCCESS; } -void TileXRComm::ResetSDMAState() +bool TileXRComm::ResetSDMAState() { + if (sdmaTransport_ == nullptr) { + return true; + } commArgs_.extraFlag &= ~ExtraFlag::SDMA; commArgs_.sdmaWorkspacePtr = nullptr; + if (UpdateCommArgsDev() != TILEXR_SUCCESS) { + return false; + } 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 @@ -886,8 +900,8 @@ TileXRComm::~TileXRComm() } FreePeerMem(commArgs_.dumpAddr); FreePeerMem(peerMem_[rank_]); + (void)ResetSDMAState(); FreePeerMem(commArgsPtr_); - ResetSDMAState(); } TileXRComm::TileXRComm(int rank, int rankSize) : rank_(rank), rankSize_(rankSize) diff --git a/src/comm/tilexr_comm.h b/src/comm/tilexr_comm.h index 8b8af295..5a4c5a42 100644 --- a/src/comm/tilexr_comm.h +++ b/src/comm/tilexr_comm.h @@ -51,6 +51,7 @@ class TileXRComm { bool IsSDMAAvailable() const; GM_ADDR GetSDMAWorkspacePtr() const; SDMAInitStatus GetSDMAInitStatus() const; + bool PrepareDestroy(); std::string PrintDFX(); friend class Lccl; friend class Lcoc; @@ -78,7 +79,7 @@ class TileXRComm { static int ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &state, void *userData); int InitSDMA(); int UpdateCommArgsDev(); - void ResetSDMAState(); + bool ResetSDMAState(); private: int rank_ = 0; // global rank id diff --git a/src/include/tilexr_sdma.h b/src/include/tilexr_sdma.h index db9fdf79..ac12be91 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; } @@ -70,12 +89,58 @@ __aicore__ inline uint64_t SDMACopyNbi( #endif } +__aicore__ inline uint64_t SDMACopyStridedNbi( + const __gm__ CommArgs* args, + __gm__ uint8_t* dst, + __gm__ uint8_t* src, + uint64_t bytes, + uint32_t copyCount, + uint64_t dstStrideBytes, + uint64_t srcStrideBytes, + uint32_t channelGroupIdx = TILEXR_SDMA_AUTO_CHANNEL_GROUP) +{ + if (copyCount == 1U) { + return SDMACopyNbi(args, dst, src, bytes, channelGroupIdx); + } +#if TILEXR_SDMA_A5_AICORE_COMPILE + if (!SDMAEnabled(args) || dst == nullptr || src == nullptr || bytes == 0U || + copyCount == 0U) { + return 0ULL; + } + const uint32_t resolvedGroup = SDMAResolveChannelGroup(channelGroupIdx); + return detail::A5SdmaCopyStridedNbi( + reinterpret_cast<__gm__ uint8_t*>(args->sdmaWorkspacePtr), + dst, src, bytes, copyCount, dstStrideBytes, srcStrideBytes, resolvedGroup); +#else + (void)args; + (void)dst; + (void)src; + (void)bytes; + (void)copyCount; + (void)dstStrideBytes; + (void)srcStrideBytes; + (void)channelGroupIdx; + return 0ULL; +#endif +} + __aicore__ inline bool SDMAWait( const __gm__ CommArgs* args, uint64_t eventHandle, 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..0592593d --- /dev/null +++ b/src/include/tilexr_sdma_a5.h @@ -0,0 +1,222 @@ +/* + * 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) +{ + return AscendC::ReadGmByPassDCache(&completion->generation); +} + +__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 A5SdmaCopyStridedNbi( + __gm__ uint8_t* workspaceAddress, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint64_t bytes, + uint32_t copyCount, + uint64_t destinationStrideBytes, + uint64_t sourceStrideBytes, + uint32_t channelIndex) +{ + __gm__ A5SdmaWorkspace* workspace = + reinterpret_cast<__gm__ A5SdmaWorkspace*>(workspaceAddress); + if (!A5SdmaWorkspaceValid(workspace) || destination == nullptr || source == nullptr || + !A5SdmaTransferLengthValid(bytes) || !A5SdmaChannelValid(channelIndex) || + copyCount == 0U || copyCount == 0xFFFFFFFFU) { + return 0ULL; + } + const uint64_t destinationAddress = reinterpret_cast(destination); + const uint64_t sourceAddress = reinterpret_cast(source); + if (copyCount > 1U && + (!A5SdmaStridedRangeValid( + destinationAddress, bytes, copyCount, destinationStrideBytes) || + !A5SdmaStridedRangeValid( + sourceAddress, bytes, copyCount, sourceStrideBytes))) { + return 0ULL; + } + + __gm__ A5SdmaChannel* channel = &workspace->channels[channelIndex]; + const uint32_t requiredEntries = copyCount + 1U; + if (channel->sqBase == 0U || channel->rtsqAddress == 0U || + channel->completionPayloadAddress == 0U || channel->completionRecordAddress == 0U || + channel->rtsqLength < sizeof(uint32_t) || + !A5SdmaQueueHasEntriesCapacity( + channel->head, channel->tail, channel->depth, requiredEntries) || + 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 firstIndex = channel->tail; + __gm__ A5SdmaCompletionLine* payload = reinterpret_cast<__gm__ A5SdmaCompletionLine*>( + channel->completionPayloadAddress); + payload->generation = generation; + + __gm__ A5SdmaSqe* sqBase = reinterpret_cast<__gm__ A5SdmaSqe*>(channel->sqBase); + for (uint32_t copy = 0U; copy < copyCount; ++copy) { + const uint32_t sqeIndex = A5SdmaAdvanceTailBy(firstIndex, copy, channel->depth); + const uint64_t sourceOffset = static_cast(copy) * sourceStrideBytes; + const uint64_t destinationOffset = + static_cast(copy) * destinationStrideBytes; + A5SdmaBuildSqe( + sqBase + sqeIndex, channel->streamId, + A5SdmaAdvanceTaskIdBy(channel->taskId, copy), + sourceAddress + sourceOffset, destinationAddress + destinationOffset, + static_cast(bytes)); + } + const uint32_t completionIndex = + A5SdmaAdvanceTailBy(firstIndex, copyCount, channel->depth); + A5SdmaBuildSqe( + sqBase + completionIndex, channel->streamId, + A5SdmaAdvanceTaskIdBy(channel->taskId, copyCount), + channel->completionPayloadAddress, + channel->completionRecordAddress, + TILEXR_SDMA_A5_COMPLETION_BYTES); + const uint32_t newTail = + A5SdmaAdvanceTailBy(firstIndex, requiredEntries, channel->depth); + + channel->generation = generation; + channel->tail = newTail; + channel->taskId = A5SdmaAdvanceTaskIdBy(channel->taskId, requiredEntries); + pipe_barrier(PIPE_ALL); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(payload)); + for (uint32_t entry = 0U; entry < requiredEntries; ++entry) { + const uint32_t sqeIndex = A5SdmaAdvanceTailBy(firstIndex, entry, channel->depth); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(sqBase + sqeIndex)); + } + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel)); + A5SdmaCleanCacheLine(reinterpret_cast<__gm__ uint8_t*>(channel) + 64U); + pipe_barrier(PIPE_ALL); + dsb(DSB_DDR); + A5SdmaRingDoorbell(channel->rtsqAddress, newTail); + return A5SdmaEncodeEvent(channelIndex, generation); +} + +__aicore__ inline uint64_t A5SdmaCopyNbi(__gm__ uint8_t* workspaceAddress, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint64_t bytes, + uint32_t channelIndex) +{ + return A5SdmaCopyStridedNbi( + workspaceAddress, destination, source, bytes, 1U, 0U, 0U, channelIndex); +} + +__aicore__ inline bool A5SdmaWaitEvent(__gm__ uint8_t* workspaceAddress, + uint64_t event, + uint32_t expectedChannel) +{ + 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); + while (A5SdmaReadCompletion(completion) != generation) { + } + 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..59d091be --- /dev/null +++ b/src/include/tilexr_sdma_a5_types.h @@ -0,0 +1,247 @@ +/* + * 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 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 static_cast((static_cast(tail) + 2U) % depth); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTailBy( + uint32_t tail, uint32_t entries, uint32_t depth) +{ + return static_cast( + (static_cast(tail) + static_cast(entries)) % depth); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaQueueDistance( + uint32_t head, uint32_t tail, uint32_t depth) +{ + return (tail + depth - head) % depth; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasEntriesCapacity( + uint32_t head, uint32_t tail, uint32_t depth, uint32_t requiredEntries) +{ + if (depth < 3U || head >= depth || tail >= depth) { + return false; + } + if (requiredEntries == 0U || requiredEntries >= depth) { + return false; + } + const uint32_t used = A5SdmaQueueDistance(head, tail, depth); + return requiredEntries <= depth - used - 1U; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaQueueHasCapacity( + uint32_t head, uint32_t tail, uint32_t depth) +{ + return A5SdmaQueueHasEntriesCapacity(head, tail, depth, 2U); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskId(uint32_t taskId) +{ + return (taskId + 2U) & 0xFFFFU; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaAdvanceTaskIdBy( + uint32_t taskId, uint32_t entries) +{ + return static_cast( + (static_cast(taskId) + static_cast(entries)) & 0xFFFFULL); +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE bool A5SdmaStridedRangeValid( + uint64_t address, uint64_t bytes, uint32_t copyCount, uint64_t strideBytes) +{ + if (address == 0U || bytes == 0U || copyCount == 0U) { + return false; + } + if (copyCount == 1U) { + return bytes - 1U <= 0xFFFFFFFFFFFFFFFFULL - address; + } + if (strideBytes < bytes) { + return false; + } + const uint64_t lastIndex = static_cast(copyCount - 1U); + if (strideBytes > 0xFFFFFFFFFFFFFFFFULL / lastIndex) { + return false; + } + const uint64_t lastOffset = strideBytes * lastIndex; + if (lastOffset > 0xFFFFFFFFFFFFFFFFULL - address) { + return false; + } + const uint64_t lastAddress = address + lastOffset; + return bytes - 1U <= 0xFFFFFFFFFFFFFFFFULL - lastAddress; +} + +TILEXR_SDMA_A5_HOST_DEVICE_INLINE uint32_t A5SdmaNextGeneration(uint32_t generation) +{ + return generation == 0xFFFFFFFFU ? 1U : generation + 1U; +} + +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..6efa6f99 100644 --- a/tests/sdma/CMakeLists.txt +++ b/tests/sdma/CMakeLists.txt @@ -3,8 +3,17 @@ project(TileXR_SDMA_Tests) set(CMAKE_CXX_STANDARD 14) set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_SKIP_BUILD_RPATH TRUE) option(BUILD_TILEXR_SDMA_DEMO "Build TileXR SDMA data-plane demo with Ascend C kernel" ON) set(TILEXR_SDMA_DEMO_SOC_TYPE "Ascend910B" CACHE STRING "SOC type used for the TileXR SDMA demo kernel") +set_property(CACHE TILEXR_SDMA_DEMO_SOC_TYPE PROPERTY STRINGS Ascend910B Ascend950) +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}) @@ -48,8 +57,8 @@ add_executable(test_tilexr_sdma_metadata unit/test_tilexr_sdma_metadata.cpp ) target_include_directories(test_tilexr_sdma_metadata PRIVATE - ${TILEXR_ROOT}/install/include ${TILEXR_ROOT}/src/include + ${TILEXR_ROOT}/install/include ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/ ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/ ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ @@ -77,6 +86,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 +107,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 @@ -94,13 +121,6 @@ target_compile_definitions(test_tilexr_sdma_comm_wiring PRIVATE TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" ) -add_executable(test_tilexr_sdma_source_guard - unit/test_tilexr_sdma_source_guard.cpp -) -target_compile_definitions(test_tilexr_sdma_source_guard PRIVATE - TILEXR_SOURCE_ROOT="${TILEXR_ROOT}" -) - add_executable(test_tilexr_sdma_header_compile unit/test_tilexr_sdma_header_compile.cpp ) @@ -132,13 +152,15 @@ 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 +171,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 +235,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 +252,15 @@ 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" + "${CMAKE_CURRENT_SOURCE_DIR}/demo/tilexr_sdma_benchmark_types.h" "${TILEXR_ROOT}/src/include/tilexr_sdma.h" + "${TILEXR_ROOT}/src/include/tilexr_sdma_a5.h" + "${TILEXR_ROOT}/src/include/tilexr_sdma_a5_types.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_compat.h" "${TILEXR_ROOT}/src/include/tilexr_sdma_types.h" VERBATIM @@ -241,6 +271,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 @@ -259,6 +294,31 @@ elseif(BUILD_TILEXR_SDMA_DEMO) ascend_hal ) list(APPEND INSTALL_TARGETS tilexr_sdma_demo) + + if(TILEXR_SDMA_DEMO_SOC_TYPE STREQUAL "Ascend950") + add_executable(tilexr_sdma_benchmark + demo/tilexr_sdma_benchmark.cpp + ) + add_dependencies(tilexr_sdma_benchmark tilexr_sdma_demo_kernel) + target_include_directories(tilexr_sdma_benchmark PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/demo + ${TILEXR_ROOT}/install/include + ${TILEXR_ROOT}/src/include + ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/ + ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/ + ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ + ${ASCEND_DRIVER_PATH}/kernel/inc + ) + target_link_directories(tilexr_sdma_benchmark PRIVATE ${CMAKE_CURRENT_BINARY_DIR}) + target_link_libraries(tilexr_sdma_benchmark + "${TILEXR_SDMA_DEMO_KERNEL_SO}" + ${TILEXR_LIB} + ascendcl + runtime + ascend_hal + ) + list(APPEND INSTALL_TARGETS tilexr_sdma_benchmark) + endif() install(FILES "${TILEXR_SDMA_DEMO_KERNEL_SO}" DESTINATION ${CMAKE_INSTALL_PREFIX}/lib) message(STATUS "TileXR SDMA demo enabled with ${BISHENG_EXECUTABLE}, SOC=${TILEXR_SDMA_DEMO_SOC_TYPE}") endif() diff --git a/tests/sdma/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_benchmark.cpp b/tests/sdma/demo/tilexr_sdma_benchmark.cpp new file mode 100644 index 00000000..9eed7e9d --- /dev/null +++ b/tests/sdma/demo/tilexr_sdma_benchmark.cpp @@ -0,0 +1,362 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "acl/acl.h" +#include "tilexr_api.h" +#include "tilexr_sdma_a5_types.h" +#include "tilexr_sdma_benchmark_types.h" +#include "tilexr_types.h" + +extern "C" void launch_tilexr_sdma_benchmark( + void* stream, + GM_ADDR commArgs, + GM_ADDR dst, + GM_ADDR src, + GM_ADDR samples, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots); + +namespace { +constexpr uint32_t kDefaultBytes = 4096U; +constexpr uint32_t kDefaultWarmup = 20U; +constexpr uint32_t kDefaultIterations = 100U; +constexpr uint32_t kDefaultSamples = 10U; +constexpr uint32_t kAlignmentBytes = 64U; +constexpr uint32_t kMaxSamples = 100U; +constexpr uint32_t kMaxBatchCopies = 32U; +constexpr size_t kMinWorkingSetBytes = 64U * 1024U * 1024U; +constexpr int kDeviceId = 0; +constexpr double kAscend950CyclesPerUs = 1000.0; + +bool CheckAcl(const std::string& label, aclError ret) +{ + if (ret == ACL_SUCCESS) { + return true; + } + std::cerr << "ERROR: " << label << " failed with " << ret << std::endl; + return false; +} + +bool CheckTileXR(const std::string& label, int ret) +{ + if (ret == TileXR::TILEXR_SUCCESS) { + return true; + } + std::cerr << "ERROR: " << label << " failed with " << ret << std::endl; + return false; +} + +bool ParseUint32(const char* text, uint32_t* value) +{ + if (text == nullptr || value == nullptr || text[0] == '\0') { + return false; + } + errno = 0; + char* end = nullptr; + const unsigned long long parsed = std::strtoull(text, &end, 10); + if (errno != 0 || end == text || *end != '\0' || + parsed > std::numeric_limits::max()) { + return false; + } + *value = static_cast(parsed); + return true; +} + +void FillPattern(std::vector& data) +{ + for (size_t index = 0U; index < data.size(); ++index) { + data[index] = static_cast((index * 37U + 11U) & 0xFFU); + } +} + +void Cleanup(uint8_t* src, + uint8_t* dst, + TileXR::test::SdmaBenchmarkSample* samples, + TileXRCommPtr comm, + aclrtStream stream, + bool deviceSet, + bool aclInitialized) +{ + if (src != nullptr) { + (void)aclrtFree(src); + } + if (dst != nullptr) { + (void)aclrtFree(dst); + } + if (samples != nullptr) { + (void)aclrtFree(samples); + } + if (comm != nullptr) { + (void)TileXRCommDestroy(comm); + } + if (stream != nullptr) { + (void)aclrtDestroyStream(stream); + } + if (deviceSet) { + (void)aclrtResetDevice(kDeviceId); + } + if (aclInitialized) { + (void)aclFinalize(); + } +} + +const char* StatusName(uint32_t status) +{ + switch (status) { + case TileXR::test::TILEXR_SDMA_BENCHMARK_OK: + return "ok"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_DISABLED: + return "sdma_disabled"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED: + return "warmup_submit_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED: + return "warmup_wait_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED: + return "submit_failed"; + case TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED: + return "wait_failed"; + default: + return "unknown"; + } +} +} // namespace + +int main(int argc, char** argv) +{ + uint32_t bytes = kDefaultBytes; + uint32_t warmupIterations = kDefaultWarmup; + uint32_t measuredIterations = kDefaultIterations; + uint32_t sampleCount = kDefaultSamples; + uint32_t channel = 0U; + uint32_t phaseProfile = 0U; + uint32_t batchCopies = 1U; + if (argc > 8 || + (argc >= 2 && !ParseUint32(argv[1], &bytes)) || + (argc >= 3 && !ParseUint32(argv[2], &warmupIterations)) || + (argc >= 4 && !ParseUint32(argv[3], &measuredIterations)) || + (argc >= 5 && !ParseUint32(argv[4], &sampleCount)) || + (argc >= 6 && !ParseUint32(argv[5], &channel)) || + (argc >= 7 && !ParseUint32(argv[6], &phaseProfile)) || + (argc == 8 && !ParseUint32(argv[7], &batchCopies)) || + bytes == 0U || (bytes % kAlignmentBytes) != 0U || measuredIterations == 0U || + sampleCount == 0U || sampleCount > kMaxSamples || + channel >= TileXR::detail::TILEXR_SDMA_A5_CHANNEL_COUNT || phaseProfile > 1U || + batchCopies == 0U || batchCopies > kMaxBatchCopies || + measuredIterations > std::numeric_limits::max() / batchCopies || + static_cast(bytes) * batchCopies > std::numeric_limits::max()) { + std::cerr << "ERROR: usage: tilexr_sdma_benchmark " + "[aligned-bytes [warmup [iterations [samples [channel " + "[phase-profile [batch-copies]]]]]]]" + << std::endl; + return 1; + } + + (void)setenv("TILEXR_ENABLE_SDMA", "1", 1); + bool aclInitialized = false; + bool deviceSet = false; + TileXRCommPtr comm = nullptr; + aclrtStream stream = nullptr; + uint8_t* src = nullptr; + uint8_t* dst = nullptr; + TileXR::test::SdmaBenchmarkSample* deviceSamples = nullptr; + + if (!CheckAcl("aclInit", aclInit(nullptr))) { + return 1; + } + aclInitialized = true; + if (!CheckAcl("aclrtSetDevice", aclrtSetDevice(kDeviceId))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + deviceSet = true; + const char* socName = aclrtGetSocName(); + if (socName == nullptr || std::strstr(socName, "Ascend950") == nullptr) { + std::cerr << "ERROR: device-cycle conversion is only defined here for Ascend950, got " + << (socName == nullptr ? "" : socName) << std::endl; + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + if (!CheckAcl("aclrtCreateStream", aclrtCreateStream(&stream)) || + !CheckTileXR("TileXRCommInitRankLocal", TileXRCommInitRankLocal(1, 0, &comm))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + bool sdmaAvailable = false; + GM_ADDR sdmaWorkspace = nullptr; + GM_ADDR commArgsDev = nullptr; + if (!CheckTileXR("TileXRSDMAAvailable", TileXRSDMAAvailable(comm, &sdmaAvailable)) || + !CheckTileXR("TileXRGetSDMAWorkspaceDev", TileXRGetSDMAWorkspaceDev(comm, &sdmaWorkspace)) || + !CheckTileXR("TileXRGetCommArgsDev", TileXRGetCommArgsDev(comm, commArgsDev)) || + !sdmaAvailable || sdmaWorkspace == nullptr || commArgsDev == nullptr) { + std::cerr << "ERROR: TileXR direct SDMA is unavailable" << std::endl; + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const size_t transferBytes = static_cast(bytes) * batchCopies; + const size_t workingSetSlots = std::max(1U, kMinWorkingSetBytes / transferBytes); + const size_t allocationBytes = transferBytes * workingSetSlots; + const size_t sampleBytes = static_cast(sampleCount) * + sizeof(TileXR::test::SdmaBenchmarkSample); + if (!CheckAcl("aclrtMalloc src", aclrtMalloc( + reinterpret_cast(&src), allocationBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + !CheckAcl("aclrtMalloc dst", aclrtMalloc( + reinterpret_cast(&dst), allocationBytes, ACL_MEM_MALLOC_HUGE_FIRST)) || + !CheckAcl("aclrtMalloc samples", aclrtMalloc( + reinterpret_cast(&deviceSamples), sampleBytes, ACL_MEM_MALLOC_HUGE_FIRST))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + std::vector hostSrc(allocationBytes); + std::vector hostDst(allocationBytes, 0U); + std::vector hostSamples(sampleCount); + FillPattern(hostSrc); + if (!CheckAcl("aclrtMemcpy H2D src", aclrtMemcpy( + src, allocationBytes, hostSrc.data(), hostSrc.size(), ACL_MEMCPY_HOST_TO_DEVICE)) || + !CheckAcl("aclrtMemcpy H2D dst", aclrtMemcpy( + dst, allocationBytes, hostDst.data(), hostDst.size(), ACL_MEMCPY_HOST_TO_DEVICE)) || + !CheckAcl("aclrtMemcpy H2D samples", aclrtMemcpy( + deviceSamples, sampleBytes, hostSamples.data(), sampleBytes, ACL_MEMCPY_HOST_TO_DEVICE))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + launch_tilexr_sdma_benchmark( + stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), + reinterpret_cast(deviceSamples), bytes, channel, + warmupIterations, measuredIterations, sampleCount, phaseProfile, batchCopies, + static_cast(workingSetSlots)); + if (!CheckAcl("aclrtSynchronizeStream", aclrtSynchronizeStream(stream)) || + !CheckAcl("aclrtMemcpy D2H dst", aclrtMemcpy( + hostDst.data(), hostDst.size(), dst, allocationBytes, ACL_MEMCPY_DEVICE_TO_HOST)) || + !CheckAcl("aclrtMemcpy D2H samples", aclrtMemcpy( + hostSamples.data(), sampleBytes, deviceSamples, sampleBytes, ACL_MEMCPY_DEVICE_TO_HOST))) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const uint64_t submittedBatches = static_cast(warmupIterations) + + static_cast(measuredIterations) * sampleCount; + const size_t touchedSlots = static_cast(std::min( + static_cast(workingSetSlots), submittedBatches)); + const size_t touchedBytes = touchedSlots * transferBytes; + bool valid = std::equal(hostDst.begin(), hostDst.begin() + touchedBytes, hostSrc.begin()) && + std::all_of(hostDst.begin() + touchedBytes, hostDst.end(), + [](uint8_t value) { return value == 0U; }); + std::vector cyclesPerCopy; + cyclesPerCopy.reserve(sampleCount); + const uint32_t expectedCopies = measuredIterations * batchCopies; + for (uint32_t index = 0U; index < sampleCount; ++index) { + const auto& sample = hostSamples[index]; + if (sample.status != TileXR::test::TILEXR_SDMA_BENCHMARK_OK || + sample.completed != expectedCopies || sample.cycles == 0U) { + std::cerr << "ERROR: sample " << index << " status=" << StatusName(sample.status) + << " completed=" << sample.completed << "/" << expectedCopies + << " cycles=" << sample.cycles << std::endl; + valid = false; + continue; + } + cyclesPerCopy.push_back( + static_cast(sample.cycles) / static_cast(expectedCopies)); + } + if (!valid || cyclesPerCopy.size() != sampleCount) { + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 1; + } + + const double meanCycles = std::accumulate(cyclesPerCopy.begin(), cyclesPerCopy.end(), 0.0) / + static_cast(cyclesPerCopy.size()); + double squaredDifferenceSum = 0.0; + for (double cycles : cyclesPerCopy) { + const double difference = cycles - meanCycles; + squaredDifferenceSum += difference * difference; + } + const double stddevCycles = std::sqrt( + squaredDifferenceSum / static_cast(cyclesPerCopy.size())); + const auto minmaxCycles = std::minmax_element(cyclesPerCopy.begin(), cyclesPerCopy.end()); + const double meanUs = meanCycles / kAscend950CyclesPerUs; + const double stddevUs = stddevCycles / kAscend950CyclesPerUs; + const double minUs = *minmaxCycles.first / kAscend950CyclesPerUs; + const double maxUs = *minmaxCycles.second / kAscend950CyclesPerUs; + const double bandwidthGBps = static_cast(bytes) / meanUs / 1000.0; + const double bandwidthGiBps = static_cast(bytes) / meanUs * 1000000.0 / + (1024.0 * 1024.0 * 1024.0); + double meanSubmitUs = 0.0; + double meanWaitUs = 0.0; + double meanCompletionUs = 0.0; + double meanReleaseUs = 0.0; + if (phaseProfile != 0U) { + for (const auto& sample : hostSamples) { + meanSubmitUs += static_cast(sample.submitCycles); + meanWaitUs += static_cast(sample.waitCycles); + meanCompletionUs += static_cast(sample.completionCycles); + meanReleaseUs += static_cast(sample.releaseCycles); + } + const double phaseDivisor = static_cast(sampleCount) * + static_cast(expectedCopies) * kAscend950CyclesPerUs; + meanSubmitUs /= phaseDivisor; + meanWaitUs /= phaseDivisor; + meanCompletionUs /= phaseDivisor; + meanReleaseUs /= phaseDivisor; + } + + std::cout << std::fixed << std::setprecision(6); + for (uint32_t index = 0U; index < sampleCount; ++index) { + std::cout << "SAMPLE index=" << index + << " total_cycles=" << hostSamples[index].cycles + << " cycles_per_copy=" << cyclesPerCopy[index] + << " latency_us=" << cyclesPerCopy[index] / kAscend950CyclesPerUs << std::endl; + } + std::cout << "RESULT soc=" << socName + << " bytes=" << bytes + << " channel=" << channel + << " warmup=" << warmupIterations + << " batches=" << measuredIterations + << " batch_copies=" << batchCopies + << " copies=" << expectedCopies + << " allocation_bytes=" << allocationBytes + << " working_set_slots=" << workingSetSlots + << " verified_slots=" << touchedSlots + << " samples=" << sampleCount + << " mean_cycles=" << meanCycles + << " mean_us=" << meanUs + << " stddev_us=" << stddevUs + << " min_us=" << minUs + << " max_us=" << maxUs + << " bandwidth_GBps=" << bandwidthGBps + << " bandwidth_GiBps=" << bandwidthGiBps << std::endl; + if (phaseProfile != 0U) { + std::cout << "PHASE submit_us=" << meanSubmitUs + << " wait_us=" << meanWaitUs + << " completion_us=" << meanCompletionUs + << " release_us=" << meanReleaseUs + << " combined_us=" << (meanSubmitUs + meanWaitUs) << std::endl; + } + std::cout << "PASS direct SDMA benchmark data verification" << std::endl; + + Cleanup(src, dst, deviceSamples, comm, stream, deviceSet, aclInitialized); + return 0; +} diff --git a/tests/sdma/demo/tilexr_sdma_benchmark_types.h b/tests/sdma/demo/tilexr_sdma_benchmark_types.h new file mode 100644 index 00000000..79c07f94 --- /dev/null +++ b/tests/sdma/demo/tilexr_sdma_benchmark_types.h @@ -0,0 +1,38 @@ +/* + * Copyright (c) 2024-2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_SDMA_BENCHMARK_TYPES_H +#define TILEXR_SDMA_BENCHMARK_TYPES_H + +#include + +namespace TileXR { +namespace test { + +constexpr uint32_t TILEXR_SDMA_BENCHMARK_OK = 0U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_DISABLED = 1U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED = 2U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED = 3U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED = 4U; +constexpr uint32_t TILEXR_SDMA_BENCHMARK_WAIT_FAILED = 5U; + +struct alignas(64) SdmaBenchmarkSample { + uint64_t cycles; + uint32_t completed; + uint32_t status; + uint64_t submitCycles; + uint64_t waitCycles; + uint64_t completionCycles; + uint64_t releaseCycles; + uint64_t reserved[2]; +}; + +static_assert(sizeof(SdmaBenchmarkSample) == 64U, + "SDMA benchmark sample must occupy one cache line"); + +} // namespace test +} // namespace TileXR + +#endif // TILEXR_SDMA_BENCHMARK_TYPES_H diff --git a/tests/sdma/demo/tilexr_sdma_demo.cpp b/tests/sdma/demo/tilexr_sdma_demo.cpp index cd0f0620..cadbf0e1 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,14 +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; +#ifndef TILEXR_SDMA_DEMO_A5 +#define TILEXR_SDMA_DEMO_A5 0 +#endif + bool CheckAcl(const std::string& label, aclError ret) { std::cout << label << " ret=" << ret << std::endl; @@ -69,6 +76,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 +143,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 +169,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 +238,62 @@ 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); + blocks, stream, commArgsDev, reinterpret_cast(dst), reinterpret_cast(src), + reinterpret_cast(debug), bytes, firstChannel, iterations); if (!CheckAcl("aclrtSynchronizeStream", aclrtSynchronizeStream(stream))) { 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 +303,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 +311,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..15e114d7 100644 --- a/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp +++ b/tests/sdma/demo/tilexr_sdma_demo_kernel.cpp @@ -12,38 +12,102 @@ #include "kernel_operator.h" #include "tilexr_sdma.h" +#include "tilexr_sdma_benchmark_types.h" + +__aicore__ inline uint64_t BenchmarkSubmit( + __gm__ TileXR::CommArgs* args, + __gm__ uint8_t* destination, + __gm__ uint8_t* source, + uint32_t bytes, + uint32_t channel, + uint32_t batchCopies) +{ + return TileXR::SDMACopyStridedNbi( + args, destination, source, static_cast(bytes), batchCopies, + static_cast(bytes), static_cast(bytes), channel); +} extern "C" __global__ __aicore__ void tilexr_sdma_copy_kernel( GM_ADDR commArgsGM, 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 +119,169 @@ 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, firstChannel, iterations); +} + +extern "C" __global__ __aicore__ void tilexr_sdma_benchmark_kernel( + GM_ADDR commArgsGM, + GM_ADDR dstGM, + GM_ADDR srcGM, + GM_ADDR samplesGM, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots) +{ + if ASCEND_IS_AIV { + auto args = reinterpret_cast<__gm__ TileXR::CommArgs*>(commArgsGM); + auto dst = reinterpret_cast<__gm__ uint8_t*>(dstGM); + auto src = reinterpret_cast<__gm__ uint8_t*>(srcGM); + auto samples = reinterpret_cast<__gm__ TileXR::test::SdmaBenchmarkSample*>(samplesGM); + const uint64_t batchBytes = static_cast(bytes) * batchCopies; + uint64_t operationIndex = 0U; + + uint32_t status = TileXR::SDMAEnabled(args) && workingSetSlots != 0U + ? TileXR::test::TILEXR_SDMA_BENCHMARK_OK + : TileXR::test::TILEXR_SDMA_BENCHMARK_DISABLED; + for (uint32_t iteration = 0U; + iteration < warmupIterations && status == TileXR::test::TILEXR_SDMA_BENCHMARK_OK; + ++iteration) { + const uint64_t offset = (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_SUBMIT_FAILED; + } else if (!TileXR::SDMAWait(args, event, channel)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WARMUP_WAIT_FAILED; + } + } + + pipe_barrier(PIPE_ALL); + for (uint32_t sampleIndex = 0U; sampleIndex < sampleCount; ++sampleIndex) { + uint32_t completed = 0U; + uint64_t startCycle = 0U; + uint64_t endCycle = 0U; + uint64_t submitCycles = 0U; + uint64_t waitCycles = 0U; + uint64_t completionCycles = 0U; + uint64_t releaseCycles = 0U; + if (status == TileXR::test::TILEXR_SDMA_BENCHMARK_OK) { + startCycle = static_cast(AscendC::GetSystemCycle()); + if (phaseProfile != 0U) { + for (uint32_t iteration = 0U; iteration < measuredIterations; ++iteration) { + const uint64_t phaseStart = + static_cast(AscendC::GetSystemCycle()); + const uint64_t offset = + (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + const uint64_t submitEnd = + static_cast(AscendC::GetSystemCycle()); + submitCycles += submitEnd - phaseStart; + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED; + break; + } +#if TILEXR_SDMA_A5_AICORE_COMPILE + uint32_t eventChannel = 0U; + uint32_t eventGeneration = 0U; + if (!TileXR::detail::A5SdmaDecodeEvent( + event, eventChannel, eventGeneration)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + auto workspace = reinterpret_cast< + __gm__ TileXR::detail::A5SdmaWorkspace*>(args->sdmaWorkspacePtr); + auto completion = reinterpret_cast< + __gm__ TileXR::detail::A5SdmaCompletionLine*>( + workspace->channels[eventChannel].completionRecordAddress); + while (TileXR::detail::A5SdmaReadCompletion(completion) != + eventGeneration) { + } + const uint64_t completionEnd = + static_cast(AscendC::GetSystemCycle()); +#else + const uint64_t completionEnd = submitEnd; +#endif + const bool waitOk = TileXR::SDMAWait(args, event, channel); + const uint64_t waitEnd = + static_cast(AscendC::GetSystemCycle()); + waitCycles += waitEnd - submitEnd; + completionCycles += completionEnd - submitEnd; + releaseCycles += waitEnd - completionEnd; + if (!waitOk) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + completed += batchCopies; + } + } else { + for (uint32_t iteration = 0U; iteration < measuredIterations; ++iteration) { + const uint64_t offset = + (operationIndex++ % workingSetSlots) * batchBytes; + const uint64_t event = BenchmarkSubmit( + args, dst + offset, src + offset, bytes, channel, batchCopies); + if (event == 0U) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_SUBMIT_FAILED; + break; + } + if (!TileXR::SDMAWait(args, event, channel)) { + status = TileXR::test::TILEXR_SDMA_BENCHMARK_WAIT_FAILED; + break; + } + completed += batchCopies; + } + } + pipe_barrier(PIPE_ALL); + endCycle = static_cast(AscendC::GetSystemCycle()); + } + samples[sampleIndex].cycles = endCycle - startCycle; + samples[sampleIndex].completed = completed; + samples[sampleIndex].status = status; + samples[sampleIndex].submitCycles = submitCycles; + samples[sampleIndex].waitCycles = waitCycles; + samples[sampleIndex].completionCycles = completionCycles; + samples[sampleIndex].releaseCycles = releaseCycles; + } + +#if TILEXR_SDMA_A5_AICORE_COMPILE + pipe_barrier(PIPE_ALL); + for (uint32_t sampleIndex = 0U; sampleIndex < sampleCount; ++sampleIndex) { + TileXR::detail::A5SdmaCleanCacheLine( + reinterpret_cast<__gm__ uint8_t*>(samples + sampleIndex)); + } + pipe_barrier(PIPE_ALL); + dsb(DSB_DDR); +#endif + } +} + +extern "C" void launch_tilexr_sdma_benchmark( + void* stream, + GM_ADDR commArgs, + GM_ADDR dst, + GM_ADDR src, + GM_ADDR samples, + uint32_t bytes, + uint32_t channel, + uint32_t warmupIterations, + uint32_t measuredIterations, + uint32_t sampleCount, + uint32_t phaseProfile, + uint32_t batchCopies, + uint32_t workingSetSlots) { - tilexr_sdma_copy_kernel<<>>(commArgs, dst, src, debug, bytes); + tilexr_sdma_benchmark_kernel<<<1, nullptr, stream>>>( + commArgs, dst, src, samples, bytes, channel, + warmupIterations, measuredIterations, sampleCount, phaseProfile, + batchCopies, workingSetSlots); } diff --git a/tests/sdma/run_tests.sh b/tests/sdma/run_tests.sh index 793cb9bb..d179edd7 100755 --- a/tests/sdma/run_tests.sh +++ b/tests/sdma/run_tests.sh @@ -63,8 +63,8 @@ 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" if [ "${HAL_AVAILABLE}" -eq 1 ]; then 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..b76c9a92 --- /dev/null +++ b/tests/sdma/unit/test_tilexr_sdma_a5_validation.cpp @@ -0,0 +1,265 @@ +#include +#include +#include +#include + +#include "sdma/tilexr_sdma_a5_cleanup.h" +#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* Handle(uintptr_t value) +{ + return reinterpret_cast(value); +} + +struct FakeCleanupRuntime { + std::string failedOperation; + void* failedHandle = nullptr; + bool failureConsumed = false; + std::vector calls; + + int Call(const char* operation, void* handle) + { + calls.push_back(operation); + if (!failureConsumed && failedOperation == operation && + failedHandle == handle) { + failureConsumed = true; + return 1; + } + return 0; + } +}; + +int FakeSetContext(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("set", handle); +} + +int FakeDestroyStream(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("stream", handle); +} + +int FakeDestroyContext(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("context", handle); +} + +int FakeFreeDevice(void* opaque, void* handle) +{ + return static_cast(opaque)->Call("free", handle); +} + +int FakeDestroyTensor(void* opaque, const void* handle) +{ + return static_cast(opaque)->Call( + "tensor", const_cast(handle)); +} + +TileXR::detail::A5QueryCleanupOps FakeCleanupOps(FakeCleanupRuntime& runtime) +{ + TileXR::detail::A5QueryCleanupOps ops; + ops.opaque = &runtime; + ops.setCurrentContext = FakeSetContext; + ops.destroyStream = FakeDestroyStream; + ops.destroyContext = FakeDestroyContext; + ops.freeDevice = FakeFreeDevice; + ops.destroyTensor = FakeDestroyTensor; + return ops; +} + +TileXR::detail::A5PendingQueryCleanup FullCleanupState() +{ + TileXR::detail::A5PendingQueryCleanup state; + state.ownerContext = Handle(1U); + state.isolatedContext = Handle(2U); + state.queryStream = Handle(3U); + state.healthStream = Handle(4U); + state.ownerBuffers = {Handle(5U), Handle(6U)}; + state.isolatedBuffers = {Handle(7U)}; + state.tensors = {Handle(8U), Handle(9U)}; + return state; +} + +bool ContainsHandle(const TileXR::detail::A5PendingQueryCleanup& state, + void* handle) +{ + if (state.isolatedContext == handle || state.queryStream == handle || + state.healthStream == handle || state.restoreContext == handle) { + return true; + } + for (void* buffer : state.ownerBuffers) { + if (buffer == handle) { + return true; + } + } + for (void* buffer : state.isolatedBuffers) { + if (buffer == handle) { + return true; + } + } + for (const void* tensor : state.tensors) { + if (tensor == handle) { + return true; + } + } + return false; +} + +void TestCompleteAndExpectedPartialClassification() +{ + using namespace TileXR::detail; + 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); +} + +void TestCleanupFailuresRetainHandlesForRetry() +{ + struct FailureCase { + const char* operation; + void* handle; + }; + const std::vector failures = { + {"set", Handle(2U)}, + {"stream", Handle(4U)}, + {"stream", Handle(3U)}, + {"free", Handle(7U)}, + {"tensor", Handle(9U)}, + {"context", Handle(2U)}, + {"free", Handle(6U)}, + }; + for (const FailureCase& failure : failures) { + FakeCleanupRuntime runtime; + runtime.failedOperation = failure.operation; + runtime.failedHandle = failure.handle; + TileXR::detail::A5PendingQueryCleanup state = FullCleanupState(); + const TileXR::detail::A5QueryCleanupOps ops = FakeCleanupOps(runtime); + CHECK_TRUE(!TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(runtime.failureConsumed); + CHECK_TRUE(ContainsHandle(state, failure.handle)); + CHECK_TRUE(!state.Empty()); + + CHECK_TRUE(TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(state.Empty()); + } +} + +void TestCleanupRestoreFailureIsRetryable() +{ + FakeCleanupRuntime runtime; + runtime.failedOperation = "set"; + runtime.failedHandle = Handle(1U); + TileXR::detail::A5PendingQueryCleanup state; + state.ownerContext = Handle(1U); + state.isolatedContext = Handle(2U); + state.queryStream = Handle(3U); + const TileXR::detail::A5QueryCleanupOps ops = FakeCleanupOps(runtime); + + CHECK_TRUE(!TileXR::detail::CleanupA5QueryResources( + state, ops, state.ownerContext)); + CHECK_TRUE(state.restorePending); + CHECK_TRUE(state.restoreContext == state.ownerContext); + CHECK_TRUE(!state.Empty()); + + CHECK_TRUE(TileXR::detail::CleanupA5QueryResources( + state, ops, state.restoreContext)); + CHECK_TRUE(state.Empty()); +} + +} // namespace + +int main() +{ + TestCompleteAndExpectedPartialClassification(); + TestPartialClassificationFailsClosed(); + TestCompleteClassificationRequiresFinishedHeader(); + TestCleanupFailuresRetainHandlesForRetry(); + TestCleanupRestoreFailureIsRetryable(); + 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..dec5ee0d 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,82 @@ 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_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(A5SdmaAdvanceTailBy(7U, 3U, 8U), 2U); + CHECK_EQ(A5SdmaAdvanceTailBy(0xFFFFFFFEU, 3U, 0xFFFFFFFFU), 2U); + CHECK_EQ(A5SdmaQueueDistance(2U, 1U, 3U), 2U); + CHECK_TRUE(A5SdmaQueueHasCapacity(0U, 0U, 3U)); + CHECK_TRUE(A5SdmaQueueHasCapacity(1U, 2U, 8U)); + CHECK_TRUE(!A5SdmaQueueHasCapacity(0U, 2U, 3U)); + CHECK_TRUE(!A5SdmaQueueHasCapacity(3U, 0U, 3U)); + CHECK_TRUE(A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 7U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 8U)); + CHECK_TRUE(A5SdmaQueueHasEntriesCapacity(2U, 6U, 8U, 3U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(2U, 6U, 8U, 4U)); + CHECK_TRUE(!A5SdmaQueueHasEntriesCapacity(0U, 0U, 8U, 0U)); + CHECK_EQ(A5SdmaAdvanceTaskId(0xFFFFU), 1U); + CHECK_EQ(A5SdmaAdvanceTaskIdBy(0xFFFFU, 17U), 16U); + CHECK_TRUE(!A5SdmaTransferLengthValid(0U)); + CHECK_TRUE(A5SdmaTransferLengthValid(1U)); + CHECK_TRUE(A5SdmaTransferLengthValid(TILEXR_SDMA_A5_MAX_TRANSFER_BYTES)); + 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); + + constexpr uint64_t maxAddress = 0xFFFFFFFFFFFFFFFFULL; + CHECK_TRUE(A5SdmaStridedRangeValid(0x1000U, 64U, 16U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0x1000U, 64U, 16U, 63U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0x1000U, 64U, 0U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(0U, 64U, 1U, 0U)); + CHECK_TRUE(A5SdmaStridedRangeValid(maxAddress - 63U, 64U, 1U, 0U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(maxAddress - 62U, 64U, 1U, 0U)); + CHECK_TRUE(A5SdmaStridedRangeValid(maxAddress - 127U, 64U, 2U, 64U)); + CHECK_TRUE(!A5SdmaStridedRangeValid(maxAddress - 126U, 64U, 2U, 64U)); +} + +void TestA5EventHelpers() +{ + 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 +135,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 deleted file mode 100644 index cfe96e35..00000000 --- a/tests/sdma/unit/test_tilexr_sdma_source_guard.cpp +++ /dev/null @@ -1,123 +0,0 @@ -#include -#include -#include -#include -#include - -namespace { - -int g_failures = 0; - -std::string RepoPath(const std::string& path) -{ -#ifdef TILEXR_SOURCE_ROOT - return std::string(TILEXR_SOURCE_ROOT) + "/" + path; -#else - return path; -#endif -} - -std::string ReadFile(const std::string& path, bool required = true) -{ - std::ifstream input(RepoPath(path).c_str()); - if (!input.is_open()) { - if (required) { - std::cerr << "failed to open " << RepoPath(path) << std::endl; - ++g_failures; - } - return {}; - } - std::ostringstream buffer; - buffer << input.rdbuf(); - return buffer.str(); -} - -void CheckNoNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - const auto pos = text.find(needle); - if (pos != std::string::npos) { - std::cerr << "unexpected dependency in " << path << ": " << needle - << " at byte " << pos << std::endl; - ++g_failures; - } -} - -void CheckNeedle(const std::string& path, const std::string& text, const std::string& needle) -{ - if (text.find(needle) == std::string::npos) { - std::cerr << "expected text not found in " << path << ": " << needle << std::endl; - ++g_failures; - } -} - -void TestCommSourcesDoNotUseShmem() -{ - const std::vector paths = { - "src/comm/CMakeLists.txt", - "src/comm/tilexr_comm.cpp", - "src/comm/comm_wrap.cpp", - "src/comm/tilexr_comm.h", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::vector forbidden = { - "shmem", - "shmem.h", - "libshmem", - "aclshmem", - "ACLSHMEM", - }; - for (const auto& path : paths) { - const auto text = ReadFile(path); - for (const auto& needle : forbidden) { - CheckNoNeedle(path, text, needle); - } - } -} - -void TestOnlyCompatIncludesSdmaIntrinsics() -{ - const std::vector paths = { - "src/include/tilexr_sdma.h", - "src/include/tilexr_sdma_compat.h", - "src/include/tilexr_sdma_types.h", - "src/include/comm_args.h", - "src/comm/tilexr_comm.cpp", - "src/comm/tilexr_comm.h", - "src/comm/comm_wrap.cpp", - "src/comm/sdma/tilexr_sdma_transport.cpp", - "src/comm/sdma/tilexr_sdma_transport.h", - }; - const std::string ptoPrefix = "pto/npu/comm/async/sdma/"; - const std::string intrinHeader = ptoPrefix + "sdma_async_intrin.hpp"; - const std::string workspaceHeader = ptoPrefix + "sdma_workspace_manager.hpp"; - for (const auto& path : paths) { - const bool required = path != "src/include/tilexr_sdma.h"; - const auto text = ReadFile(path, required); - if (path == "src/include/tilexr_sdma_compat.h") { - CheckNeedle(path, text, intrinHeader); - CheckNoNeedle(path, text, workspaceHeader); - continue; - } - if (path == "src/comm/sdma/tilexr_sdma_transport.cpp") { - CheckNeedle(path, text, workspaceHeader); - CheckNoNeedle(path, text, intrinHeader); - continue; - } - CheckNoNeedle(path, text, ptoPrefix); - } -} - -} // namespace - -int main() -{ - TestCommSourcesDoNotUseShmem(); - TestOnlyCompatIncludesSdmaIntrinsics(); - if (g_failures != 0) { - std::cerr << g_failures << " SDMA source guard checks failed" << std::endl; - return 1; - } - std::cout << "TileXR SDMA source guard checks passed" << std::endl; - return 0; -} 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;