From ade7f81367f218a1c65930b8b46a4456dc966306 Mon Sep 17 00:00:00 2001 From: Kur0x Date: Tue, 28 Jul 2026 11:15:53 +0800 Subject: [PATCH] feat(ccu): add direct runtime engine --- CMakeLists.txt | 1 - docs/CCU_ASCENDC_HANDOFF.md | 222 ++ src/comm/CMakeLists.txt | 4 + .../ccu/tilexr_ccu_direct_orchestrator.cpp | 1199 ++++++++++ src/comm/ccu/tilexr_ccu_direct_orchestrator.h | 137 ++ src/comm/ccu/tilexr_ccu_direct_runtime.cpp | 1529 ++++++++++++ src/comm/ccu/tilexr_ccu_direct_runtime.h | 152 ++ .../test_tilexr_ccu_direct_orchestrator.py | 1724 ++++++++++++++ ...est_tilexr_ccu_ra_custom_channel_loader.py | 2121 +++++++++++++++++ 9 files changed, 7088 insertions(+), 1 deletion(-) create mode 100644 docs/CCU_ASCENDC_HANDOFF.md create mode 100644 src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp create mode 100644 src/comm/ccu/tilexr_ccu_direct_orchestrator.h create mode 100644 src/comm/ccu/tilexr_ccu_direct_runtime.cpp create mode 100644 src/comm/ccu/tilexr_ccu_direct_runtime.h create mode 100644 tests/ccu/test_tilexr_ccu_direct_orchestrator.py create mode 100644 tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py diff --git a/CMakeLists.txt b/CMakeLists.txt index 6f80a96c..1c300f5f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -58,7 +58,6 @@ include_directories( ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/profiling/ ${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/ ${ASCEND_HOME_PATH}/${ARCH}-linux/include/ - ${ASCEND_HOME_PATH}/${ARCH}-linux/include/hccl/ ${ASCEND_HOME_PATH}/${ARCH}-linux/include/experiment ${ASCEND_HOME_PATH}/${ARCH}-linux/include/experiment/runtime/ ${ASCEND_HOME_PATH}/${ARCH}-linux/include/experiment/msprof/ diff --git a/docs/CCU_ASCENDC_HANDOFF.md b/docs/CCU_ASCENDC_HANDOFF.md new file mode 100644 index 00000000..b38f699f --- /dev/null +++ b/docs/CCU_ASCENDC_HANDOFF.md @@ -0,0 +1,222 @@ +# TileXR Direct CCU Handoff + +This document is the maintained handoff for TileXR-owned direct CCU support. +It replaces the earlier investigation log. Keep it factual and update it only +when the production path, public API, or validation gate changes. + +## Current Status + +TileXR can prepare, install, submit, and validate a two-rank direct CCU flow +without a production dependency on hcomm, HCCL private CCU producers, or +`libmc2_client.so`. + +Validated path: + +- Direct CCU repository and mission installation. +- Mission/key/taskInfo generation for `rtCCULaunch`. +- XN, CKE, channel, PFE, and jetty lower-layer resource installation. +- Barrier smoke through CCU synchronization instructions. +- P2P data-plane validation through CCU memory-copy microcode. + +Completion evidence from the 950 validation server: + +```text +tilexr_ccu_direct_evening_smoke finalStatus prepare=pass submit=pass barrier=pass p2p=pass +tilexr_ccu_direct_smoke p2pCcuCopy ... mismatches=0 ... passed=1 +TileXRDirectCcuTrace program.sync[5] decoded=TransRmtMemToLocMem ... +TileXR CCU dependency guard passed: no hcomm/HCCL private CCU dependency or symbol reference +``` + +## Production Boundary + +`src/comm` must not link or include hcomm or HCCL private CCU producer APIs. +The allowed production dependency surface is CANN runtime/ACL plus TileXR-owned +CCU code under `src/comm/ccu`. + +Keep these out of production code: + +- `libhcomm.so`, `libhccl_v2.so`, `libhccl_fwk.so`, `libmc2_client.so` +- `HcclGetCcuTaskInfo`, `HcomGetCcuTaskInfo` +- hcomm private resource allocators, repositories, and channel abstractions +- AscendC kernel-side `Hccl` integration until TileXR has + a stable TileXR-owned context producer for that exact ABI + +Run the dependency guard after every CCU production edit: + +```bash +source scripts/common_env.sh +bash tests/ccu/check_tile_comm_no_hcomm_deps.sh build/src/comm/libtile-comm.so +``` + +## Architecture + +Main modules: + +- `tilexr_ccu_direct_runtime.*`: owns direct runtime interaction and basic CCU + information discovery. +- `tilexr_ccu_driver_adapter.*`: wraps the low-level custom-channel driver + operations used for CCU resource and repository installation. +- `tilexr_ccu_hccp_loader.*`: resolves TileXR-owned HCCP/RA entry points. +- `tilexr_ccu_install_provider.*`: installs lower-layer CCU resources, + repository images, and missions. +- `tilexr_ccu_repository.*`: builds and uploads the instruction repository. +- `tilexr_ccu_resource_allocator.*`: reserves mission, instruction, XN, CKE, + GSA, and channel ranges. +- `tilexr_ccu_producer_plan.*`: builds the direct CCU program and submit tasks. +- `tilexr_ccu_barrier_program.*`: emits synchronization microcode. +- `tilexr_ccu_memory_program.*`: emits CCU memory-copy microcode. +- `tilexr_ccu_direct_orchestrator.*`: joins allocation, lower-layer planning, + repository installation, mission installation, and launch-package creation. + +Host integration: + +- `TileXRComm::PrepareDirectCcuInstallAttempt(...)` prepares generic direct CCU + tasks. +- `TileXRComm::PrepareDirectCcuMemoryCopyInstallAttempt(...)` prepares P2P CCU + copy tasks. +- `TileXRCommPrepareDirectCcu(...)` and + `TileXRCommPrepareDirectCcuMemoryCopy(...)` expose the public C API. + +Runtime launch: + +- The final launch path uses `TileXRDirectCcuSubmitPrepared(...)` or + `TileXRDirectCcuSubmitPreparedTask(...)`. +- `rtCCULaunch` is reached through `libruntime.so`; no hcomm launch wrapper is + required. + +## Public API + +`src/include/tilexr_api.h` exposes: + +```c +int TileXRCommPrepareDirectCcu( + TileXRCommPtr comm, + const TileXRDirectCcuPrepareOptions* options, + TileXRDirectCcuPreparedTasksPtr* prepared, + TileXRDirectCcuPrepareReport* report); + +int TileXRCommPrepareDirectCcuMemoryCopy( + TileXRCommPtr comm, + const TileXRDirectCcuMemoryCopyPrepareOptions* options, + TileXRDirectCcuPreparedTasksPtr* prepared, + TileXRDirectCcuPrepareReport* report); + +int TileXRDirectCcuSubmitPrepared( + TileXRDirectCcuPreparedTasksPtr prepared, + aclrtStream stream, + TileXRDirectCcuSubmitReport* report); + +int TileXRDirectCcuSubmitPreparedTask( + TileXRDirectCcuPreparedTasksPtr prepared, + uint32_t taskIndex, + aclrtStream stream, + TileXRDirectCcuSubmitReport* report); + +int TileXRDirectCcuDestroyPrepared(TileXRDirectCcuPreparedTasksPtr prepared); +``` + +Memory-copy direction constants: + +```c +TILEXR_DIRECT_CCU_MEMORY_COPY_REMOTE_TO_LOCAL +TILEXR_DIRECT_CCU_MEMORY_COPY_LOCAL_TO_REMOTE +``` + +The direct P2P smoke currently validates remote-to-local transfer on both ranks. + +## Key Lessons + +- The AscendC CCU `HcclCombineOpParam` route is not a validated TileXR + integration path. Public MC2 allocation probes returned AICPU/MC2 context + shapes, not the CCU context consumed by `HCCL_SERVER_TYPE_CCU`. +- `rtCcuTaskInfo_t` and `rtCCULaunch` are reachable through CANN runtime, but + the hard part is producing the repository, mission key, lower-layer resources, + and task fields correctly. +- Repository installation requires the RA context resource-window route. A zero + or synthetic resource-window token can let lower-layer calls appear partially + valid while `SET_INSTRUCTION` still fails. +- The passing lower-layer route uses hcomm-compatible semantics without linking + hcomm: reverse endpoint EID and imported peer TPN are important. +- P2P validation must use CCU copy microcode. The old marker/IPC harness did not + prove the direct CCU data plane and has been removed from the maintained smoke. +- Process bring-up overrides for runtime task fields, submit-task arguments, + peer remote-XN/CKE bindings, and hcomm-trace resource remapping were removed + from maintained code. The formal path relies on exchanged TileXR resource + evidence. +- CCU direct work is independent of UDMA. Do not use UDMA code as a reference + for CCU resource installation or data movement. +- CANN 9.1 runtime declares `rtUbDevQueryInfo` in + `runtime/rts/rts_device.h`; the `tile-comm` target must include + `${ASCEND_HOME_PATH}/${ARCH}-linux/pkg_inc/runtime/runtime/`. + +## Maintained Validation + +Local source/unit checks: + +```bash +python -m unittest \ + tests.ccu.test_tilexr_ccu_direct_smoke_probe \ + tests.ccu.test_tilexr_ccu_public_comm_api \ + tests.ccu.test_tilexr_ccu_direct_smoke_runner \ + tests.ccu.test_tilexr_ccu_direct_orchestrator \ + tests.ccu.test_tilexr_ccu_memory_program \ + tests.ccu.test_tilexr_ccu_microcode +``` + +Broader CCU checks: + +```bash +python -m unittest discover tests/ccu +``` + +Remote build and no-hcomm guard: + +```bash +source scripts/common_env.sh +cmake --build build --target tile-comm -j2 +bash tests/ccu/check_tile_comm_no_hcomm_deps.sh build/src/comm/libtile-comm.so +``` + +Hardware smoke: + +```bash +source scripts/common_env.sh +env \ + TILEXR_CCU_SMOKE_DEVICES=0,1 \ + TILEXR_CCU_SMOKE_ALLOW_UNHEALTHY_NPU=1 \ + TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE=ra_ctx \ + TILEXR_CCU_DIRECT_EVENING_SMOKE_STAGES=prepare,submit,barrier,p2p \ + TILEXR_CCU_DIRECT_EVENING_PREPARE_PROFILES='ra_ctx_full:acl:full_repository:instruction_bytes:lower_layer_first' \ + TILEXR_CCU_EVENING_TOTAL_TIMEOUT=180 \ + bash tests/ccu/run_tilexr_ccu_direct_evening_smoke.sh +``` + +Expected final line: + +```text +tilexr_ccu_direct_evening_smoke finalStatus prepare=pass submit=pass barrier=pass p2p=pass +``` + +Expected P2P evidence in rank logs: + +```text +tilexr_ccu_direct_smoke p2pCcuCopy ... mismatches=0 ... passed=1 +decoded=TransRmtMemToLocMem +``` + +## Development Rules + +- Keep direct CCU code under `src/comm/ccu` and public API declarations under + `src/include/tilexr_api.h`. +- Keep hardware smoke scripts default-safe; they must require explicit env + opt-in before touching NPU devices. +- Do not reintroduce marker-based P2P success criteria. It does not prove CCU + transfer. +- Do not add runtime env overrides that mutate prepared task fields, prepared + task arguments, or peer binding proof. They make the smoke hard to interpret + and are not a production integration path. +- Keep verbose CCU instruction/resource traces behind env gates. They are useful + for hardware failures, but should not be required for normal validation. +- Prefer small unit tests for microcode, repository layout, resource allocation, + and submit-task packaging. Use hardware smoke only for the final device-plane + proof. diff --git a/src/comm/CMakeLists.txt b/src/comm/CMakeLists.txt index f9fcf19a..50519bf7 100644 --- a/src/comm/CMakeLists.txt +++ b/src/comm/CMakeLists.txt @@ -108,6 +108,10 @@ set(TILEXR_SOURCE_FILE tilexr_comm.cpp ccu/tilexr_ccu_abi_constants.h ccu/tilexr_ccu_barrier_program.h ccu/tilexr_ccu_barrier_program.cpp + ccu/tilexr_ccu_direct_orchestrator.h + ccu/tilexr_ccu_direct_orchestrator.cpp + ccu/tilexr_ccu_direct_runtime.h + ccu/tilexr_ccu_direct_runtime.cpp ccu/tilexr_ccu_driver_adapter.h ccu/tilexr_ccu_driver_adapter.cpp ccu/tilexr_ccu_hccp_types.h diff --git a/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp b/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp new file mode 100644 index 00000000..f3d8aef5 --- /dev/null +++ b/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp @@ -0,0 +1,1199 @@ +/* + * Copyright (c) 2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#include "ccu/tilexr_ccu_direct_orchestrator.h" + +#include "ccu/tilexr_ccu_runtime.h" + +#include +#include +#include +#include +#include + +namespace TileXR { +namespace { + +constexpr const char* TILEXR_CCU_DIRECT_KNOWN_MISSING_INSTALL_SURFACES = + "remote XN install provider is missing"; +constexpr uint16_t TILEXR_CCU_TRACE_LOAD_SQE_ARGS_TO_X_HEADER = 0x0001U; +constexpr uint16_t TILEXR_CCU_TRACE_LOAD_IMD_TO_GSA_HEADER = 0x0002U; +constexpr uint16_t TILEXR_CCU_TRACE_LOAD_IMD_TO_XN_HEADER = 0x0003U; +constexpr uint16_t TILEXR_CCU_TRACE_SET_CKE_HEADER = 0x0802U; +constexpr uint16_t TILEXR_CCU_TRACE_CLEAR_CKE_HEADER = 0x0804U; +constexpr uint16_t TILEXR_CCU_TRACE_TRANS_RMT_MEM_TO_LOC_MEM_HEADER = 0x1008U; +constexpr uint16_t TILEXR_CCU_TRACE_TRANS_LOC_MEM_TO_RMT_MEM_HEADER = 0x1009U; +constexpr uint16_t TILEXR_CCU_TRACE_SYNC_CKE_HEADER = 0x100bU; +constexpr uint16_t TILEXR_CCU_TRACE_SYNC_XN_HEADER = 0x100dU; +constexpr uint64_t TILEXR_CCU_PACKED_TOKEN_VALID_SHIFT = 52ULL; +constexpr uint64_t TILEXR_CCU_PACKED_TOKEN_ID_SHIFT = 32ULL; +constexpr uint64_t TILEXR_CCU_PACKED_TOKEN_ID_MASK = 0xfffffULL; +constexpr uint64_t TILEXR_CCU_PACKED_TOKEN_VALUE_MASK = 0xffffffffULL; +constexpr uint32_t TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT = 7U; +constexpr uint32_t TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT = 3U; +constexpr uint32_t TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_GSA_COUNT = 2U; + +void ResetReport(TileXRCcuDirectInstallReport* report) +{ + if (report != nullptr) { + *report = TileXRCcuDirectInstallReport{}; + } +} + +void ClearAttempt(TileXRCcuDirectInstallAttempt* attempt) +{ + if (attempt != nullptr) { + *attempt = TileXRCcuDirectInstallAttempt{}; + } +} + +int Fail(TileXRCcuDirectInstallAttempt* attempt, TileXRCcuDirectInstallReport* report, const std::string& message) +{ + ClearAttempt(attempt); + if (report != nullptr) { + report->message = message; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; +} + +bool HasRepositoryInstallInputs(const TileXRCcuDirectInstallOptions& options) +{ + return options.driverAdapter != nullptr && + options.repositoryMemoryOps.alloc != nullptr && + options.repositoryMemoryOps.copyHostToDevice != nullptr && + options.repositoryMemoryOps.free != nullptr; +} + +bool SyncXnMode(TileXRCcuBarrierMode mode) +{ + return mode == TileXRCcuBarrierMode::SyncXn || + mode == TileXRCcuBarrierMode::SyncXnPostOnly || + mode == TileXRCcuBarrierMode::SyncXnLoadPostOnly; +} + +bool HasNonZeroArgs(const TileXRCcuTaskWindow& window) +{ + for (uint64_t arg : window.args) { + if (arg != 0) { + return true; + } + } + return false; +} + +bool ContainsRange(uint16_t outerStart, uint16_t outerCount, uint16_t innerStart, uint32_t innerCount) +{ + if (outerCount == 0 || innerCount == 0) { + return false; + } + const uint32_t outerBegin = outerStart; + const uint32_t outerEnd = outerBegin + outerCount; + const uint32_t innerBegin = innerStart; + const uint32_t innerEnd = innerBegin + innerCount; + return innerBegin >= outerBegin && innerEnd <= outerEnd; +} + +bool RangesOverlap(uint16_t firstStart, uint32_t firstCount, uint16_t secondStart, uint32_t secondCount) +{ + const uint32_t firstEnd = static_cast(firstStart) + firstCount; + const uint32_t secondEnd = static_cast(secondStart) + secondCount; + return static_cast(firstStart) < secondEnd && static_cast(secondStart) < firstEnd; +} + +TileXRCcuRange MakeRange(uint8_t dieId, uint16_t startId, uint16_t count) +{ + TileXRCcuRange range; + range.dieId = dieId; + range.startId = startId; + range.num = count; + return range; +} + +uint64_t PackCcuSqeToken(uint32_t tokenId, uint32_t tokenValue, bool valid) +{ + const uint64_t validBits = valid ? 1ULL : 0ULL; + return (validBits << TILEXR_CCU_PACKED_TOKEN_VALID_SHIFT) | + ((static_cast(tokenId) & TILEXR_CCU_PACKED_TOKEN_ID_MASK) << TILEXR_CCU_PACKED_TOKEN_ID_SHIFT) | + (static_cast(tokenValue) & TILEXR_CCU_PACKED_TOKEN_VALUE_MASK); +} + +uint16_t ReadLe16(const uint8_t* raw, uint32_t offset) +{ + return static_cast(raw[offset]) | + static_cast(static_cast(raw[offset + 1U]) << 8U); +} + +uint64_t DecodeChannelRemoteCcuVa(const TileXRCcuChannelCtxDataV1& ctx) +{ + const uint16_t word28 = ReadLe16(ctx.raw, 28); + const uint16_t word34 = ReadLe16(ctx.raw, 34); + const uint64_t dstVa = + ((static_cast(word28) >> 8U) & 0xffULL) | + (static_cast(ReadLe16(ctx.raw, 30)) << 8U) | + (static_cast(ReadLe16(ctx.raw, 32)) << 24U) | + ((static_cast(word34) & 0x1ULL) << 40U); + return dstVa << TILEXR_CCU_REMOTE_CCU_VA_SHIFT; +} + +int PopulateHcommStyleSqeTaskArgs( + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (attempt == nullptr || + attempt->plan.taskWindows.size() < 2 || + !SyncXnMode(attempt->plan.barrierMode)) { + return TILEXR_SUCCESS; + } + + TileXRCcuTaskWindow& sqeLoadTask = attempt->plan.taskWindows[0]; + if (sqeLoadTask.argSize != TILEXR_CCU_SQE_ARGS_LEN || + sqeLoadTask.instCnt == 0 || + attempt->specInfo.resourceAddr == 0) { + return TILEXR_SUCCESS; + } + if (HasNonZeroArgs(sqeLoadTask)) { + return TILEXR_SUCCESS; + } + + sqeLoadTask.args.assign(TILEXR_CCU_SQE_ARGS_LEN, 0); + sqeLoadTask.args[0] = attempt->specInfo.resourceAddr; + sqeLoadTask.args[1] = attempt->specInfo.resourceAddr; + + if (!attempt->preparedLowerLayerPlan.msidTokens.empty()) { + const auto& token = attempt->preparedLowerLayerPlan.msidTokens[0]; + sqeLoadTask.args[2] = PackCcuSqeToken(token.tokenId, token.tokenValue, true); + } + for (const auto& channel : attempt->preparedLowerLayerPlan.channels) { + const uint64_t remoteCcuVa = DecodeChannelRemoteCcuVa(channel.ctx); + if (remoteCcuVa != 0) { + sqeLoadTask.args[3] = remoteCcuVa; + break; + } + } + + if (report != nullptr) { + report->message.clear(); + } + return TILEXR_SUCCESS; +} + +bool DirectTraceEnabled() +{ + const char* value = std::getenv("TILEXR_CCU_DIRECT_TRACE"); + return value != nullptr && value[0] != '\0' && value[0] != '0'; +} + +uint16_t TraceSlot(uint64_t word, uint32_t slot) +{ + return static_cast((word >> (slot * 16U)) & 0xffffU); +} + +uint16_t TraceRead16(const uint8_t* raw, uint32_t offset) +{ + return ReadLe16(raw, offset); +} + +uint64_t TraceReadDoorbellVa(const TileXRCcuLocalJettyCtxData& ctx) +{ + uint64_t value = 0; + for (uint32_t word = 0; word < 4U; ++word) { + value |= static_cast(TraceRead16(ctx.raw, word * 2U)) << (word * 16U); + } + return value; +} + +uint64_t TraceLoadImmediate(const TileXRCcuInstr& instr) +{ + return (instr.words[0] >> 32U) | ((instr.words[1] & 0xffffffffULL) << 32U); +} + +void TraceDecodedPfeCtx(size_t index, const TileXRCcuPfeInstall& pfe) +{ + const uint16_t word = TraceRead16(pfe.ctx.raw, 2); + const uint16_t jettyCountMinusOne = word & 0x7fU; + std::cerr << "TileXRDirectCcuTrace lowerLayerPfe[" << index << "]" + << " decoded=PfeCtx" + << " dieId=" << static_cast(pfe.dieId) + << " pfeOffset=" << pfe.pfeOffset + << " startTaJettyId=" << TraceRead16(pfe.ctx.raw, 0) + << " jettyCount=" << static_cast(jettyCountMinusOne) + 1U + << " jettyCountMinusOne=" << jettyCountMinusOne + << " startLocalJettyCtxId=" << ((word >> 7U) & 0x7fU) + << "\n"; +} + +void TraceDecodedLocalJettyCtx(size_t jettyIndex, size_t ctxIndex, const TileXRCcuLocalJettyCtxData& ctx) +{ + const uint16_t word8 = TraceRead16(ctx.raw, 8); + const uint16_t word10 = TraceRead16(ctx.raw, 10); + const uint16_t word14 = TraceRead16(ctx.raw, 14); + const uint16_t word22 = TraceRead16(ctx.raw, 22); + const uint16_t word24 = TraceRead16(ctx.raw, 24); + const uint32_t tokenId = + ((word8 >> 8U) & 0xffU) | + ((static_cast(word10) & 0xfffU) << 8U); + const uint32_t tokenValue = + ((word10 >> 12U) & 0xfU) | + (static_cast(TraceRead16(ctx.raw, 12)) << 4U) | + ((static_cast(word14) & 0xfffU) << 20U); + const uint32_t wqeBasicBlockShift = (word14 >> 12U) & 0xfU; + const uint32_t wqeBasicBlockCount = 1U << wqeBasicBlockShift; + const uint32_t wqeBasicBlockStartId = + ((word22 >> 12U) & 0xfU) | + ((static_cast(word24) & 0xffU) << 4U); + + std::cerr << "TileXRDirectCcuTrace lowerLayerJettyCtx[" << jettyIndex << "," << ctxIndex << "]" + << " decoded=LocalJettyCtx" + << " doorbellVa=" << std::hex << std::showbase << TraceReadDoorbellVa(ctx) + << " doorbellTokenId=" << tokenId + << " doorbellTokenValue=" << tokenValue + << std::dec << std::noshowbase + << " pfeId=" << (word8 & 0xfU) + << " ioDieId=" << ((word8 >> 4U) & 0x1U) + << " doorbellAddrType=" << ((word8 >> 5U) & 0x1U) + << " tokenValueValid=" << ((word8 >> 6U) & 0x1U) + << " sqeBasicBlockLeftShifts=" << wqeBasicBlockShift + << " wqeBasicBlockCount=" << wqeBasicBlockCount + << " inferredSqDepth=" << (wqeBasicBlockCount / 4U) + << " wqeBasicBlockStartId=" << wqeBasicBlockStartId + << " pi=" << TraceRead16(ctx.raw, 16) + << " ci=" << TraceRead16(ctx.raw, 18) + << " maxCi=" << TraceRead16(ctx.raw, 20) + << " oooCqeCnt=" << (word22 & 0xfffU) + << " doorbellSendState=" << ((word24 >> 8U) & 0x3U) + << "\n"; +} + +void TraceDecodedChannelCtxV1(size_t index, const TileXRCcuChannelInstall& channel) +{ + const uint16_t word16 = TraceRead16(channel.ctx.raw, 16); + const uint16_t word18 = TraceRead16(channel.ctx.raw, 18); + const uint16_t word20 = TraceRead16(channel.ctx.raw, 20); + const uint16_t word22 = TraceRead16(channel.ctx.raw, 22); + const uint16_t word24 = TraceRead16(channel.ctx.raw, 24); + const uint16_t word28 = TraceRead16(channel.ctx.raw, 28); + const uint16_t word34 = TraceRead16(channel.ctx.raw, 34); + + const uint32_t tpn = word16 | ((static_cast(word18) & 0xffU) << 16U); + const uint16_t startJettyId = + ((word18 >> 12U) & 0xfU) | + static_cast((word20 & 0xfffU) << 4U); + const uint16_t jettyCountMinusOne = + ((word20 >> 12U) & 0xfU) | + static_cast((word22 & 0x7U) << 4U); + const uint32_t tokenId = + ((static_cast(word22) >> 4U) & 0xfffU) | + ((static_cast(word24) & 0xffU) << 12U); + const uint32_t tokenValue = + ((static_cast(word24) >> 8U) & 0xffU) | + (static_cast(TraceRead16(channel.ctx.raw, 26)) << 8U) | + ((static_cast(word28) & 0xffU) << 24U); + const uint64_t dstVa = + ((static_cast(word28) >> 8U) & 0xffULL) | + (static_cast(TraceRead16(channel.ctx.raw, 30)) << 8U) | + (static_cast(TraceRead16(channel.ctx.raw, 32)) << 24U) | + ((static_cast(word34) & 0x1ULL) << 40U); + const uint64_t remoteCcuVa = dstVa << TILEXR_CCU_REMOTE_CCU_VA_SHIFT; + + std::cerr << "TileXRDirectCcuTrace lowerLayerChannel[" << index << "]" + << " decoded=ChannelCtxV1" + << " dieId=" << static_cast(channel.dieId) + << " channelId=" << channel.channelId + << " tpn=" << std::hex << std::showbase << tpn + << " memoryTokenId=" << tokenId + << " memoryTokenValue=" << tokenValue + << " dstVaShifted=" << dstVa + << " remoteCcuVa=" << remoteCcuVa + << std::dec << std::noshowbase + << " sourcePfeId=" << ((word18 >> 8U) & 0xfU) + << " startTaJettyId=" << startJettyId + << " jettyCount=" << static_cast(jettyCountMinusOne) + 1U + << " jettyCountMinusOne=" << jettyCountMinusOne + << " ioDieId=" << ((word22 >> 3U) & 0x1U) + << " tokenValueValid=" << ((word34 >> 1U) & 0x1U) + << " remoteEid="; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + std::cerr << std::hex << std::setw(2) << std::setfill('0') << std::noshowbase + << static_cast(channel.ctx.raw[i]); + } + std::cerr << std::dec << std::setfill(' ') << "\n"; +} + +void TraceDecodedInstr(const char* label, size_t index, const TileXRCcuInstr& instr) +{ + const uint16_t opcode = TraceSlot(instr.words[0], 0); + std::cerr << "TileXRDirectCcuTrace " << label << "[" << index << "] "; + switch (opcode) { + case TILEXR_CCU_TRACE_LOAD_SQE_ARGS_TO_X_HEADER: + std::cerr << "decoded=LoadSqeArgsToX" + << " xnId=" << TraceSlot(instr.words[0], 1) + << " sqeArgId=" << TraceSlot(instr.words[0], 2); + break; + case TILEXR_CCU_TRACE_LOAD_IMD_TO_XN_HEADER: + std::cerr << "decoded=LoadImdToXn" + << " xnId=" << TraceSlot(instr.words[0], 1) + << " immediate=" << std::hex << std::showbase << TraceLoadImmediate(instr) + << std::dec << std::noshowbase + << " secFlag=" << TraceSlot(instr.words[1], 2); + break; + case TILEXR_CCU_TRACE_LOAD_IMD_TO_GSA_HEADER: + std::cerr << "decoded=LoadImdToGSA" + << " gsaId=" << TraceSlot(instr.words[0], 1) + << " immediate=" << std::hex << std::showbase << TraceLoadImmediate(instr) + << std::dec << std::noshowbase; + break; + case TILEXR_CCU_TRACE_TRANS_RMT_MEM_TO_LOC_MEM_HEADER: { + const uint16_t control = TraceSlot(instr.words[1], 3); + const uint16_t flags = TraceSlot(instr.words[2], 3); + std::cerr << "decoded=TransRmtMemToLocMem" + << " remoteGsa=" << TraceSlot(instr.words[0], 3) + << " remoteXn=" << TraceSlot(instr.words[1], 0) + << " localGsa=" << TraceSlot(instr.words[0], 1) + << " localXn=" << TraceSlot(instr.words[0], 2) + << " lengthXn=" << TraceSlot(instr.words[1], 1) + << " channelId=" << TraceSlot(instr.words[1], 2) + << " clearType=" << (flags & 0x1U) + << " lengthEn=" << ((flags >> 1U) & 0x1U) + << " reduceEn=" << ((flags >> 2U) & 0x1U) + << " reduceDataType=" << ((control >> 8U) & 0xfU) + << " reduceOpCode=" << ((control >> 12U) & 0xfU) + << " setCkeId=" << TraceSlot(instr.words[3], 0) + << " setCkeMask=" << TraceSlot(instr.words[3], 1) + << " waitCkeId=" << TraceSlot(instr.words[3], 2) + << " waitCkeMask=" << TraceSlot(instr.words[3], 3); + break; + } + case TILEXR_CCU_TRACE_TRANS_LOC_MEM_TO_RMT_MEM_HEADER: { + const uint16_t control = TraceSlot(instr.words[1], 3); + const uint16_t flags = TraceSlot(instr.words[2], 3); + std::cerr << "decoded=TransLocMemToRmtMem" + << " localGsa=" << TraceSlot(instr.words[0], 3) + << " localXn=" << TraceSlot(instr.words[1], 0) + << " remoteGsa=" << TraceSlot(instr.words[0], 1) + << " remoteXn=" << TraceSlot(instr.words[0], 2) + << " lengthXn=" << TraceSlot(instr.words[1], 1) + << " channelId=" << TraceSlot(instr.words[1], 2) + << " clearType=" << (flags & 0x1U) + << " lengthEn=" << ((flags >> 1U) & 0x1U) + << " reduceEn=" << ((flags >> 2U) & 0x1U) + << " reduceDataType=" << ((control >> 8U) & 0xfU) + << " reduceOpCode=" << ((control >> 12U) & 0xfU) + << " setCkeId=" << TraceSlot(instr.words[3], 0) + << " setCkeMask=" << TraceSlot(instr.words[3], 1) + << " waitCkeId=" << TraceSlot(instr.words[3], 2) + << " waitCkeMask=" << TraceSlot(instr.words[3], 3); + break; + } + case TILEXR_CCU_TRACE_SYNC_XN_HEADER: + std::cerr << "decoded=SyncXn" + << " remoteXn=" << TraceSlot(instr.words[0], 1) + << " localXn=" << TraceSlot(instr.words[0], 2) + << " channelId=" << TraceSlot(instr.words[1], 0) + << " notifyCke=" << TraceSlot(instr.words[1], 1) + << " notifyMask=" << TraceSlot(instr.words[1], 2) + << " traceFlag=" << std::hex << std::showbase << instr.words[2] + << std::dec << std::noshowbase + << " setCkeId=" << TraceSlot(instr.words[3], 0) + << " setCkeMask=" << TraceSlot(instr.words[3], 1) + << " waitCkeId=" << TraceSlot(instr.words[3], 2) + << " waitCkeMask=" << TraceSlot(instr.words[3], 3); + break; + case TILEXR_CCU_TRACE_SYNC_CKE_HEADER: + std::cerr << "decoded=SyncCke" + << " remoteCke=" << TraceSlot(instr.words[0], 1) + << " localCke=" << TraceSlot(instr.words[0], 2) + << " localCkeMask=" << TraceSlot(instr.words[0], 3) + << " channelId=" << TraceSlot(instr.words[1], 0) + << " clearType=" << TraceSlot(instr.words[2], 3) + << " setCkeId=" << TraceSlot(instr.words[3], 0) + << " setCkeMask=" << TraceSlot(instr.words[3], 1) + << " waitCkeId=" << TraceSlot(instr.words[3], 2) + << " waitCkeMask=" << TraceSlot(instr.words[3], 3); + break; + case TILEXR_CCU_TRACE_SET_CKE_HEADER: + std::cerr << "decoded=SetCke" + << " clearType=" << TraceSlot(instr.words[0], 1) + << " ckeId=" << TraceSlot(instr.words[0], 2) + << " mask=" << TraceSlot(instr.words[0], 3) + << " waitCkeId=" << TraceSlot(instr.words[1], 0) + << " waitMask=" << TraceSlot(instr.words[1], 1); + break; + case TILEXR_CCU_TRACE_CLEAR_CKE_HEADER: + std::cerr << "decoded=ClearCke" + << " clearType=" << TraceSlot(instr.words[0], 1) + << " ckeId=" << TraceSlot(instr.words[0], 2) + << " mask=" << TraceSlot(instr.words[0], 3) + << " waitCkeId=" << TraceSlot(instr.words[1], 0) + << " waitMask=" << TraceSlot(instr.words[1], 1); + break; + default: + std::cerr << "decoded=Unknown opcode=" << std::hex << std::showbase << opcode + << std::dec << std::noshowbase; + break; + } + std::cerr << "\n"; +} + +void TraceInstr(const char* label, size_t index, const TileXRCcuInstr& instr) +{ + std::cerr << "TileXRDirectCcuTrace " << label << "[" << index << "] words=" + << std::hex << std::showbase + << instr.words[0] << "," << instr.words[1] << "," + << instr.words[2] << "," << instr.words[3] + << std::dec << std::noshowbase << "\n"; + TraceDecodedInstr(label, index, instr); +} + +void TraceDirectInstallAttempt(const TileXRCcuDirectInstallAttempt& attempt) +{ + if (!DirectTraceEnabled()) { + return; + } + + std::cerr << "TileXRDirectCcuTrace begin" + << " missionId=" << static_cast(attempt.plan.mission.missionId) + << " missionKey=" << std::hex << std::showbase << attempt.plan.mission.key + << std::dec << std::noshowbase + << " dieId=" << static_cast(attempt.plan.mission.dieId) + << "\n"; + for (size_t i = 0; i < attempt.plan.syncResources.size(); ++i) { + const auto& resource = attempt.plan.syncResources[i]; + std::cerr << "TileXRDirectCcuTrace syncResource[" << i << "]" + << " dieId=" << static_cast(resource.dieId) + << " channelId=" << resource.channelId + << " localXn=" << resource.localXn + << " remoteXn=" << resource.remoteXn + << " notifyCke=" << resource.notifyCke + << " localWaitCke=" << resource.localWaitCke + << " localWaitMask=" << resource.localWaitMask + << " remoteNotifyMask=" << resource.remoteNotifyMask + << " sourceCke=" << resource.sourceCke + << " sourceCkeMask=" << resource.sourceCkeMask + << "\n"; + } + for (size_t i = 0; i < attempt.preparedLowerLayerPlan.remoteXnBindings.size(); ++i) { + const auto& proof = attempt.preparedLowerLayerPlan.remoteXnBindings[i]; + std::cerr << "TileXRDirectCcuTrace remoteXnBinding[" << i << "]" + << " dieId=" << static_cast(proof.dieId) + << " channelId=" << proof.channelId + << " localXn=" << proof.localXn + << " remoteXn=" << proof.remoteXn + << " notifyCke=" << proof.notifyCke + << " localWaitCke=" << proof.localWaitCke + << " peerRank=" << proof.peerRank + << " peerExchangeObserved=" << (proof.peerExchangeObserved ? 1 : 0) + << " endpointRouteVerified=" << (proof.endpointRouteVerified ? 1 : 0) + << " channelResourceOwnerVerified=" << (proof.channelResourceOwnerVerified ? 1 : 0) + << " transportResourceExchangeVerified=" << (proof.transportResourceExchangeVerified ? 1 : 0) + << "\n"; + } + for (size_t i = 0; i < attempt.preparedLowerLayerPlan.pfes.size(); ++i) { + const auto& pfe = attempt.preparedLowerLayerPlan.pfes[i]; + std::cerr << "TileXRDirectCcuTrace lowerLayerPfe[" << i << "]" + << " dieId=" << static_cast(pfe.dieId) + << " pfeOffset=" << pfe.pfeOffset + << " ctx="; + for (uint8_t byte : pfe.ctx.raw) { + std::cerr << std::hex << std::setw(2) << std::setfill('0') + << static_cast(byte); + } + std::cerr << std::dec << std::setfill(' ') << "\n"; + TraceDecodedPfeCtx(i, pfe); + } + for (size_t i = 0; i < attempt.preparedLowerLayerPlan.channels.size(); ++i) { + const auto& channel = attempt.preparedLowerLayerPlan.channels[i]; + std::cerr << "TileXRDirectCcuTrace lowerLayerChannel[" << i << "]" + << " dieId=" << static_cast(channel.dieId) + << " channelId=" << channel.channelId + << " ctx="; + for (uint8_t byte : channel.ctx.raw) { + std::cerr << std::hex << std::setw(2) << std::setfill('0') + << static_cast(byte); + } + std::cerr << std::dec << std::setfill(' ') << "\n"; + TraceDecodedChannelCtxV1(i, channel); + } + for (size_t i = 0; i < attempt.preparedLowerLayerPlan.jettys.size(); ++i) { + const auto& jetty = attempt.preparedLowerLayerPlan.jettys[i]; + std::cerr << "TileXRDirectCcuTrace lowerLayerJetty[" << i << "]" + << " dieId=" << static_cast(jetty.dieId) + << " startJettyCtxId=" << jetty.startJettyCtxId + << " ctxCount=" << jetty.ctxs.size() + << "\n"; + for (size_t j = 0; j < jetty.ctxs.size(); ++j) { + std::cerr << "TileXRDirectCcuTrace lowerLayerJettyCtx[" << i << "," << j << "] ctx="; + for (uint8_t byte : jetty.ctxs[j].raw) { + std::cerr << std::hex << std::setw(2) << std::setfill('0') + << static_cast(byte); + } + std::cerr << std::dec << std::setfill(' ') << "\n"; + TraceDecodedLocalJettyCtx(i, j, jetty.ctxs[j]); + } + } + for (size_t i = 0; i < attempt.plan.taskWindows.size(); ++i) { + const auto& window = attempt.plan.taskWindows[i]; + std::cerr << "TileXRDirectCcuTrace taskWindow[" << i << "]" + << " dieId=" << static_cast(window.dieId) + << " instStartId=" << window.instStartId + << " instCnt=" << window.instCnt + << " argSize=" << window.argSize + << " args="; + for (size_t arg = 0; arg < window.args.size(); ++arg) { + if (arg != 0) { + std::cerr << ","; + } + std::cerr << std::hex << std::showbase << window.args[arg] + << std::dec << std::noshowbase; + } + std::cerr << "\n"; + } + const auto& tracedTasks = attempt.submitTasks.empty() ? attempt.package.tasks : attempt.submitTasks; + for (size_t i = 0; i < tracedTasks.size(); ++i) { + const auto& task = tracedTasks[i]; + std::cerr << "TileXRDirectCcuTrace task[" << i << "]" + << " dieId=" << static_cast(task.dieId) + << " missionId=" << static_cast(task.missionId) + << " timeout=" << task.timeout + << " instStartId=" << task.instStartId + << " instCnt=" << task.instCnt + << " argSize=" << task.argSize + << " key=" << std::hex << std::showbase << task.key + << std::dec << std::noshowbase + << " args="; + for (uint32_t arg = 0; arg < TILEXR_CCU_SQE_ARGS_LEN; ++arg) { + if (arg != 0) { + std::cerr << ","; + } + std::cerr << std::hex << std::showbase << task.args[arg] + << std::dec << std::noshowbase; + } + std::cerr << "\n"; + } + for (size_t i = 0; i < attempt.package.program.sqeLoad.size(); ++i) { + TraceInstr("program.sqeLoad", i, attempt.package.program.sqeLoad[i]); + } + for (size_t i = 0; i < attempt.package.program.sync.size(); ++i) { + TraceInstr("program.sync", i, attempt.package.program.sync[i]); + } + std::cerr << "TileXRDirectCcuTrace end\n"; +} + +void ApplySplitCkeOptions( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuResourceSpec* resourceSpec) +{ + if (resourceSpec == nullptr) { + return; + } + if (options.localWaitCkeStartId != 0 || options.localWaitCkeCount != 0) { + resourceSpec->localWaitCkeStartId = + options.localWaitCkeStartId == 0 ? options.ckeStartId : options.localWaitCkeStartId; + resourceSpec->localWaitCkeCount = + options.localWaitCkeCount == 0 ? resourceSpec->ckeCount : options.localWaitCkeCount; + } + if (options.remoteNotifyCkeStartId != 0 || options.remoteNotifyCkeCount != 0) { + resourceSpec->remoteNotifyCkeStartId = + options.remoteNotifyCkeStartId == 0 ? options.ckeStartId : options.remoteNotifyCkeStartId; + resourceSpec->remoteNotifyCkeCount = + options.remoteNotifyCkeCount == 0 ? resourceSpec->ckeCount : options.remoteNotifyCkeCount; + } +} + +void ApplyRemoteXnOptions( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuResourceSpec* resourceSpec) +{ + if (resourceSpec == nullptr) { + return; + } + if (options.remoteXnStartId != 0 || options.remoteXnCount != 0) { + resourceSpec->remoteXnStartId = options.remoteXnStartId; + resourceSpec->remoteXnCount = + options.remoteXnCount == 0 ? resourceSpec->xnCount : options.remoteXnCount; + } +} + +int PrepareLowerLayerPlanIfNeeded( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (options.prepareLowerLayerPlan == nullptr) { + return TILEXR_SUCCESS; + } + if (attempt == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + TileXRCcuLowerLayerInstallPlan plan; + TileXRCcuLowerLayerPlanBuilderReport planReport; + const int ret = options.prepareLowerLayerPlan( + attempt->allocation, + &plan, + &planReport, + options.lowerLayerPlanUserData); + attempt->lowerLayerPlanReport = planReport; + if (ret != TILEXR_SUCCESS) { + if (report != nullptr) { + report->message = planReport.message.empty() ? + "failed to prepare direct CCU lower-layer install plan" : + planReport.message; + } + return ret; + } + attempt->preparedLowerLayerPlan = plan; + return TILEXR_SUCCESS; +} + +int ReconcileProducerPlanWithLowerLayerProof( + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (attempt == nullptr || attempt->preparedLowerLayerPlan.remoteXnBindings.empty()) { + return TILEXR_SUCCESS; + } + if (attempt->preparedLowerLayerPlan.remoteXnBindings.size() != attempt->plan.syncResources.size()) { + if (report != nullptr) { + report->message = "lower-layer remote XN proof count does not match producer sync resources"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + for (auto& resource : attempt->plan.syncResources) { + bool matched = false; + for (const auto& proof : attempt->preparedLowerLayerPlan.remoteXnBindings) { + if (proof.dieId != resource.dieId || + proof.channelId != resource.channelId || + proof.localXn != resource.localXn || + !proof.peerExchangeObserved) { + continue; + } + resource.remoteXn = proof.remoteXn; + resource.notifyCke = proof.notifyCke; + if (proof.localWaitCke != 0) { + resource.localWaitCke = proof.localWaitCke; + } + matched = true; + break; + } + if (!matched) { + if (report != nullptr) { + report->message = "lower-layer remote XN proof does not cover producer sync resource"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + } + return TILEXR_SUCCESS; +} + +int ConfigureDirectMemoryCopyResources( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (attempt == nullptr || attempt->plan.syncResources.size() != 1 || attempt->plan.taskWindows.size() != 1) { + if (report != nullptr) { + report->message = "memory copy direct CCU plan requires one sync resource and one task"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + if (options.gsaStartId == 0 || attempt->resourceSpec.gsaCount < TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_GSA_COUNT) { + if (report != nullptr) { + report->message = "memory copy direct CCU requires a kernel-local GSA resource window"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + const uint8_t dieId = attempt->specInfo.dieId; + const uint16_t localXnStart = attempt->allocation.localXn.startId; + if (!ContainsRange( + attempt->resourceSpec.xnStartId, + attempt->resourceSpec.xnCount, + localXnStart, + TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT)) { + if (report != nullptr) { + report->message = "memory copy direct CCU local XN window is too small"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + uint16_t remoteXnStart = attempt->allocation.remoteXn.startId; + if (RangesOverlap( + localXnStart, + TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT, + remoteXnStart, + 1U)) { + remoteXnStart = static_cast(localXnStart + TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT); + } + if (!ContainsRange(attempt->resourceSpec.xnStartId, attempt->resourceSpec.xnCount, remoteXnStart, 1U)) { + if (report != nullptr) { + report->message = "memory copy direct CCU remote XN window is outside the XN resource range"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + attempt->allocation.localXn = + MakeRange(dieId, localXnStart, static_cast(TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT)); + attempt->allocation.localGsa = + MakeRange(dieId, options.gsaStartId, static_cast(TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_GSA_COUNT)); + attempt->allocation.remoteXn = MakeRange(dieId, remoteXnStart, 1U); + attempt->plan.kernelLocalXn = attempt->allocation.localXn; + attempt->plan.kernelLocalGsa = attempt->allocation.localGsa; + attempt->plan.syncResources[0].remoteXn = remoteXnStart; + attempt->plan.taskWindows[0].instCnt = + static_cast(std::max( + attempt->plan.taskWindows[0].instCnt, + TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT)); + return TILEXR_SUCCESS; +} + +int BuildDirectMemoryCopyLaunchPackage( + const TileXRCcuDirectMemoryCopySpec& memoryCopy, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (attempt == nullptr || attempt->plan.syncResources.empty()) { + if (report != nullptr) { + report->message = "missing direct CCU memory copy producer resources"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + const TileXRCcuSyncResource& resource = attempt->plan.syncResources[0]; + if (attempt->plan.kernelLocalGsa.num < TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_GSA_COUNT || + attempt->plan.kernelLocalXn.num < TILEXR_CCU_DIRECT_MEMORY_COPY_LOCAL_XN_COUNT) { + if (report != nullptr) { + report->message = "missing direct CCU memory copy GSA/XN resources"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + TileXRCcuMemoryCopySpec copySpec; + copySpec.direction = memoryCopy.direction; + copySpec.localGsa = attempt->plan.kernelLocalGsa.startId; + copySpec.remoteGsa = static_cast(attempt->plan.kernelLocalGsa.startId + 1U); + copySpec.localXn = attempt->plan.kernelLocalXn.startId; + copySpec.remoteXn = static_cast(attempt->plan.kernelLocalXn.startId + 1U); + copySpec.lengthXn = static_cast(attempt->plan.kernelLocalXn.startId + 2U); + copySpec.localAddr = memoryCopy.localAddr; + copySpec.localToken = memoryCopy.localToken; + copySpec.remoteAddr = memoryCopy.remoteAddr; + copySpec.remoteToken = memoryCopy.remoteToken; + copySpec.lengthBytes = memoryCopy.lengthBytes; + copySpec.channelId = resource.channelId; + copySpec.completionCke = resource.localWaitCke == 0 ? resource.notifyCke : resource.localWaitCke; + copySpec.completionMask = resource.localWaitMask == 0 ? 1U : resource.localWaitMask; + + TileXRCcuProgram program; + TileXRCcuMemoryProgramReport memoryReport; + if (TileXRCcuBuildMemoryCopyProgram(copySpec, &program.sync, &memoryReport) != TILEXR_SUCCESS) { + if (report != nullptr) { + report->message = memoryReport.message; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + TileXRCcuRepositoryImage repository; + TileXRCcuRepositoryReport repositoryReport; + if (TileXRCcuBuildRepositoryImage(attempt->plan, program, &repository, &repositoryReport) != TILEXR_SUCCESS) { + if (report != nullptr) { + report->message = repositoryReport.message; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + std::vector tasks; + TileXRCcuProducerPlanReport planReport; + if (TileXRCcuBuildTasks(attempt->plan, &tasks, &planReport) != TILEXR_SUCCESS) { + if (report != nullptr) { + report->message = planReport.message; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + attempt->package.plan = attempt->plan; + attempt->package.program = program; + attempt->package.repository = repository; + attempt->package.tasks = tasks; + attempt->package.installScope = TileXRCcuLaunchInstallScope {}; + attempt->package.requiresHardwareInstall = true; + return TILEXR_SUCCESS; +} + +void FillReportFromAttempt(const TileXRCcuDirectInstallAttempt& attempt, TileXRCcuDirectInstallReport* report) +{ + if (report == nullptr) { + return; + } + report->pipelineBuilt = true; + report->installAttempted = attempt.installReport.installAttempted; + report->installSucceeded = attempt.installReport.installSucceeded; + report->submitReady = attempt.providerReport.submitReady; + report->requiredInstallSurfaceCount = attempt.installReport.requiredInstallSurfaceCount; + report->publicVerifiedInstallSurfaceCount = attempt.installReport.publicVerifiedInstallSurfaceCount; + report->missingInstallSurfaceCount = attempt.installReport.missingInstallSurfaceCount; + report->taskCount = static_cast(attempt.package.tasks.size()); + report->submitTaskCount = static_cast(attempt.submitTasks.size()); +} + +void ResetSubmitReport(TileXRCcuDirectSubmitReport* report) +{ + if (report != nullptr) { + *report = TileXRCcuDirectSubmitReport{}; + } +} + +void TraceFinalRuntimeTask(size_t taskIndex, const TileXRCcuTask& task) +{ + if (!DirectTraceEnabled()) { + return; + } + std::cerr << "TileXRDirectCcuTrace finalRuntimeTask[" << taskIndex << "]" + << " dieId=" << static_cast(task.dieId) + << " missionId=" << static_cast(task.missionId) + << " timeout=" << task.timeout + << " instStartId=" << task.instStartId + << " instCnt=" << task.instCnt + << " key=0x" << std::hex << std::nouppercase << task.key + << std::dec + << " argSize=" << task.argSize; + for (uint32_t arg = 0; arg < TILEXR_CCU_SQE_ARGS_LEN; ++arg) { + std::cerr << " args[" << arg << "]=0x" + << std::hex << std::nouppercase << task.args[arg] + << std::dec; + } + std::cerr << "\n"; +} + +std::string FormatSubmitTaskFailure( + size_t taskIndex, + int ret, + const TileXRCcuTask& task, + const TileXRCcuRuntimeSubmitReport* runtimeReport) +{ + const TileXRCcuTask& diagnosticTask = + (runtimeReport != nullptr && runtimeReport->finalTaskCaptured) ? + runtimeReport->finalTask : + task; + std::ostringstream oss; + oss << "direct CCU submit failed task=" << taskIndex + << " ret=" << ret; + if (runtimeReport != nullptr && runtimeReport->runtimeLaunchAttempted) { + oss << " rtRet=" << runtimeReport->runtimeRet; + } + oss << " dieId=" << static_cast(diagnosticTask.dieId) + << " missionId=" << static_cast(diagnosticTask.missionId) + << " timeout=" << diagnosticTask.timeout + << " instStartId=" << diagnosticTask.instStartId + << " instCnt=" << diagnosticTask.instCnt + << " key=0x" << std::hex << std::nouppercase << diagnosticTask.key + << std::dec + << " argSize=" << diagnosticTask.argSize; + for (uint32_t arg = 0; arg < TILEXR_CCU_SQE_ARGS_LEN; ++arg) { + oss << " args[" << arg << "]=0x" + << std::hex << std::nouppercase << diagnosticTask.args[arg] + << std::dec; + } + return oss.str(); +} + +int ReturnWithAttemptStatus( + int ret, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (attempt != nullptr) { + FillReportFromAttempt(*attempt, report); + if (report != nullptr) { + if (attempt->providerReport.submitReady) { + report->message = attempt->providerReport.message; + } else if (!attempt->installReport.message.empty()) { + report->message = attempt->installReport.message; + } else { + report->message = attempt->providerReport.message.empty() ? + TILEXR_CCU_DIRECT_KNOWN_MISSING_INSTALL_SURFACES : + attempt->providerReport.message; + } + } + } + return ret; +} + +} // namespace + +int TileXRCcuSubmitPreparedTasks( + const std::vector& submitTasks, + void* stream, + TileXRCcuTaskSubmitFn submitFn, + void* submitUserData, + TileXRCcuDirectSubmitReport* report) +{ + ResetSubmitReport(report); + if (submitTasks.empty()) { + if (report != nullptr) { + report->message = "missing prepared direct CCU submit tasks"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + if (stream == nullptr) { + if (report != nullptr) { + report->taskCount = static_cast(submitTasks.size()); + report->message = "missing runtime stream for direct CCU submit"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + uint32_t submitted = 0; + for (size_t i = 0; i < submitTasks.size(); ++i) { + TileXRCcuTask taskForSubmit = submitTasks[i]; + + TileXRCcuRuntimeSubmitReport runtimeReport; + const bool useDefaultSubmit = submitFn == nullptr; + const int ret = useDefaultSubmit ? + TileXRCcuSubmitTaskWithReport(taskForSubmit, stream, &runtimeReport) : + submitFn(taskForSubmit, stream, submitUserData); + if (useDefaultSubmit && runtimeReport.finalTaskCaptured) { + TraceFinalRuntimeTask(i, runtimeReport.finalTask); + } + if (ret != TILEXR_SUCCESS) { + if (report != nullptr) { + report->taskCount = static_cast(submitTasks.size()); + report->submittedTaskCount = submitted; + report->message = FormatSubmitTaskFailure( + i, + ret, + taskForSubmit, + useDefaultSubmit ? &runtimeReport : nullptr); + } + return ret; + } + ++submitted; + } + + if (report != nullptr) { + report->submitted = true; + report->taskCount = static_cast(submitTasks.size()); + report->submittedTaskCount = submitted; + report->message = "direct CCU prepared tasks submitted"; + } + return TILEXR_SUCCESS; +} + +int RunDirectInstallAttemptImpl( + const TileXRCcuDirectInstallOptions& options, + const TileXRCcuDirectMemoryCopySpec* memoryCopy, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + ResetReport(report); + if (attempt == nullptr) { + return Fail(nullptr, report, "missing output direct CCU install attempt"); + } + ClearAttempt(attempt); + + if (options.basicInfo == nullptr) { + return Fail(attempt, report, "missing direct CCU basic info"); + } + if (options.provider.empty()) { + return Fail(attempt, report, "missing direct CCU install provider"); + } + if (!options.offlineOnly && !HasRepositoryInstallInputs(options)) { + return Fail(attempt, report, "missing direct CCU repository install inputs"); + } + attempt->repositoryMemoryOps = options.repositoryMemoryOps; + attempt->repositoryMemoryUserData = options.repositoryMemoryUserData; + + TileXRCcuSpecsReport specsReport; + int ret = TileXRCcuDecodeBasicInfo(*options.basicInfo, &attempt->specInfo, &specsReport); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, specsReport.message); + } + + ret = TileXRCcuBuildResourceSpec( + attempt->specInfo, + options.missionStartId, + options.instructionStartId, + options.xnStartId, + options.ckeStartId, + options.channelStartId, + &attempt->resourceSpec, + &specsReport, + options.gsaStartId); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, specsReport.message); + } + attempt->resourceSpec.missionInstructionStartId = options.missionInstructionStartId; + ApplyRemoteXnOptions(options, &attempt->resourceSpec); + ApplySplitCkeOptions(options, &attempt->resourceSpec); + + attempt->resourceRequest.sqeArgCount = memoryCopy == nullptr ? options.sqeArgCount : 0U; + attempt->resourceRequest.syncResourceCount = memoryCopy == nullptr ? options.syncResourceCount : 1U; + attempt->resourceRequest.syncInstructionCount = memoryCopy == nullptr ? + options.syncInstructionCount : + std::max(options.syncInstructionCount, TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT); + attempt->resourceRequest.bindingsPerSyncResource = options.bindingsPerSyncResource; + attempt->resourceRequest.barrierMode = options.barrierMode; + + TileXRCcuResourceAllocator allocator; + if (allocator.Init(attempt->resourceSpec) != TILEXR_SUCCESS) { + return Fail(attempt, report, "failed to initialize direct CCU resource allocator"); + } + + TileXRCcuResourceAllocatorReport allocatorReport; + ret = allocator.Allocate( + attempt->resourceRequest, + &attempt->plan, + &attempt->allocation, + &allocatorReport); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, allocatorReport.message); + } + + if (memoryCopy != nullptr) { + ret = ConfigureDirectMemoryCopyResources(options, attempt, report); + if (ret != TILEXR_SUCCESS) { + return Fail( + attempt, + report, + report == nullptr || report->message.empty() ? + "failed to configure direct CCU memory copy resources" : + report->message); + } + } + + ret = PrepareLowerLayerPlanIfNeeded(options, attempt, report); + if (ret != TILEXR_SUCCESS) { + return Fail( + attempt, + report, + report == nullptr || report->message.empty() ? + "failed to prepare direct CCU lower-layer install plan" : + report->message); + } + + ret = ReconcileProducerPlanWithLowerLayerProof(attempt, report); + if (ret != TILEXR_SUCCESS) { + return Fail( + attempt, + report, + report == nullptr || report->message.empty() ? + "failed to reconcile direct CCU lower-layer peer resources" : + report->message); + } + if (memoryCopy == nullptr) { + ret = PopulateHcommStyleSqeTaskArgs(attempt, report); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, "failed to populate direct CCU SQE task arguments"); + } + } + + TileXRCcuLaunchPackageReport packageReport; + ret = memoryCopy == nullptr ? + TileXRCcuBuildLaunchPackage(attempt->plan, &attempt->package, &packageReport) : + BuildDirectMemoryCopyLaunchPackage(*memoryCopy, attempt, report); + if (ret != TILEXR_SUCCESS) { + return Fail( + attempt, + report, + memoryCopy == nullptr ? packageReport.message : + (report == nullptr || report->message.empty() ? + "failed to build direct CCU memory copy launch package" : + report->message)); + } + + TraceDirectInstallAttempt(*attempt); + + ret = TileXRCcuBindLaunchPackageInstallScope( + &attempt->package, + options.deviceId, + options.rank, + options.provider); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, "failed to bind direct CCU launch install scope"); + } + + TileXRCcuInstallManifestReport manifestReport; + ret = TileXRCcuBuildInstallManifest(attempt->package, &attempt->manifest, &manifestReport); + if (ret != TILEXR_SUCCESS) { + return Fail(attempt, report, manifestReport.message); + } + + TileXRCcuInstallRequest installRequest; + installRequest.package = &attempt->package; + installRequest.manifest = &attempt->manifest; + installRequest.deviceId = options.deviceId; + installRequest.rank = options.rank; + installRequest.provider = options.provider; + installRequest.offlineOnly = options.offlineOnly; + installRequest.driverAdapter = options.driverAdapter; + installRequest.repositoryMemoryOps = options.repositoryMemoryOps; + installRequest.repositoryMemoryUserData = options.repositoryMemoryUserData; + installRequest.repositoryInstallOptions = options.repositoryInstallOptions; + installRequest.repositoryReceipt = &attempt->repositoryReceipt; + installRequest.installOrder = options.installOrder; + installRequest.lowerLayerPlan = + options.prepareLowerLayerPlan == nullptr ? options.lowerLayerPlan : &attempt->preparedLowerLayerPlan; + + const int installRet = TileXRCcuInstallHardware( + installRequest, + &attempt->evidence, + &attempt->installReport); + + const int submitRet = TileXRCcuPrepareSubmitTasks( + attempt->package, + attempt->evidence, + &attempt->submitTasks, + &attempt->providerReport); + if (submitRet == TILEXR_SUCCESS) { + return ReturnWithAttemptStatus(TILEXR_SUCCESS, attempt, report); + } + if (installRet != TILEXR_SUCCESS) { + return ReturnWithAttemptStatus(installRet, attempt, report); + } + return ReturnWithAttemptStatus(submitRet, attempt, report); +} + +int TileXRCcuRunDirectInstallAttempt( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + return RunDirectInstallAttemptImpl(options, nullptr, attempt, report); +} + +int TileXRCcuRunDirectMemoryCopyInstallAttempt( + const TileXRCcuDirectInstallOptions& options, + const TileXRCcuDirectMemoryCopySpec& memoryCopy, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report) +{ + if (memoryCopy.localAddr == 0 || memoryCopy.localToken == 0 || + memoryCopy.remoteAddr == 0 || memoryCopy.remoteToken == 0 || + memoryCopy.lengthBytes == 0) { + ResetReport(report); + ClearAttempt(attempt); + if (report != nullptr) { + report->message = "invalid direct CCU memory copy address/token inputs"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + return RunDirectInstallAttemptImpl(options, &memoryCopy, attempt, report); +} + +int TileXRCcuReleaseDirectInstallAttemptResources(TileXRCcuDirectInstallAttempt& attempt) +{ + if (attempt.repositoryReceipt.deviceInstructionPtr == nullptr) { + attempt.repositoryReceipt = TileXRCcuRepositoryInstallReceipt{}; + attempt.repositoryReleaseReport = TileXRCcuRepositoryReport{}; + attempt.repositoryReleaseReport.message = "ok"; + return TILEXR_SUCCESS; + } + return TileXRCcuReleaseRepositoryInstallReceipt( + attempt.repositoryReceipt, + attempt.repositoryMemoryOps, + attempt.repositoryMemoryUserData, + &attempt.repositoryReleaseReport); +} + +} // namespace TileXR diff --git a/src/comm/ccu/tilexr_ccu_direct_orchestrator.h b/src/comm/ccu/tilexr_ccu_direct_orchestrator.h new file mode 100644 index 00000000..60530cf7 --- /dev/null +++ b/src/comm/ccu/tilexr_ccu_direct_orchestrator.h @@ -0,0 +1,137 @@ +/* + * Copyright (c) 2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_CCU_DIRECT_ORCHESTRATOR_H +#define TILEXR_CCU_DIRECT_ORCHESTRATOR_H + +#include "ccu/tilexr_ccu_install_provider.h" +#include "ccu/tilexr_ccu_lower_layer_plan_builder.h" +#include "ccu/tilexr_ccu_memory_program.h" +#include "ccu/tilexr_ccu_specs.h" + +#include +#include +#include + +namespace TileXR { + +using TileXRCcuLowerLayerPlanPrepareFn = int (*)( + const TileXRCcuResourceAllocation& allocation, + TileXRCcuLowerLayerInstallPlan* plan, + TileXRCcuLowerLayerPlanBuilderReport* report, + void* userData); + +struct TileXRCcuDirectInstallOptions { + const TileXRCcuBasicInfo* basicInfo = nullptr; + uint32_t sqeArgCount = 0; + uint32_t syncResourceCount = 0; + uint32_t syncInstructionCount = 0; + uint32_t bindingsPerSyncResource = 1; + uint16_t missionStartId = 0; + uint16_t instructionStartId = 0; + uint16_t missionInstructionStartId = 0; + uint16_t xnStartId = 0; + uint16_t gsaStartId = 0; + uint16_t remoteXnStartId = 0; + uint16_t remoteXnCount = 0; + uint16_t ckeStartId = 0; + uint16_t channelStartId = 0; + uint16_t localWaitCkeStartId = 0; + uint16_t localWaitCkeCount = 0; + uint16_t remoteNotifyCkeStartId = 0; + uint16_t remoteNotifyCkeCount = 0; + TileXRCcuBarrierMode barrierMode = TileXRCcuBarrierMode::SyncXn; + uint32_t deviceId = 0; + uint32_t rank = 0; + std::string provider; + bool offlineOnly = false; + const TileXRCcuDriverAdapter* driverAdapter = nullptr; + TileXRCcuDeviceMemoryOps repositoryMemoryOps; + void* repositoryMemoryUserData = nullptr; + TileXRCcuRepositoryInstallOptions repositoryInstallOptions; + TileXRCcuRepositoryMemoryAllocMode repositoryMemoryAllocMode = TileXRCcuRepositoryMemoryAllocMode::Acl; + TileXRCcuInstallOrder installOrder = TileXRCcuInstallOrder::InstallLowerLayerFirst; + const TileXRCcuLowerLayerInstallPlan* lowerLayerPlan = nullptr; + TileXRCcuLowerLayerPlanPrepareFn prepareLowerLayerPlan = nullptr; + void* lowerLayerPlanUserData = nullptr; +}; + +struct TileXRCcuDirectMemoryCopySpec { + TileXRCcuMemoryCopyDirection direction = TileXRCcuMemoryCopyDirection::RemoteToLocal; + uint64_t localAddr = 0; + uint64_t localToken = 0; + uint64_t remoteAddr = 0; + uint64_t remoteToken = 0; + uint64_t lengthBytes = 0; +}; + +struct TileXRCcuDirectInstallAttempt { + TileXRCcuSpecInfo specInfo; + TileXRCcuResourceSpec resourceSpec; + TileXRCcuResourceRequest resourceRequest; + TileXRCcuResourceAllocation allocation; + TileXRCcuProducerPlan plan; + TileXRCcuLaunchPackage package; + TileXRCcuInstallManifest manifest; + TileXRCcuHardwareInstallEvidence evidence; + TileXRCcuInstallProviderReport installReport; + TileXRCcuProviderReport providerReport; + TileXRCcuRepositoryInstallReceipt repositoryReceipt; + TileXRCcuRepositoryReport repositoryReleaseReport; + TileXRCcuDeviceMemoryOps repositoryMemoryOps; + void* repositoryMemoryUserData = nullptr; + TileXRCcuLowerLayerInstallPlan preparedLowerLayerPlan; + TileXRCcuLowerLayerPlanBuilderReport lowerLayerPlanReport; + std::vector submitTasks; +}; + +struct TileXRCcuDirectInstallReport { + bool pipelineBuilt = false; + bool installAttempted = false; + bool installSucceeded = false; + bool submitReady = false; + uint32_t requiredInstallSurfaceCount = 0; + uint32_t publicVerifiedInstallSurfaceCount = 0; + uint32_t missingInstallSurfaceCount = 0; + uint32_t taskCount = 0; + uint32_t submitTaskCount = 0; + std::string message; +}; + +struct TileXRCcuDirectSubmitReport { + bool submitted = false; + uint32_t taskCount = 0; + uint32_t submittedTaskCount = 0; + std::string message; +}; + +using TileXRCcuTaskSubmitFn = int (*)( + const TileXRCcuTask& task, + void* stream, + void* userData); + +int TileXRCcuRunDirectInstallAttempt( + const TileXRCcuDirectInstallOptions& options, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report); + +int TileXRCcuRunDirectMemoryCopyInstallAttempt( + const TileXRCcuDirectInstallOptions& options, + const TileXRCcuDirectMemoryCopySpec& memoryCopy, + TileXRCcuDirectInstallAttempt* attempt, + TileXRCcuDirectInstallReport* report); + +int TileXRCcuReleaseDirectInstallAttemptResources(TileXRCcuDirectInstallAttempt& attempt); + +int TileXRCcuSubmitPreparedTasks( + const std::vector& submitTasks, + void* stream, + TileXRCcuTaskSubmitFn submitFn, + void* submitUserData, + TileXRCcuDirectSubmitReport* report); + +} // namespace TileXR + +#endif // TILEXR_CCU_DIRECT_ORCHESTRATOR_H diff --git a/src/comm/ccu/tilexr_ccu_direct_runtime.cpp b/src/comm/ccu/tilexr_ccu_direct_runtime.cpp new file mode 100644 index 00000000..c1291d26 --- /dev/null +++ b/src/comm/ccu/tilexr_ccu_direct_runtime.cpp @@ -0,0 +1,1529 @@ +/* + * Copyright (c) 2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#include "ccu/tilexr_ccu_direct_runtime.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace TileXR { +namespace { + +constexpr uint32_t TILEXR_CCU_DEFAULT_DIRECT_SQ_DEPTH = 8; +constexpr uint32_t TILEXR_CCU_DIRECT_CCUM_SQE_BYTES = 64; +constexpr uint32_t TILEXR_CCU_DIRECT_SQ_EBB_WORDS = 4; +constexpr uint32_t TILEXR_CCU_DIRECT_LOOP_JETTY_ID = 1024; +constexpr uint32_t TILEXR_CCU_DIRECT_LOOP_JETTY_CTX_ID = 0; +constexpr uint64_t TILEXR_CCU_V1_WQE_BASIC_BLOCK_OFFSET = TILEXR_CCU_V1_CCUM_OFFSET + 0x800000ULL; +constexpr uint64_t TILEXR_CCU_DIRECT_SQ_BUFFER_BYTES = 256ULL * 1024ULL; +constexpr uint32_t TILEXR_CCU_DIRECT_CCU_POLL_CQ_DEPTH = 64; +constexpr uint32_t TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_ASYNC_MAX_POLLS = 1000; +constexpr uint32_t TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_ASYNC_SLEEP_US = 1000; +constexpr uint32_t TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_MAX_ATTEMPTS = 8; +constexpr int TILEXR_CCU_DIRECT_MAX_RANK_SIZE = 128; +constexpr int TILEXR_CCU_HCCP_JFC_MODE_CCU_POLL = 2; +constexpr int TILEXR_CCU_HCCP_ASYNC_EAGAIN = 128301; +constexpr const char* TILEXR_CCU_DIRECT_HDC_TYPE_ENV = "TILEXR_CCU_DIRECT_HDC_TYPE"; +constexpr const char* TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_ID_ENV = + "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_ID"; +constexpr const char* TILEXR_CCU_DIRECT_RESOURCE_WINDOW_RAW_TOKEN_ID_ENV = + "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_RAW_TOKEN_ID"; +constexpr const char* TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_VALUE_ENV = + "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_VALUE"; +constexpr const char* TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE_ENV = + "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE"; +constexpr const char* TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE_ENV = + "TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE"; +constexpr const char* TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE_ENV = + "TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE"; +constexpr const char* TILEXR_CCU_DIRECT_TRUST_SYNTHETIC_ENDPOINT_ROUTE_ENV = + "TILEXR_CCU_DIRECT_TRUST_SYNTHETIC_ENDPOINT_ROUTE"; +constexpr const char* TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_EXCHANGE_MODE_ENV = + "TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_EXCHANGE_MODE"; +constexpr const char* TILEXR_CCU_DIRECT_REMOTE_CCU_VA_OFFSET_ENV = + "TILEXR_CCU_DIRECT_REMOTE_CCU_VA_OFFSET"; +constexpr uint64_t TILEXR_CCU_UB_MEM_PAGE_SIZE = 4096ULL; +constexpr uint32_t TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT = 8; + +struct TileXRCcuEndpointTpHandleExchange { + uint64_t tpHandles[TILEXR_CCU_DIRECT_MAX_RANK_SIZE] = {}; + uint32_t psn = 0; +}; + +void ResetReport(TileXRCcuDirectRuntimeReport* report) +{ + if (report != nullptr) { + *report = TileXRCcuDirectRuntimeReport{}; + } +} + +int Fail(TileXRCcuDirectRuntimeReport* report, const std::string& message, int code = TILEXR_ERROR_NOT_FOUND) +{ + if (report != nullptr) { + report->message = message; + } + return code; +} + +bool IsEmptyEid(const std::array& eid) +{ + return std::all_of(eid.begin(), eid.end(), [](uint8_t value) { return value == 0; }); +} + +std::array ReverseEndpointEid( + const std::array& eid) +{ + std::array reversed {}; + std::reverse_copy(eid.begin(), eid.end(), reversed.begin()); + return reversed; +} + +std::string FormatEndpointEid(const std::array& eid) +{ + std::ostringstream text; + text << std::hex << std::setfill('0'); + for (uint8_t byte : eid) { + text << std::setw(2) << static_cast(byte); + } + return text.str(); +} + +bool HasCompleteEndpointRoute(const TileXRCcuLowerLayerTransportRoute& route) +{ + return !IsEmptyEid(route.remoteEid) && + route.doorbellVa != 0 && + route.doorbellTokenId != 0 && + route.sqDepth != 0; +} + +bool UseImportedPeerEndpointRoute() +{ + const char* mode = std::getenv(TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_EXCHANGE_MODE_ENV); + return mode == nullptr || mode[0] == '\0' || std::strcmp(mode, "imported_peer") == 0; +} + +uint64_t SelectResourceWindowBytes(const TileXRCcuBasicInfo& basicInfo) +{ + (void)basicInfo; + return TILEXR_CCU_RESOURCE_WINDOW_BYTES; +} + +std::string RankEnvName(const char* base, int rank) +{ + std::ostringstream name; + name << base << "_RANK" << rank; + return name.str(); +} + +const char* SelectRankedEnv(const char* base, int rank) +{ + const std::string rankedName = RankEnvName(base, rank); + const char* ranked = std::getenv(rankedName.c_str()); + if (ranked != nullptr && ranked[0] != '\0') { + return ranked; + } + const char* value = std::getenv(base); + return value != nullptr && value[0] != '\0' ? value : nullptr; +} + +bool ParseUnsignedEnv(const char* value, uint64_t* out) +{ + if (value == nullptr || value[0] == '\0' || out == nullptr) { + return false; + } + errno = 0; + char* end = nullptr; + const unsigned long long parsed = std::strtoull(value, &end, 0); + if (errno != 0 || end == value || *end != '\0') { + return false; + } + *out = static_cast(parsed); + return true; +} + +uint64_t SelectRemoteCcuVaOffset() +{ + uint64_t value = 0; + return ParseUnsignedEnv(std::getenv(TILEXR_CCU_DIRECT_REMOTE_CCU_VA_OFFSET_ENV), &value) ? value : 0; +} + +bool ParseU32RankedEnv(const char* base, int rank, uint32_t* out) +{ + uint64_t value = 0; + if (!ParseUnsignedEnv(SelectRankedEnv(base, rank), &value) || value > 0xffffffffULL || out == nullptr) { + return false; + } + *out = static_cast(value); + return true; +} + +bool HasRankedEnv(const char* base, int rank) +{ + return SelectRankedEnv(base, rank) != nullptr; +} + +bool IsRaCtxResourceWindowRegistrationMode() +{ + const char* value = std::getenv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE_ENV); + if (value == nullptr || value[0] == '\0') { + return false; + } + const std::string mode(value); + return mode == "ra_ctx" || mode == "ractx" || mode == "public_ra_ctx" || mode == "1"; +} + +bool TrustSyntheticEndpointRouteForDiagnostics() +{ + const char* value = std::getenv(TILEXR_CCU_DIRECT_TRUST_SYNTHETIC_ENDPOINT_ROUTE_ENV); + return value != nullptr && value[0] != '\0' && value[0] != '0'; +} + +bool TraceEndpointRoute() +{ + const char* value = std::getenv(TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE_ENV); + return value != nullptr && value[0] != '\0' && value[0] != '0'; +} + +void TraceEndpointRouteStep(const std::string& message) +{ + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute " << message << std::endl; + } +} + +bool IsRaCtxLoopEndpointRouteCollectionMode() +{ + const char* value = std::getenv(TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE_ENV); + if (value == nullptr || value[0] == '\0') { + return false; + } + const std::string mode(value); + return mode == "ra_ctx_loop" || mode == "ractx_loop" || mode == "public_ra_ctx_loop"; +} + +uint64_t AlignResourceWindowAddr(uint64_t addr) +{ + return addr & ~(TILEXR_CCU_UB_MEM_PAGE_SIZE - 1ULL); +} + +bool HasRaCtxResourceWindowSymbols(const TileXRCcuHccpLoader& loader) +{ + return loader.RaGetDevEidInfoNum != nullptr && + loader.RaGetDevEidInfoList != nullptr && + loader.RaCtxInit != nullptr && + loader.RaCtxDeinit != nullptr && + loader.RaCtxTokenIdAlloc != nullptr && + loader.RaCtxTokenIdFree != nullptr && + loader.RaGetSecRandom != nullptr && + loader.RaCtxLmemRegister != nullptr && + loader.RaCtxLmemUnregister != nullptr; +} + +bool HasRaCtxEndpointRouteSymbols(const TileXRCcuHccpLoader& loader) +{ + return loader.RaCtxCqCreate != nullptr && + loader.RaCtxCqDestroy != nullptr && + loader.RaCtxQpCreate != nullptr && + loader.RaCtxQpDestroy != nullptr && + loader.RaCtxQpImport != nullptr && + loader.RaCtxQpUnimport != nullptr && + loader.RaGetTpInfoListAsync != nullptr && + loader.RaGetAsyncReqResult != nullptr; +} + +uint32_t SelectEndpointRouteSqDepth() +{ + return TILEXR_CCU_DEFAULT_DIRECT_SQ_DEPTH; +} + +uint32_t SelectEndpointRouteSqBytes(uint32_t sqDepth) +{ + return sqDepth * TILEXR_CCU_DIRECT_SQ_EBB_WORDS * TILEXR_CCU_DIRECT_CCUM_SQE_BYTES; +} + +uint64_t SelectEndpointRouteSqVa(const TileXRCcuLocalResourceWindowInfo& localResourceWindow) +{ + return localResourceWindow.addr + TILEXR_CCU_V1_WQE_BASIC_BLOCK_OFFSET + + static_cast(TILEXR_CCU_DIRECT_LOOP_JETTY_CTX_ID) * TILEXR_CCU_DIRECT_SQ_BUFFER_BYTES; +} + +int WaitRaCtxAsyncRequest(TileXRCcuHccpLoader& loader, void* reqHandle) +{ + if (loader.RaGetAsyncReqResult == nullptr || reqHandle == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + for (uint32_t poll = 0; poll < TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_ASYNC_MAX_POLLS; ++poll) { + int reqResult = 0; + const int ret = loader.RaGetAsyncReqResult(reqHandle, &reqResult); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute asyncPoll poll=" << poll + << " ret=" << ret + << " reqResult=" << reqResult << std::endl; + } + if (ret == TILEXR_CCU_HCCP_ASYNC_EAGAIN) { + std::this_thread::sleep_for( + std::chrono::microseconds(TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_ASYNC_SLEEP_US)); + continue; + } + if (ret != 0 || reqResult != 0) { + return TILEXR_ERROR_MKIRT; + } + return TILEXR_SUCCESS; + } + return TILEXR_ERROR_TIMEOUT; +} + +bool ApplyResourceWindowTokenOverride( + int rank, + TileXRCcuLocalResourceWindowInfo* window) +{ + if (window == nullptr || !HasRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_ID_ENV, rank)) { + return false; + } + uint32_t tokenId = 0; + if (!ParseU32RankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_ID_ENV, rank, &tokenId) || tokenId == 0) { + return false; + } + + uint32_t rawTokenId = tokenId; + if (HasRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_RAW_TOKEN_ID_ENV, rank) && + !ParseU32RankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_RAW_TOKEN_ID_ENV, rank, &rawTokenId)) { + return false; + } + + uint32_t tokenValue = 0; + if (HasRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_VALUE_ENV, rank) && + !ParseU32RankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_VALUE_ENV, rank, &tokenValue)) { + return false; + } else if (!HasRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_VALUE_ENV, rank)) { + tokenValue = window->tokenValue; + } + + window->tokenId = tokenId; + window->rawTokenId = rawTokenId; + window->tokenValue = tokenValue; + return true; +} + +int SelectDirectCcuHdcType() +{ + uint64_t value = 0; + if (!ParseUnsignedEnv(std::getenv(TILEXR_CCU_DIRECT_HDC_TYPE_ENV), &value)) { + return TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2; + } + if (value == TILEXR_CCU_HDC_SERVICE_TYPE_RDMA || value == TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2) { + return static_cast(value); + } + return TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2; +} + +bool ParseHexBytePair(char high, char low, uint8_t* out) +{ + auto hexValue = [](char c) -> int { + if (c >= '0' && c <= '9') { + return c - '0'; + } + if (c >= 'a' && c <= 'f') { + return c - 'a' + 10; + } + if (c >= 'A' && c <= 'F') { + return c - 'A' + 10; + } + return -1; + }; + const int highValue = hexValue(high); + const int lowValue = hexValue(low); + if (highValue < 0 || lowValue < 0 || out == nullptr) { + return false; + } + *out = static_cast((highValue << 4U) | lowValue); + return true; +} + +bool ParseEndpointEid(const char* value, std::array* eid) +{ + if (value == nullptr || eid == nullptr) { + return false; + } + std::string text(value); + if (text.size() >= 2 && text[0] == '0' && (text[1] == 'x' || text[1] == 'X')) { + text.erase(0, 2); + } + text.erase(std::remove_if(text.begin(), text.end(), [](char c) { + return c == ':' || c == '-' || c == '_' || c == ' '; + }), text.end()); + if (text.size() != TILEXR_CCU_EID_BYTES * 2U) { + return false; + } + std::array parsed {}; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + if (!ParseHexBytePair(text[i * 2U], text[i * 2U + 1U], &parsed[i])) { + return false; + } + } + *eid = parsed; + return true; +} + +int CollectLocalEndpointRouteFromEnv( + int rank, + TileXRCcuLowerLayerTransportRoute* route, + TileXRCcuDirectRuntimeReport* report) +{ + if (route == nullptr) { + return Fail(report, "missing output direct CCU local endpoint route", TILEXR_ERROR_PARA_CHECK_FAIL); + } + + const char* eidEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_EID", rank); + const char* tpnEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_TPN", rank); + const char* doorbellVaEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_VA", rank); + const char* tokenIdEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_TOKEN_ID", rank); + const char* tokenValueEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_TOKEN_VALUE", rank); + const char* sqDepthEnv = SelectRankedEnv("TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_SQ_DEPTH", rank); + + if (eidEnv == nullptr && tpnEnv == nullptr && doorbellVaEnv == nullptr && + tokenIdEnv == nullptr && tokenValueEnv == nullptr && sqDepthEnv == nullptr) { + return Fail(report, "direct CCU local endpoint route collector is not configured", TILEXR_ERROR_NOT_FOUND); + } + if (eidEnv == nullptr || tpnEnv == nullptr || doorbellVaEnv == nullptr || + tokenIdEnv == nullptr || sqDepthEnv == nullptr) { + return Fail(report, "direct CCU local endpoint route env is incomplete", TILEXR_ERROR_PARA_CHECK_FAIL); + } + + TileXRCcuLowerLayerTransportRoute parsed; + if (!ParseEndpointEid(eidEnv, &parsed.remoteEid)) { + return Fail(report, "invalid direct CCU local endpoint EID env", TILEXR_ERROR_PARA_CHECK_FAIL); + } + uint64_t value = 0; + if (!ParseUnsignedEnv(tpnEnv, &value) || value == 0 || value > 0xffffffffULL) { + return Fail(report, "invalid direct CCU local endpoint TPN env", TILEXR_ERROR_PARA_CHECK_FAIL); + } + parsed.tpn = static_cast(value); + if (!ParseUnsignedEnv(doorbellVaEnv, &parsed.doorbellVa)) { + return Fail(report, "invalid direct CCU local endpoint doorbell VA env", TILEXR_ERROR_PARA_CHECK_FAIL); + } + if (!ParseUnsignedEnv(tokenIdEnv, &value) || value == 0 || value > 0xffffffffULL) { + return Fail(report, "invalid direct CCU local endpoint doorbell token id env", TILEXR_ERROR_PARA_CHECK_FAIL); + } + parsed.doorbellTokenId = static_cast(value); + if (tokenValueEnv != nullptr) { + if (!ParseUnsignedEnv(tokenValueEnv, &value) || value > 0xffffffffULL) { + return Fail(report, "invalid direct CCU local endpoint doorbell token value env", + TILEXR_ERROR_PARA_CHECK_FAIL); + } + parsed.doorbellTokenValue = static_cast(value); + } + if (!ParseUnsignedEnv(sqDepthEnv, &value) || value == 0 || value > 0xffffffffULL) { + return Fail(report, "invalid direct CCU local endpoint SQ depth env", TILEXR_ERROR_PARA_CHECK_FAIL); + } + parsed.sqDepth = static_cast(value); + parsed.endpointRouteVerified = true; + *route = parsed; + return TILEXR_SUCCESS; +} + +void FillProviderResourceWindow( + const TileXRCcuLocalResourceWindowInfo& localResourceWindow, + TileXRCcuEndpointRouteProviderResourceWindow* providerWindow) +{ + if (providerWindow == nullptr) { + return; + } + providerWindow->addr = localResourceWindow.addr; + providerWindow->bytes = localResourceWindow.bytes; + providerWindow->tokenId = localResourceWindow.tokenId; + providerWindow->rawTokenId = localResourceWindow.rawTokenId; + providerWindow->tokenValue = localResourceWindow.tokenValue; +} + +void CopyProviderRoute( + const TileXRCcuEndpointRouteProviderRoute& providerRoute, + TileXRCcuLowerLayerTransportRoute* route) +{ + if (route == nullptr) { + return; + } + TileXRCcuLowerLayerTransportRoute copied; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + copied.remoteEid[i] = providerRoute.remoteEid[i]; + } + copied.tpn = providerRoute.tpn; + copied.doorbellVa = providerRoute.doorbellVa; + copied.doorbellTokenId = providerRoute.doorbellTokenId; + copied.doorbellTokenValue = providerRoute.doorbellTokenValue; + copied.sqDepth = providerRoute.sqDepth; + copied.endpointRouteVerified = providerRoute.endpointRouteVerified; + *route = copied; +} + +} // namespace + +int TileXRCcuDirectRuntime::Init( + const TileXRCcuDirectRuntimeOptions& options, + TileXRCcuDirectRuntimeReport* report) +{ + ResetReport(report); + Shutdown(); + options_ = options; + + TileXRCcuHccpLoaderReport loaderReport; + int ret = loader_.Load(&loaderReport); + if (ret != TILEXR_SUCCESS) { + return Fail(report, loaderReport.message); + } + + ret = ResolveDevicePhyId(&devicePhyId_, report); + if (ret != TILEXR_SUCCESS) { + Shutdown(); + return ret; + } + const int hdcType = SelectDirectCcuHdcType(); + if (report != nullptr) { + report->hdcType = hdcType; + } + TileXRCcuHccpLoaderReport raReport; + ret = loader_.InitRaHdc(devicePhyId_, hdcType, true, &raReport); + if (ret != TILEXR_SUCCESS) { + Shutdown(); + return Fail(report, raReport.message, ret); + } + if (report != nullptr) { + report->raInitialized = true; + } + + TileXRCcuHccpLoaderReport tlvReport; + ret = loader_.InitCcuTlv(devicePhyId_, &tlvReport); + if (ret != TILEXR_SUCCESS) { + Shutdown(); + return Fail(report, tlvReport.message, ret); + } + + initialized_ = true; + if (report != nullptr) { + report->initialized = true; + report->raInitialized = true; + report->ccuTlvInitialized = true; + report->logicDevId = static_cast(options_.devId); + report->devicePhyId = devicePhyId_; + report->hdcType = hdcType; + report->message = "ok"; + } + return TILEXR_SUCCESS; +} + +void TileXRCcuDirectRuntime::Shutdown() +{ + ReleaseLocalEndpointRoute(); + ReleaseRegisteredResourceWindow(); + loader_.Unload(); + cachedBasicInfo_ = TileXRCcuBasicInfo{}; + localResourceWindow_ = TileXRCcuLocalResourceWindowInfo{}; + localVerifiedEndpointRoute_ = TileXRCcuLowerLayerTransportRoute{}; + cachedBasicInfoValid_ = false; + resourceWindowRegistered_ = false; + localVerifiedEndpointRouteValid_ = false; + endpointChanHandle_ = nullptr; + endpointCqHandle_ = nullptr; + endpointQpHandle_ = nullptr; + endpointRemoteQpHandle_ = nullptr; + endpointPeerRemoteQpHandles_.clear(); + endpointQpKey_ = TileXRCcuHccpQpKey{}; + endpointQpKeyValid_ = false; + endpointRouteBound_ = false; + endpointPsn_ = 1; + devicePhyId_ = 0; + initialized_ = false; +} + +bool TileXRCcuDirectRuntime::IsAvailable() const +{ + return initialized_ && loader_.IsLoaded() && loader_.RaCustomChannel != nullptr; +} + +int TileXRCcuDirectRuntime::ResolveDevicePhyId(uint32_t* devicePhyId, TileXRCcuDirectRuntimeReport* report) const +{ + if (devicePhyId == nullptr) { + return Fail(report, "missing output CCU physical device id", TILEXR_ERROR_PARA_CHECK_FAIL); + } + TileXRCcuHccpLoaderReport loaderReport; + const int ret = loader_.ResolveDevicePhyId( + static_cast(options_.devId), + devicePhyId, + &loaderReport); + if (ret != TILEXR_SUCCESS) { + return Fail(report, loaderReport.message.empty() ? + "failed to resolve CCU physical device id" : loaderReport.message, ret); + } + if (report != nullptr) { + report->logicDevId = static_cast(options_.devId); + report->devicePhyId = *devicePhyId; + } + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::CreateDriverAdapter( + TileXRCcuDriverAdapter* adapter, + TileXRCcuDriverAdapterReport* report) +{ + if (!IsAvailable()) { + if (report != nullptr) { + *report = TileXRCcuDriverAdapterReport{}; + report->message = "direct CCU runtime is unavailable"; + } + return TILEXR_ERROR_NOT_FOUND; + } + + TileXRCcuRaCustomChannelProviderReport providerReport; + int ret = raCustomChannelProvider_.Init(devicePhyId_, loader_.RaCustomChannel, &providerReport); + if (ret != TILEXR_SUCCESS) { + if (report != nullptr) { + *report = TileXRCcuDriverAdapterReport{}; + report->message = providerReport.message; + } + return ret; + } + return raCustomChannelProvider_.CreateAdapter(adapter, report); +} + +int TileXRCcuDirectRuntime::QueryBasicInfo( + uint8_t dieId, + TileXRCcuBasicInfo* basicInfo, + TileXRCcuDriverAdapterReport* report) +{ + if (basicInfo == nullptr) { + if (report != nullptr) { + *report = TileXRCcuDriverAdapterReport{}; + report->message = "missing output CCU basic info"; + } + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + TileXRCcuDriverAdapter adapter; + int ret = CreateDriverAdapter(&adapter, report); + if (ret != TILEXR_SUCCESS) { + return ret; + } + bool enabled = false; + ret = adapter.GetDieEnabled(dieId, &enabled, report); + if (ret != TILEXR_SUCCESS) { + return ret; + } + if (!enabled) { + if (report != nullptr) { + report->message = "direct CCU die is not enabled"; + } + return TILEXR_ERROR_NOT_FOUND; + } + ret = adapter.GetBasicInfo(dieId, basicInfo, report); + if (ret == TILEXR_SUCCESS) { + cachedBasicInfo_ = *basicInfo; + cachedBasicInfoValid_ = true; + } + return ret; +} + +int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBuffer(uint64_t resourceAddr) +{ + if (!IsAvailable()) { + return TILEXR_ERROR_NOT_INITIALIZED; + } + ReleaseRegisteredResourceWindow(); + if (!cachedBasicInfoValid_ || resourceAddr == 0) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + if (cachedBasicInfo_.resourceAddr != 0 && cachedBasicInfo_.resourceAddr != resourceAddr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + const uint64_t resourceBytes = SelectResourceWindowBytes(cachedBasicInfo_); + if (IsRaCtxResourceWindowRegistrationMode()) { + return RegisterCcuResourceRmaBufferWithRaCtx(resourceAddr, resourceBytes); + } + + localResourceWindow_ = TileXRCcuLocalResourceWindowInfo{}; + localResourceWindow_.addr = resourceAddr; + localResourceWindow_.bytes = resourceBytes; + localResourceWindow_.tokenId = cachedBasicInfo_.msidToken.tokenId; + localResourceWindow_.rawTokenId = cachedBasicInfo_.msidToken.tokenId; + localResourceWindow_.tokenValue = cachedBasicInfo_.msidToken.tokenValue; + if (HasRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_TOKEN_ID_ENV, options_.rank) && + !ApplyResourceWindowTokenOverride(options_.rank, &localResourceWindow_)) { + localResourceWindow_ = TileXRCcuLocalResourceWindowInfo{}; + resourceWindowRegistered_ = false; + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + resourceWindowRegistered_ = true; + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBufferWithRaCtx( + uint64_t resourceAddr, + uint64_t resourceBytes) +{ + if (!HasRaCtxResourceWindowSymbols(loader_)) { + return TILEXR_ERROR_NOT_FOUND; + } + + TileXRCcuRaInfo raInfo {}; + raInfo.mode = TILEXR_CCU_NETWORK_OFFLINE; + raInfo.phyId = devicePhyId_; + + uint32_t eidNum = 0; + int ret = loader_.RaGetDevEidInfoNum(raInfo, &eidNum); + if (ret != 0 || eidNum == 0) { + return ret == 0 ? TILEXR_ERROR_NOT_FOUND : TILEXR_ERROR_MKIRT; + } + + std::vector eidInfos(eidNum); + uint32_t queriedEidNum = eidNum; + ret = loader_.RaGetDevEidInfoList(raInfo, eidInfos.data(), &queriedEidNum); + if (ret != 0 || queriedEidNum == 0) { + return ret == 0 ? TILEXR_ERROR_NOT_FOUND : TILEXR_ERROR_MKIRT; + } + + void* ctxHandle = nullptr; + void* tokenIdHandle = nullptr; + void* lmemHandle = nullptr; + + TileXRCcuHccpCtxInitCfg ctxCfg {}; + ctxCfg.mode = TILEXR_CCU_NETWORK_OFFLINE; + ctxCfg.rdma.disabledLiteThread = false; + TileXRCcuHccpCtxInitAttr ctxAttr {}; + ctxAttr.phyId = devicePhyId_; + ctxAttr.ub.eidIndex = eidInfos[0].eidIndex; + ctxAttr.ub.eid = eidInfos[0].eid; + + ret = loader_.RaCtxInit(&ctxCfg, &ctxAttr, &ctxHandle); + if (ret != 0 || ctxHandle == nullptr) { + return TILEXR_ERROR_MKIRT; + } + + TileXRCcuHccpTokenId allocatedToken {}; + ret = loader_.RaCtxTokenIdAlloc(ctxHandle, &allocatedToken, &tokenIdHandle); + if (ret != 0 || tokenIdHandle == nullptr) { + if (loader_.RaCtxDeinit != nullptr) { + (void)loader_.RaCtxDeinit(ctxHandle); + } + return TILEXR_ERROR_MKIRT; + } + + uint32_t tokenValue = 0; + TileXRCcuRaInfo randomInfo {}; + randomInfo.mode = TILEXR_CCU_NETWORK_OFFLINE; + randomInfo.phyId = devicePhyId_; + ret = loader_.RaGetSecRandom(&randomInfo, &tokenValue); + if (ret != 0) { + if (loader_.RaCtxTokenIdFree != nullptr) { + (void)loader_.RaCtxTokenIdFree(ctxHandle, tokenIdHandle); + } + if (loader_.RaCtxDeinit != nullptr) { + (void)loader_.RaCtxDeinit(ctxHandle); + } + return TILEXR_ERROR_MKIRT; + } + + const uint64_t alignedAddr = AlignResourceWindowAddr(resourceAddr); + const uint64_t alignedBytes = resourceBytes + (resourceAddr - alignedAddr); + TileXRCcuHccpMrRegInfo mr {}; + mr.in.mem.addr = alignedAddr; + mr.in.mem.size = alignedBytes; + mr.in.ub.flags.value = 0; + mr.in.ub.flags.bs.tokenPolicy = TILEXR_CCU_HCCP_TOKEN_POLICY_PLAIN_TEXT; + mr.in.ub.flags.bs.tokenIdValid = 1; + mr.in.ub.flags.bs.access = TILEXR_CCU_HCCP_MEM_SEG_ACCESS_DEFAULT; + mr.in.ub.flags.bs.nonPin = 1; + mr.in.ub.tokenValue = tokenValue; + mr.in.ub.tokenIdHandle = tokenIdHandle; + + ret = loader_.RaCtxLmemRegister(ctxHandle, &mr, &lmemHandle); + if (ret != 0 || lmemHandle == nullptr) { + if (lmemHandle != nullptr && loader_.RaCtxLmemUnregister != nullptr) { + (void)loader_.RaCtxLmemUnregister(ctxHandle, lmemHandle); + } + if (loader_.RaCtxTokenIdFree != nullptr) { + (void)loader_.RaCtxTokenIdFree(ctxHandle, tokenIdHandle); + } + if (loader_.RaCtxDeinit != nullptr) { + (void)loader_.RaCtxDeinit(ctxHandle); + } + return TILEXR_ERROR_MKIRT; + } + + const uint32_t rawTokenId = mr.out.ub.tokenId != 0 ? mr.out.ub.tokenId : allocatedToken.tokenId; + localResourceWindow_ = TileXRCcuLocalResourceWindowInfo{}; + localResourceWindow_.addr = resourceAddr; + localResourceWindow_.bytes = resourceBytes; + localResourceWindow_.rawTokenId = rawTokenId; + localResourceWindow_.tokenId = rawTokenId >> TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT; + localResourceWindow_.tokenValue = tokenValue; + localResourceWindow_.targetSegHandle = mr.out.ub.targetSegHandle; + localResourceWindow_.raCtxHandle = ctxHandle; + localResourceWindow_.tokenIdHandle = tokenIdHandle; + localResourceWindow_.lmemHandle = lmemHandle; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + localResourceWindow_.eid[i] = eidInfos[0].eid.raw[i]; + } + localResourceWindow_.eidIndex = eidInfos[0].eidIndex; + localResourceWindow_.raCtxRegistered = true; + resourceWindowRegistered_ = true; + return TILEXR_SUCCESS; +} + +void TileXRCcuDirectRuntime::ReleaseRegisteredResourceWindow() +{ + ReleaseLocalEndpointRoute(); + if (localResourceWindow_.raCtxRegistered) { + if (localResourceWindow_.lmemHandle != nullptr && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxLmemUnregister != nullptr) { + (void)loader_.RaCtxLmemUnregister(localResourceWindow_.raCtxHandle, localResourceWindow_.lmemHandle); + } + if (localResourceWindow_.tokenIdHandle != nullptr && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxTokenIdFree != nullptr) { + (void)loader_.RaCtxTokenIdFree(localResourceWindow_.raCtxHandle, localResourceWindow_.tokenIdHandle); + } + if (localResourceWindow_.raCtxHandle != nullptr && loader_.RaCtxDeinit != nullptr) { + (void)loader_.RaCtxDeinit(localResourceWindow_.raCtxHandle); + } + } + localResourceWindow_ = TileXRCcuLocalResourceWindowInfo{}; + localVerifiedEndpointRoute_ = TileXRCcuLowerLayerTransportRoute{}; + resourceWindowRegistered_ = false; + localVerifiedEndpointRouteValid_ = false; +} + +void TileXRCcuDirectRuntime::ReleasePeerEndpointImports() +{ + if (localResourceWindow_.raCtxHandle != nullptr && loader_.RaCtxQpUnimport != nullptr) { + for (void* handle : endpointPeerRemoteQpHandles_) { + if (handle != nullptr) { + (void)loader_.RaCtxQpUnimport(localResourceWindow_.raCtxHandle, handle); + } + } + } + endpointPeerRemoteQpHandles_.clear(); +} + +void TileXRCcuDirectRuntime::ReleaseLocalEndpointRoute() +{ + ReleasePeerEndpointImports(); + if (endpointRouteBound_ && endpointQpHandle_ != nullptr && loader_.RaCtxQpUnbind != nullptr) { + (void)loader_.RaCtxQpUnbind(endpointQpHandle_); + } + if (endpointRemoteQpHandle_ != nullptr && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxQpUnimport != nullptr) { + (void)loader_.RaCtxQpUnimport(localResourceWindow_.raCtxHandle, endpointRemoteQpHandle_); + } + if (endpointQpHandle_ != nullptr && loader_.RaCtxQpDestroy != nullptr) { + (void)loader_.RaCtxQpDestroy(endpointQpHandle_); + } + if (endpointCqHandle_ != nullptr && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxCqDestroy != nullptr) { + (void)loader_.RaCtxCqDestroy(localResourceWindow_.raCtxHandle, endpointCqHandle_); + } + if (endpointChanHandle_ != nullptr && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxChanDestroy != nullptr) { + (void)loader_.RaCtxChanDestroy(localResourceWindow_.raCtxHandle, endpointChanHandle_); + } + endpointChanHandle_ = nullptr; + endpointCqHandle_ = nullptr; + endpointQpHandle_ = nullptr; + endpointRemoteQpHandle_ = nullptr; + endpointQpKey_ = TileXRCcuHccpQpKey{}; + endpointQpKeyValid_ = false; + endpointRouteBound_ = false; + localVerifiedEndpointRoute_ = TileXRCcuLowerLayerTransportRoute{}; + localVerifiedEndpointRouteValid_ = false; +} + +int TileXRCcuDirectRuntime::CollectLocalEndpointRouteWithRaCtx(TileXRCcuLowerLayerTransportRoute* route) +{ + if (route == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *route = TileXRCcuLowerLayerTransportRoute{}; + int lastRet = TILEXR_ERROR_NOT_FOUND; + for (uint32_t attempt = 0; attempt < TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_MAX_ATTEMPTS; ++attempt) { + TileXRCcuLowerLayerTransportRoute attemptRoute; + bool asyncWaitFailed = false; + const int ret = CollectLocalEndpointRouteWithRaCtxOnce(&attemptRoute, &asyncWaitFailed); + if (ret == TILEXR_SUCCESS) { + *route = attemptRoute; + return TILEXR_SUCCESS; + } + lastRet = ret; + if (!asyncWaitFailed || attempt + 1 >= TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_MAX_ATTEMPTS) { + return ret; + } + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute retryAfterAsyncFailure" + << " attempt=" << (attempt + 1) + << " ret=" << ret << std::endl; + } + } + return lastRet; +} + +int TileXRCcuDirectRuntime::CollectLocalEndpointRouteWithRaCtxOnce( + TileXRCcuLowerLayerTransportRoute* route, + bool* asyncWaitFailed) +{ + if (route == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + if (asyncWaitFailed != nullptr) { + *asyncWaitFailed = false; + } + *route = TileXRCcuLowerLayerTransportRoute{}; + if (!localResourceWindow_.raCtxRegistered || + localResourceWindow_.raCtxHandle == nullptr || + localResourceWindow_.tokenIdHandle == nullptr || + IsEmptyEid(localResourceWindow_.eid) || + !HasRaCtxEndpointRouteSymbols(loader_)) { + TraceEndpointRouteStep("raCtxLoop precheck failed"); + return TILEXR_ERROR_NOT_FOUND; + } + + ReleaseLocalEndpointRoute(); + const uint32_t sqDepth = SelectEndpointRouteSqDepth(); + const uint64_t sqVa = SelectEndpointRouteSqVa(localResourceWindow_); + const uint32_t sqBytes = SelectEndpointRouteSqBytes(sqDepth); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute begin" + << " ctx=" << localResourceWindow_.raCtxHandle + << " sqDepth=" << sqDepth + << " sqVa=0x" << std::hex << sqVa + << " sqBytes=0x" << sqBytes + << std::dec << std::endl; + } + + TileXRCcuHccpCqInfo cqInfo {}; + cqInfo.in.chanHandle = nullptr; + cqInfo.in.depth = TILEXR_CCU_DIRECT_CCU_POLL_CQ_DEPTH; + cqInfo.in.ub.userCtx = 0; + cqInfo.in.ub.mode = TILEXR_CCU_HCCP_JFC_MODE_CCU_POLL; + cqInfo.in.ub.ceqn = 0; + cqInfo.in.ub.flag.value = 0; + int ret = loader_.RaCtxCqCreate(localResourceWindow_.raCtxHandle, &cqInfo, &endpointCqHandle_); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute cqCreate ret=" << ret + << " cq=" << endpointCqHandle_ << std::endl; + } + if (ret != 0 || endpointCqHandle_ == nullptr) { + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_MKIRT; + } + + TileXRCcuHccpQpCreateAttr qpAttr {}; + qpAttr.scqHandle = endpointCqHandle_; + qpAttr.rcqHandle = endpointCqHandle_; + qpAttr.srqHandle = endpointCqHandle_; + qpAttr.sqDepth = sqDepth; + qpAttr.rqDepth = TILEXR_CCU_HCCP_RQ_DEPTH_DEFAULT; + qpAttr.transportMode = TILEXR_CCU_HCCP_TRANSPORT_MODE_RM; + qpAttr.ub.mode = static_cast(TILEXR_CCU_HCCP_JETTY_MODE_CCU); + qpAttr.ub.jettyId = TILEXR_CCU_DIRECT_LOOP_JETTY_ID; + qpAttr.ub.tokenIdHandle = localResourceWindow_.tokenIdHandle; + qpAttr.ub.tokenValue = localResourceWindow_.tokenValue; + qpAttr.ub.flag.value = 0; + qpAttr.ub.flag.bs.shareJfr = 1; + qpAttr.ub.jfsFlag.bs.errorSuspend = 1; + qpAttr.ub.priority = 2; + qpAttr.ub.rnrRetry = TILEXR_CCU_HCCP_RNR_RETRY_DEFAULT; + qpAttr.ub.extMode.cstmFlag.value = 0; + qpAttr.ub.extMode.cstmFlag.bs.sqCstm = 1; + qpAttr.ub.extMode.sq.buffVa = sqVa; + qpAttr.ub.extMode.sq.buffSize = sqBytes; + qpAttr.ub.extMode.sqebbNum = sqDepth; + + TileXRCcuHccpQpCreateInfo qpInfo {}; + ret = loader_.RaCtxQpCreate(localResourceWindow_.raCtxHandle, &qpAttr, &qpInfo, &endpointQpHandle_); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute qpCreate ret=" << ret + << " qp=" << endpointQpHandle_ + << " keySize=" << static_cast(qpInfo.key.size) + << " dbAddr=0x" << std::hex << qpInfo.ub.dbAddr + << " dbTokenId=0x" << qpInfo.ub.dbTokenId + << std::dec << std::endl; + } + if (ret != 0 || endpointQpHandle_ == nullptr) { + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_MKIRT; + } + endpointQpKey_ = qpInfo.key; + endpointQpKeyValid_ = qpInfo.key.size != 0; + + TileXRCcuHccpGetTpCfg tpCfg {}; + tpCfg.flag.value = 0; + tpCfg.flag.bs.rtp = 1; + tpCfg.transMode = TILEXR_CCU_HCCP_TRANSPORT_MODE_RM; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + tpCfg.localEid.raw[i] = localResourceWindow_.eid[i]; + tpCfg.peerEid.raw[i] = localResourceWindow_.eid[i]; + } + TileXRCcuHccpTpInfo tpInfo {}; + uint32_t tpInfoNum = 1; + void* reqHandle = nullptr; + ret = loader_.RaGetTpInfoListAsync( + localResourceWindow_.raCtxHandle, + &tpCfg, + &tpInfo, + &tpInfoNum, + &reqHandle); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute getTpInfoAsync ret=" << ret + << " req=" << reqHandle + << " num=" << tpInfoNum + << " tpHandle=0x" << std::hex << tpInfo.tpHandle + << std::dec << std::endl; + } + if (ret != 0 || reqHandle == nullptr || tpInfoNum == 0) { + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_MKIRT; + } + ret = WaitRaCtxAsyncRequest(loader_, reqHandle); + if (ret != TILEXR_SUCCESS) { + ReleaseLocalEndpointRoute(); + if (asyncWaitFailed != nullptr) { + *asyncWaitFailed = true; + } + return ret; + } + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute getTpInfoDone" + << " tpHandle=0x" << std::hex << tpInfo.tpHandle + << std::dec << std::endl; + } + if (tpInfo.tpHandle == 0) { + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_NOT_FOUND; + } + + TileXRCcuHccpQpImportInfo importInfo {}; + importInfo.in.key = qpInfo.key; + importInfo.in.ub.mode = TILEXR_CCU_HCCP_JETTY_IMPORT_MODE_EXP; + importInfo.in.ub.tokenValue = localResourceWindow_.tokenValue; + importInfo.in.ub.policy = TILEXR_CCU_HCCP_JETTY_GRP_POLICY_RR; + importInfo.in.ub.type = TILEXR_CCU_HCCP_TARGET_TYPE_JETTY; + importInfo.in.ub.flag.value = 0; + importInfo.in.ub.flag.bs.tokenPolicy = TILEXR_CCU_HCCP_TOKEN_POLICY_PLAIN_TEXT; + importInfo.in.ub.expImportCfg.tpHandle = tpInfo.tpHandle; + importInfo.in.ub.expImportCfg.peerTpHandle = tpInfo.tpHandle; + importInfo.in.ub.expImportCfg.txPsn = endpointPsn_; + importInfo.in.ub.expImportCfg.rxPsn = endpointPsn_; + importInfo.in.ub.tpType = TILEXR_CCU_HCCP_TP_TYPE_RTP; + + ret = loader_.RaCtxQpImport(localResourceWindow_.raCtxHandle, &importInfo, &endpointRemoteQpHandle_); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute qpImport ret=" << ret + << " remoteQp=" << endpointRemoteQpHandle_ + << " tpn=0x" << std::hex << importInfo.out.ub.tpn + << std::dec << std::endl; + } + if (ret != 0 || endpointRemoteQpHandle_ == nullptr) { + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_MKIRT; + } + ++endpointPsn_; + + TileXRCcuLowerLayerTransportRoute collected; + collected.remoteEid = localResourceWindow_.eid; + collected.tpn = importInfo.out.ub.tpn; + collected.doorbellVa = qpInfo.ub.dbAddr; + collected.doorbellTokenId = qpInfo.ub.dbTokenId >> TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT; + collected.doorbellTokenValue = localResourceWindow_.tokenValue; + collected.sqDepth = sqDepth; + collected.endpointRouteVerified = true; + if (!HasCompleteEndpointRoute(collected)) { + TraceEndpointRouteStep("collected route incomplete"); + ReleaseLocalEndpointRoute(); + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute verified" + << " tpn=0x" << std::hex << collected.tpn + << " doorbellVa=0x" << collected.doorbellVa + << " doorbellTokenId=0x" << collected.doorbellTokenId + << std::dec + << " sqDepth=" << collected.sqDepth << std::endl; + } + *route = collected; + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::RefreshLocalVerifiedEndpointRoute(TileXRCcuDirectRuntimeReport* report) +{ + ResetReport(report); + localVerifiedEndpointRoute_ = TileXRCcuLowerLayerTransportRoute{}; + localVerifiedEndpointRouteValid_ = false; + if (!resourceWindowRegistered_) { + return Fail(report, "direct CCU resource window is not registered for endpoint route collection", + TILEXR_ERROR_NOT_INITIALIZED); + } + + TileXRCcuLowerLayerTransportRoute route; + int ret = TILEXR_ERROR_NOT_FOUND; + if (options_.localEndpointRouteCollector != nullptr) { + ret = options_.localEndpointRouteCollector( + devicePhyId_, + localResourceWindow_, + &route, + options_.localEndpointRouteCollectorUserData); + if (ret != TILEXR_SUCCESS) { + return Fail(report, "direct CCU local endpoint route collector failed", ret); + } + } else if (IsRaCtxLoopEndpointRouteCollectionMode()) { + ret = CollectLocalEndpointRouteWithRaCtx(&route); + if (ret != TILEXR_SUCCESS) { + return Fail(report, "direct CCU RA ctx loop endpoint route collection failed", ret); + } + } else { + TileXRCcuHccpLoaderReport providerReport; + ret = loader_.LoadEndpointRouteProviderFromEnv(&providerReport); + if (ret == TILEXR_SUCCESS && loader_.CollectLocalEndpointRoute != nullptr) { + TileXRCcuEndpointRouteProviderResourceWindow providerWindow; + FillProviderResourceWindow(localResourceWindow_, &providerWindow); + TileXRCcuEndpointRouteProviderRoute providerRoute; + ret = loader_.CollectLocalEndpointRoute(devicePhyId_, &providerWindow, &providerRoute); + if (ret != TILEXR_SUCCESS) { + return Fail(report, "direct CCU endpoint route provider failed", ret); + } + CopyProviderRoute(providerRoute, &route); + } else if (ret == TILEXR_ERROR_NOT_FOUND && !providerReport.endpointRouteProviderConfigured) { + const bool canCollectRaCtxRoute = + IsRaCtxResourceWindowRegistrationMode() && + localResourceWindow_.raCtxRegistered && + HasRaCtxEndpointRouteSymbols(loader_); + if (canCollectRaCtxRoute) { + ret = CollectLocalEndpointRouteWithRaCtx(&route); + if (ret != TILEXR_SUCCESS && ret != TILEXR_ERROR_NOT_FOUND) { + return Fail(report, "direct CCU RA ctx loop endpoint route collection failed", ret); + } + } + if (ret != TILEXR_SUCCESS) { + ret = CollectLocalEndpointRouteFromEnv(options_.rank, &route, report); + if (ret != TILEXR_SUCCESS) { + return ret; + } + } + } else { + return Fail(report, providerReport.message.empty() ? + "direct CCU endpoint route provider failed to load" : providerReport.message, ret); + } + } + + ret = ConfigureLocalVerifiedEndpointRoute(route); + if (ret != TILEXR_SUCCESS) { + return Fail(report, "direct CCU local endpoint route collector returned incomplete route", ret); + } + if (report != nullptr) { + report->initialized = initialized_; + report->logicDevId = static_cast(options_.devId); + report->devicePhyId = devicePhyId_; + report->message = "direct CCU local endpoint route collected"; + } + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::ConfigureLocalVerifiedEndpointRoute( + const TileXRCcuLowerLayerTransportRoute& route) +{ + if (!route.endpointRouteVerified || !HasCompleteEndpointRoute(route)) { + localVerifiedEndpointRoute_ = TileXRCcuLowerLayerTransportRoute{}; + localVerifiedEndpointRouteValid_ = false; + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + localVerifiedEndpointRoute_ = route; + localVerifiedEndpointRouteValid_ = true; + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::ExportLocalCcuRmaBuffer(TileXRCcuLocalResourceWindowInfo* info) const +{ + if (info == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *info = TileXRCcuLocalResourceWindowInfo{}; + if (!resourceWindowRegistered_) { + return TILEXR_ERROR_NOT_INITIALIZED; + } + *info = localResourceWindow_; + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::QueryTpHandleForPeer( + const std::array& peerEid, + uint64_t* tpHandle) +{ + if (tpHandle == nullptr || IsEmptyEid(localResourceWindow_.eid) || IsEmptyEid(peerEid) || + localResourceWindow_.raCtxHandle == nullptr || loader_.RaGetTpInfoListAsync == nullptr || + loader_.RaGetAsyncReqResult == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *tpHandle = 0; + + TileXRCcuHccpGetTpCfg tpCfg {}; + tpCfg.flag.value = 0; + tpCfg.flag.bs.rtp = 1; + tpCfg.transMode = TILEXR_CCU_HCCP_TRANSPORT_MODE_RM; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + tpCfg.localEid.raw[i] = localResourceWindow_.eid[i]; + tpCfg.peerEid.raw[i] = peerEid[i]; + } + + TileXRCcuHccpTpInfo tpInfo {}; + uint32_t tpInfoNum = 1; + void* reqHandle = nullptr; + const int ret = loader_.RaGetTpInfoListAsync( + localResourceWindow_.raCtxHandle, + &tpCfg, + &tpInfo, + &tpInfoNum, + &reqHandle); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute peerTpInfoAsync" + << " ret=" << ret + << " req=" << reqHandle + << " num=" << tpInfoNum + << " tpHandle=0x" << std::hex << tpInfo.tpHandle + << std::dec << std::endl; + } + if (ret != 0 || reqHandle == nullptr || tpInfoNum == 0) { + return TILEXR_ERROR_MKIRT; + } + const int waitRet = WaitRaCtxAsyncRequest(loader_, reqHandle); + if (waitRet != TILEXR_SUCCESS) { + return waitRet; + } + if (tpInfo.tpHandle == 0) { + return TILEXR_ERROR_NOT_FOUND; + } + *tpHandle = tpInfo.tpHandle; + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::ImportPeerEndpointRoute( + const TileXRCcuResourceWindowExchange& peerWindow, + uint64_t localTpHandle, + uint64_t peerTpHandle, + uint32_t localPsn, + uint32_t peerPsn, + TileXRCcuLowerLayerTransportRoute* importedRoute) +{ + if (importedRoute == nullptr || peerWindow.qpKey.size == 0 || localTpHandle == 0 || peerTpHandle == 0 || + localResourceWindow_.raCtxHandle == nullptr || loader_.RaCtxQpImport == nullptr || + !localVerifiedEndpointRouteValid_ || !HasCompleteEndpointRoute(localVerifiedEndpointRoute_)) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *importedRoute = TileXRCcuLowerLayerTransportRoute{}; + + TileXRCcuHccpQpImportInfo importInfo {}; + importInfo.in.key = peerWindow.qpKey; + importInfo.in.ub.mode = TILEXR_CCU_HCCP_JETTY_IMPORT_MODE_EXP; + importInfo.in.ub.tokenValue = peerWindow.tokenValue; + importInfo.in.ub.policy = TILEXR_CCU_HCCP_JETTY_GRP_POLICY_RR; + importInfo.in.ub.type = TILEXR_CCU_HCCP_TARGET_TYPE_JETTY; + importInfo.in.ub.flag.value = 0; + importInfo.in.ub.flag.bs.tokenPolicy = TILEXR_CCU_HCCP_TOKEN_POLICY_PLAIN_TEXT; + importInfo.in.ub.expImportCfg.tpHandle = localTpHandle; + importInfo.in.ub.expImportCfg.peerTpHandle = peerTpHandle; + importInfo.in.ub.expImportCfg.txPsn = localPsn; + importInfo.in.ub.expImportCfg.rxPsn = peerPsn; + importInfo.in.ub.tpType = TILEXR_CCU_HCCP_TP_TYPE_RTP; + + void* remoteQpHandle = nullptr; + const int ret = loader_.RaCtxQpImport(localResourceWindow_.raCtxHandle, &importInfo, &remoteQpHandle); + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute peerQpImport" + << " ret=" << ret + << " remoteQp=" << remoteQpHandle + << " localTpHandle=0x" << std::hex << localTpHandle + << " peerTpHandle=0x" << peerTpHandle + << " tpn=0x" << importInfo.out.ub.tpn + << std::dec + << " localPsn=" << localPsn + << " peerPsn=" << peerPsn + << " peerEid=" << FormatEndpointEid(peerWindow.remoteEid) + << " channelEid=" << FormatEndpointEid(ReverseEndpointEid(peerWindow.remoteEid)) + << std::endl; + } + if (ret != 0 || remoteQpHandle == nullptr) { + return TILEXR_ERROR_MKIRT; + } + endpointPeerRemoteQpHandles_.push_back(remoteQpHandle); + + importedRoute->remoteEid = ReverseEndpointEid(peerWindow.remoteEid); + importedRoute->tpn = importInfo.out.ub.tpn; + importedRoute->doorbellVa = localVerifiedEndpointRoute_.doorbellVa; + importedRoute->doorbellTokenId = localVerifiedEndpointRoute_.doorbellTokenId; + importedRoute->doorbellTokenValue = localVerifiedEndpointRoute_.doorbellTokenValue; + importedRoute->sqDepth = localVerifiedEndpointRoute_.sqDepth; + importedRoute->endpointRouteVerified = true; + return HasCompleteEndpointRoute(*importedRoute) ? TILEXR_SUCCESS : TILEXR_ERROR_PARA_CHECK_FAIL; +} + +int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector* buffers) +{ + if (buffers == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + buffers->clear(); + if (!resourceWindowRegistered_) { + return TILEXR_ERROR_NOT_INITIALIZED; + } + if (options_.rankSize <= 1) { + return TILEXR_SUCCESS; + } + if (options_.rank < 0 || options_.rank >= options_.rankSize || + options_.rankSize > TILEXR_CCU_DIRECT_MAX_RANK_SIZE || + options_.allGather == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + ReleasePeerEndpointImports(); + + TileXRCcuResourceWindowExchange local { + localResourceWindow_.addr, + localResourceWindow_.bytes, + localResourceWindow_.tokenId, + localResourceWindow_.rawTokenId, + localResourceWindow_.tokenValue, + }; + if (localVerifiedEndpointRouteValid_) { + local.remoteEid = localVerifiedEndpointRoute_.remoteEid; + local.tpn = localVerifiedEndpointRoute_.tpn; + local.doorbellVa = localVerifiedEndpointRoute_.doorbellVa; + local.doorbellTokenId = localVerifiedEndpointRoute_.doorbellTokenId; + local.doorbellTokenValue = localVerifiedEndpointRoute_.doorbellTokenValue; + local.sqDepth = localVerifiedEndpointRoute_.sqDepth; + local.qpKey = endpointQpKey_; + local.psn = endpointPsn_; + local.endpointRouteVerified = true; + local.channelResourceOwnerVerified = localVerifiedEndpointRoute_.channelResourceOwnerVerified; + local.transportResourceExchangeVerified = localVerifiedEndpointRoute_.transportResourceExchangeVerified; + } + std::vector all(static_cast(options_.rankSize)); + const int exportRet = options_.allGather( + &local, + sizeof(local), + all.data(), + options_.allGatherUserData); + if (exportRet != TILEXR_SUCCESS) { + return exportRet; + } + + const bool useImportedPeerRoute = UseImportedPeerEndpointRoute(); + const bool canImportPeerRoutes = + useImportedPeerRoute && + endpointQpKeyValid_ && + localVerifiedEndpointRouteValid_ && + HasCompleteEndpointRoute(localVerifiedEndpointRoute_) && + localResourceWindow_.raCtxHandle != nullptr && + loader_.RaCtxQpImport != nullptr && + loader_.RaGetTpInfoListAsync != nullptr && + loader_.RaGetAsyncReqResult != nullptr; + TileXRCcuEndpointTpHandleExchange localTpHandles {}; + std::vector allTpHandles; + bool peerTpHandlesReady = false; + if (canImportPeerRoutes) { + localTpHandles.psn = endpointPsn_; + for (int peer = 0; peer < options_.rankSize; ++peer) { + if (peer == options_.rank) { + continue; + } + const auto& peerWindow = all[peer]; + if (peerWindow.addr == 0 || IsEmptyEid(peerWindow.remoteEid)) { + continue; + } + uint64_t tpHandle = 0; + if (QueryTpHandleForPeer(peerWindow.remoteEid, &tpHandle) == TILEXR_SUCCESS) { + localTpHandles.tpHandles[peer] = tpHandle; + } + } + allTpHandles.resize(static_cast(options_.rankSize)); + const int tpExchangeRet = options_.allGather( + &localTpHandles, + sizeof(localTpHandles), + allTpHandles.data(), + options_.allGatherUserData); + peerTpHandlesReady = tpExchangeRet == TILEXR_SUCCESS; + } + + const uint64_t remoteCcuVaOffset = SelectRemoteCcuVaOffset(); + buffers->reserve(static_cast(options_.rankSize - 1)); + for (int peer = 0; peer < options_.rankSize; ++peer) { + if (peer == options_.rank) { + continue; + } + const auto& peerWindow = all[peer]; + if (peerWindow.addr == 0) { + buffers->clear(); + return TILEXR_ERROR_NOT_FOUND; + } + TileXRCcuRemoteCcuBufferInfo remote; + if (remoteCcuVaOffset > std::numeric_limits::max() - peerWindow.addr) { + buffers->clear(); + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + remote.remoteCcuVa = peerWindow.addr + remoteCcuVaOffset; + remote.peerRank = static_cast(peer); + remote.memoryTokenId = peerWindow.tokenId; + remote.rawMemoryTokenId = peerWindow.rawTokenId; + remote.memoryTokenValue = peerWindow.tokenValue; + remote.remoteEid = ReverseEndpointEid(peerWindow.remoteEid); + TileXRCcuLowerLayerTransportRoute importedRoute; + const uint64_t localTpForPeer = peerTpHandlesReady ? + localTpHandles.tpHandles[peer] : 0; + const uint64_t peerTpForLocal = peerTpHandlesReady ? + allTpHandles[static_cast(peer)].tpHandles[options_.rank] : 0; + const bool importedPeerRoute = + peerTpHandlesReady && + peerTpForLocal != 0 && + localTpForPeer != 0 && + ImportPeerEndpointRoute( + peerWindow, + localTpForPeer, + peerTpForLocal, + localTpHandles.psn, + allTpHandles[static_cast(peer)].psn, + &importedRoute) == TILEXR_SUCCESS; + if (importedPeerRoute) { + remote.remoteEid = importedRoute.remoteEid; + remote.tpn = importedRoute.tpn; + remote.doorbellVa = importedRoute.doorbellVa; + remote.doorbellTokenId = importedRoute.doorbellTokenId; + remote.doorbellTokenValue = importedRoute.doorbellTokenValue; + remote.sqDepth = importedRoute.sqDepth; + } else { + remote.tpn = peerWindow.tpn; + remote.doorbellVa = peerWindow.doorbellVa; + remote.doorbellTokenId = peerWindow.doorbellTokenId; + remote.doorbellTokenValue = peerWindow.doorbellTokenValue; + remote.sqDepth = peerWindow.sqDepth; + if (!useImportedPeerRoute && TraceEndpointRoute() && peerWindow.endpointRouteVerified) { + std::cerr << "TileXRDirectCcuTrace endpointRoute usePeerExportedRoute" + << " peerRank=" << peer + << " tpn=0x" << std::hex << peerWindow.tpn + << " peerEid=" << FormatEndpointEid(peerWindow.remoteEid) + << " channelEid=" << FormatEndpointEid(remote.remoteEid) + << " remoteCcuVa=0x" << remote.remoteCcuVa + << " remoteCcuVaBase=0x" << peerWindow.addr + << " remoteCcuVaOffset=0x" << remoteCcuVaOffset + << std::dec << std::endl; + } + } + if (localVerifiedEndpointRouteValid_) { + remote.localDoorbellVa = localVerifiedEndpointRoute_.doorbellVa; + remote.localDoorbellTokenId = localVerifiedEndpointRoute_.doorbellTokenId; + remote.localDoorbellTokenValue = localVerifiedEndpointRoute_.doorbellTokenValue; + remote.localSqDepth = localVerifiedEndpointRoute_.sqDepth; + } + remote.endpointRouteVerified = peerWindow.endpointRouteVerified; + if (importedPeerRoute) { + remote.endpointRouteVerified = true; + } + remote.channelResourceOwnerVerified = peerWindow.channelResourceOwnerVerified; + remote.transportResourceExchangeVerified = peerWindow.transportResourceExchangeVerified; + if (TraceEndpointRoute()) { + std::cerr << "TileXRDirectCcuTrace endpointRoute channelRoute" + << " peerRank=" << peer + << " importedPeerRoute=" << (importedPeerRoute ? 1 : 0) + << " peerTpHandlesReady=" << (peerTpHandlesReady ? 1 : 0) + << " localEid=" << FormatEndpointEid(localResourceWindow_.eid) + << " peerEid=" << FormatEndpointEid(peerWindow.remoteEid) + << " selectedRemoteEid=" << FormatEndpointEid(remote.remoteEid) + << " localTpHandle=0x" << std::hex << localTpForPeer + << " peerTpHandle=0x" << peerTpForLocal + << " tpn=0x" << remote.tpn + << " remoteCcuVa=0x" << remote.remoteCcuVa + << " remoteCcuVaBase=0x" << peerWindow.addr + << " remoteCcuVaOffset=0x" << remoteCcuVaOffset + << " doorbellVa=0x" << remote.doorbellVa + << " localDoorbellVa=0x" << remote.localDoorbellVa + << " memoryTokenId=0x" << remote.memoryTokenId + << " rawMemoryTokenId=0x" << remote.rawMemoryTokenId + << " memoryTokenValue=0x" << remote.memoryTokenValue + << std::dec + << " endpointRouteVerified=" << (remote.endpointRouteVerified ? 1 : 0) + << " channelResourceOwnerVerified=" << (remote.channelResourceOwnerVerified ? 1 : 0) + << " transportResourceExchangeVerified=" + << (remote.transportResourceExchangeVerified ? 1 : 0) + << std::endl; + } + buffers->push_back(remote); + } + if (peerTpHandlesReady) { + ++endpointPsn_; + } + return TILEXR_SUCCESS; +} + +int TileXRCcuDirectRuntime::ExportLowerLayerTransportSnapshot( + const TileXRCcuLowerLayerTransportSnapshot& templateSnapshot, + TileXRCcuLowerLayerTransportSnapshot* snapshot) const +{ + if (snapshot == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *snapshot = TileXRCcuLowerLayerTransportSnapshot{}; + if (!resourceWindowRegistered_) { + return TILEXR_ERROR_NOT_INITIALIZED; + } + + *snapshot = templateSnapshot; + for (uint32_t i = 0; i < snapshot->routes.size(); ++i) { + auto& route = snapshot->routes[i]; + const bool preserveVerifiedRoute = route.endpointRouteVerified && HasCompleteEndpointRoute(route); + bool filledSyntheticEndpointField = false; + if (IsEmptyEid(route.remoteEid)) { + route.remoteEid[0] = static_cast(devicePhyId_ & 0xffU); + route.remoteEid[1] = static_cast((devicePhyId_ >> 8U) & 0xffU); + route.remoteEid[2] = static_cast(route.peerRank & 0xffU); + route.remoteEid[3] = static_cast((route.peerRank >> 8U) & 0xffU); + route.remoteEid[4] = static_cast(options_.rank & 0xff); + route.remoteEid[5] = static_cast(snapshot->dieId); + route.remoteEid[6] = static_cast(route.channelId & 0xffU); + route.remoteEid[7] = static_cast((route.channelId >> 8U) & 0xffU); + filledSyntheticEndpointField = true; + } + if (!preserveVerifiedRoute && route.tpn == 0) { + route.tpn = route.channelId + 1U; + filledSyntheticEndpointField = true; + } + if (route.doorbellVa == 0) { + route.doorbellVa = localResourceWindow_.addr + TILEXR_CCU_V1_XN_RESOURCE_OFFSET + + static_cast(snapshot->xnStartId + i) * TILEXR_CCU_XN_SLOT_BYTES; + filledSyntheticEndpointField = true; + } + if (route.doorbellTokenId == 0) { + route.doorbellTokenId = localResourceWindow_.tokenId; + filledSyntheticEndpointField = true; + } + if (!preserveVerifiedRoute && route.doorbellTokenValue == 0) { + route.doorbellTokenValue = localResourceWindow_.tokenValue; + } + if (route.sqDepth == 0) { + route.sqDepth = TILEXR_CCU_DEFAULT_DIRECT_SQ_DEPTH; + filledSyntheticEndpointField = true; + } + route.endpointRouteVerified = (preserveVerifiedRoute && !filledSyntheticEndpointField) || + (filledSyntheticEndpointField && TrustSyntheticEndpointRouteForDiagnostics() && + HasCompleteEndpointRoute(route)); + } + return TILEXR_SUCCESS; +} + +} // namespace TileXR diff --git a/src/comm/ccu/tilexr_ccu_direct_runtime.h b/src/comm/ccu/tilexr_ccu_direct_runtime.h new file mode 100644 index 00000000..3e358f06 --- /dev/null +++ b/src/comm/ccu/tilexr_ccu_direct_runtime.h @@ -0,0 +1,152 @@ +/* + * Copyright (c) 2026 TileXR Project + * Licensed under the Apache License, Version 2.0 + */ + +#ifndef TILEXR_CCU_DIRECT_RUNTIME_H +#define TILEXR_CCU_DIRECT_RUNTIME_H + +#include "ccu/tilexr_ccu_driver_adapter.h" +#include "ccu/tilexr_ccu_hccp_loader.h" +#include "ccu/tilexr_ccu_lower_layer_plan_builder.h" +#include "ccu/tilexr_ccu_ra_custom_channel_provider.h" + +#include +#include +#include +#include +#include + +namespace TileXR { + +using TileXRCcuDirectAllGatherFn = int (*)( + const void* sendBuf, + size_t sendBytes, + void* recvBuf, + void* userData); + +struct TileXRCcuLocalResourceWindowInfo; + +using TileXRCcuLocalEndpointRouteCollectorFn = int (*)( + uint32_t devicePhyId, + const TileXRCcuLocalResourceWindowInfo& localResourceWindow, + TileXRCcuLowerLayerTransportRoute* route, + void* userData); + +struct TileXRCcuDirectRuntimeOptions { + int rank = 0; + int rankSize = 0; + int devId = 0; + TileXRCcuDirectAllGatherFn allGather = nullptr; + void* allGatherUserData = nullptr; + TileXRCcuLocalEndpointRouteCollectorFn localEndpointRouteCollector = nullptr; + void* localEndpointRouteCollectorUserData = nullptr; +}; + +struct TileXRCcuLocalResourceWindowInfo { + uint64_t addr = 0; + uint64_t bytes = 0; + uint32_t tokenId = 0; + uint32_t rawTokenId = 0; + uint32_t tokenValue = 0; + uint64_t targetSegHandle = 0; + void* raCtxHandle = nullptr; + void* tokenIdHandle = nullptr; + void* lmemHandle = nullptr; + std::array eid {}; + uint32_t eidIndex = 0; + bool raCtxRegistered = false; +}; + +struct TileXRCcuResourceWindowExchange { + uint64_t addr = 0; + uint64_t bytes = 0; + uint32_t tokenId = 0; + uint32_t rawTokenId = 0; + uint32_t tokenValue = 0; + std::array remoteEid {}; + uint32_t tpn = 0; + uint64_t doorbellVa = 0; + uint32_t doorbellTokenId = 0; + uint32_t doorbellTokenValue = 0; + uint32_t sqDepth = 0; + TileXRCcuHccpQpKey qpKey {}; + uint32_t psn = 0; + bool endpointRouteVerified = false; + bool channelResourceOwnerVerified = false; + bool transportResourceExchangeVerified = false; +}; + +struct TileXRCcuDirectRuntimeReport { + bool initialized = false; + bool raInitialized = false; + bool ccuTlvInitialized = false; + uint32_t logicDevId = 0; + uint32_t devicePhyId = 0; + int hdcType = 0; + std::string message; +}; + +class TileXRCcuDirectRuntime { +public: + int Init(const TileXRCcuDirectRuntimeOptions& options, TileXRCcuDirectRuntimeReport* report); + void Shutdown(); + bool IsAvailable() const; + + int QueryBasicInfo(uint8_t dieId, TileXRCcuBasicInfo* basicInfo, TileXRCcuDriverAdapterReport* report); + int CreateDriverAdapter(TileXRCcuDriverAdapter* adapter, TileXRCcuDriverAdapterReport* report); + int RegisterCcuResourceRmaBuffer(uint64_t resourceAddr); + int RefreshLocalVerifiedEndpointRoute(TileXRCcuDirectRuntimeReport* report); + int ConfigureLocalVerifiedEndpointRoute(const TileXRCcuLowerLayerTransportRoute& route); + int ExportLocalCcuRmaBuffer(TileXRCcuLocalResourceWindowInfo* info) const; + int ExportRemoteCcuRmaBuffers(std::vector* buffers); + int ExportLowerLayerTransportSnapshot( + const TileXRCcuLowerLayerTransportSnapshot& templateSnapshot, + TileXRCcuLowerLayerTransportSnapshot* snapshot) const; + +private: + int ResolveDevicePhyId(uint32_t* devicePhyId, TileXRCcuDirectRuntimeReport* report) const; + int RegisterCcuResourceRmaBufferWithRaCtx(uint64_t resourceAddr, uint64_t resourceBytes); + int CollectLocalEndpointRouteWithRaCtx(TileXRCcuLowerLayerTransportRoute* route); + int CollectLocalEndpointRouteWithRaCtxOnce( + TileXRCcuLowerLayerTransportRoute* route, + bool* asyncWaitFailed); + int QueryTpHandleForPeer( + const std::array& peerEid, + uint64_t* tpHandle); + int ImportPeerEndpointRoute( + const TileXRCcuResourceWindowExchange& peerWindow, + uint64_t localTpHandle, + uint64_t peerTpHandle, + uint32_t localPsn, + uint32_t peerPsn, + TileXRCcuLowerLayerTransportRoute* importedRoute); + void ReleasePeerEndpointImports(); + void ReleaseRegisteredResourceWindow(); + void ReleaseLocalEndpointRoute(); + + TileXRCcuDirectRuntimeOptions options_; + TileXRCcuHccpLoader loader_; + TileXRCcuRaCustomChannelProvider raCustomChannelProvider_; + TileXRCcuBasicInfo cachedBasicInfo_ = {}; + TileXRCcuLocalResourceWindowInfo localResourceWindow_ = {}; + TileXRCcuLowerLayerTransportRoute localVerifiedEndpointRoute_ = {}; + bool cachedBasicInfoValid_ = false; + bool resourceWindowRegistered_ = false; + bool localVerifiedEndpointRouteValid_ = false; + void* endpointChanHandle_ = nullptr; + void* endpointCqHandle_ = nullptr; + void* endpointQpHandle_ = nullptr; + void* endpointRemoteQpHandle_ = nullptr; + std::vector endpointPeerRemoteQpHandles_; + TileXRCcuHccpQpKey endpointQpKey_ = {}; + bool endpointQpKeyValid_ = false; + bool endpointRouteBound_ = false; + uint32_t endpointPsn_ = 1; + uint32_t devicePhyId_ = 0; + bool initialized_ = false; +}; + +} // namespace TileXR + +#endif // TILEXR_CCU_DIRECT_RUNTIME_H diff --git a/tests/ccu/test_tilexr_ccu_direct_orchestrator.py b/tests/ccu/test_tilexr_ccu_direct_orchestrator.py new file mode 100644 index 00000000..d45dc384 --- /dev/null +++ b/tests/ccu/test_tilexr_ccu_direct_orchestrator.py @@ -0,0 +1,1724 @@ +#!/usr/bin/env python3 +# +# Copyright (c) 2026 TileXR Project +# + +import os +import platform +import shutil +import subprocess +import tempfile +import textwrap +import unittest +from pathlib import Path + + +REPO_ROOT = Path(__file__).resolve().parents[2] +DIRECT_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_direct_orchestrator.h" +DIRECT_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_direct_orchestrator.cpp" +INSTALL_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_install_provider.cpp" +PROVIDER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_provider.cpp" +PACKAGE_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_launch_package.cpp" +REPOSITORY_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_repository.cpp" +DRIVER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp" +SPECS_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_specs.cpp" +ALLOCATOR_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_resource_allocator.cpp" +PRODUCER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_producer_plan.cpp" +BARRIER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_barrier_program.cpp" +MICROCODE_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_microcode.cpp" +MEMORY_PROGRAM_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_memory_program.cpp" +RUNTIME_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_runtime.cpp" +LOWER_LAYER_PLAN_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_lower_layer_plan_builder.cpp" +LOWER_LAYER_PAYLOAD_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_lower_layer_payloads.cpp" +COMM_CMAKE = REPO_ROOT / "src" / "comm" / "CMakeLists.txt" +INCLUDE_DIR = REPO_ROOT / "src" / "include" +COMM_DIR = REPO_ROOT / "src" / "comm" + + +PRIVATE_CCU_PRODUCER_NEEDLES = [ + "#include + #include + #include + #include + + using namespace TileXR; + + struct ObservedCall { + uint32_t op = 0; + uint32_t die = 0; + uint32_t offset = 0; + uint32_t dataLen = 0; + uint32_t arraySize = 0; + }; + + struct FakeState { + std::vector deviceBytes; + std::vector calls; + bool freed = false; + }; + + int FakeAlloc(uint64_t bytes, void** ptr, void* userData) + { + auto* state = static_cast(userData); + state->deviceBytes.assign(static_cast(bytes), 0); + *ptr = state->deviceBytes.data(); + return 0; + } + + int FakeCopy(void* dst, uint64_t dstBytes, const void* src, uint64_t bytes, void* userData) + { + auto* state = static_cast(userData); + if (dst != state->deviceBytes.data() || dstBytes != state->deviceBytes.size() || + bytes != state->deviceBytes.size()) { + return -1; + } + std::memcpy(dst, src, static_cast(bytes)); + return 0; + } + + int FakeFree(void* ptr, void* userData) + { + auto* state = static_cast(userData); + if (ptr != state->deviceBytes.data()) { + return -1; + } + state->freed = true; + state->deviceBytes.clear(); + return 0; + } + + int FakeCustomChannel( + uint32_t, + const TileXRCcuCustomChannelIn& in, + TileXRCcuCustomChannelOut* out, + void* userData) + { + auto* state = static_cast(userData); + ObservedCall call; + call.op = in.op; + call.die = in.data.dataInfo.udieIdx; + call.offset = in.offsetStartIdx; + call.dataLen = in.data.dataInfo.dataLen; + call.arraySize = in.data.dataInfo.dataArraySize; + state->calls.push_back(call); + out->opRet = 0; + return 0; + } + + TileXRCcuLowerLayerInstallPlan MakeLowerLayer() + { + TileXRCcuLowerLayerInstallPlan lowerLayer; + lowerLayer.msidTokens.push_back({1, 0x45, 0x1234, 0x5678}); + + TileXRCcuPfeInstall pfe; + pfe.dieId = 1; + pfe.pfeOffset = 7; + pfe.ctx.raw[0] = 0xa1; + lowerLayer.pfes.push_back(pfe); + + TileXRCcuJettyInstall jetty; + jetty.dieId = 1; + jetty.startJettyCtxId = 9; + jetty.ctxs.resize(1); + jetty.ctxs[0].raw[0] = 0xb2; + lowerLayer.jettys.push_back(jetty); + + TileXRCcuChannelInstall channel; + channel.dieId = 1; + channel.channelId = 11; + channel.ctx.raw[0] = 0xc3; + lowerLayer.channels.push_back(channel); + lowerLayer.xnClears.push_back({1, 32, 1}); + lowerLayer.ckeClears.push_back({1, 16, 1}); + return lowerLayer; + } + + int main() + { + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.sqeArgCount = TILEXR_CCU_SQE_ARGS_LEN; + options.syncResourceCount = 3; + options.syncInstructionCount = 143; + options.bindingsPerSyncResource = 3; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-direct-install-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.ckeStartId = 332; + options.channelStartId = 2; + + TileXRCcuLowerLayerInstallPlan lowerLayer = MakeLowerLayer(); + options.lowerLayerPlan = &lowerLayer; + + FakeState state; + TileXRCcuDriverAdapter adapter; + TileXRCcuDriverAdapterReport driverReport; + if (adapter.Init(5, FakeCustomChannel, &state, &driverReport) != TILEXR_SUCCESS) { + std::cerr << "adapter init failed\n"; + return 1; + } + TileXRCcuDeviceMemoryOps memoryOps; + memoryOps.alloc = FakeAlloc; + memoryOps.copyHostToDevice = FakeCopy; + memoryOps.free = FakeFree; + options.driverAdapter = &adapter; + options.repositoryMemoryOps = memoryOps; + options.repositoryMemoryUserData = &state; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + if (ret != TILEXR_ERROR_NOT_FOUND) { + std::cerr << "direct install attempt should stop at known missing surfaces: " + << report.message << "\n"; + return 2; + } + if (!report.pipelineBuilt || !report.installAttempted || report.submitReady || + report.submitTaskCount != 0 || report.missingInstallSurfaceCount != 2 || + report.publicVerifiedInstallSurfaceCount != 4) { + std::cerr << "direct install report mismatch\n"; + return 3; + } + if (report.message.find("remote XN install provider is missing") == + std::string::npos) { + std::cerr << "direct install diagnostic is weak: " << report.message << "\n"; + return 4; + } + if (attempt.package.tasks.size() != 2 || attempt.manifest.requirements.size() != 6 || + attempt.allocation.receiptId == 0 || + attempt.repositoryReceipt.instructionStartId != attempt.package.repository.missionStartId || + attempt.repositoryReceipt.instructionCount != attempt.package.repository.missionCount || + attempt.package.tasks[0].instStartId != attempt.package.repository.missionStartId || + attempt.package.tasks[0].instCnt != 13 || + attempt.package.tasks[1].instStartId != + attempt.package.repository.missionStartId + attempt.package.repository.sqeLoadCount || + attempt.package.tasks[1].instCnt != + attempt.package.repository.missionCount - attempt.package.repository.sqeLoadCount) { + std::cerr << "attempt artifacts mismatch" + << " taskSize=" << attempt.package.tasks.size() + << " reqSize=" << attempt.manifest.requirements.size() + << " receipt=" << attempt.allocation.receiptId + << " receiptStart=" << attempt.repositoryReceipt.instructionStartId + << " missionStart=" << attempt.package.repository.missionStartId + << " receiptCount=" << attempt.repositoryReceipt.instructionCount + << " missionCount=" << attempt.package.repository.missionCount + << " task0Start=" << attempt.package.tasks[0].instStartId + << " task0Cnt=" << attempt.package.tasks[0].instCnt + << " task1Start=" << attempt.package.tasks[1].instStartId + << " expectedTask1Start=" + << attempt.package.repository.missionStartId + attempt.package.repository.sqeLoadCount + << " task1Cnt=" << attempt.package.tasks[1].instCnt + << " expectedTask1Cnt=" + << attempt.package.repository.missionCount - attempt.package.repository.sqeLoadCount + << "\n"; + return 5; + } + if (!attempt.evidence.repositoryInstalled || !attempt.evidence.missionInstalled || + !attempt.evidence.localXnInstalled || + !attempt.evidence.notifyCkeInstalled || + attempt.evidence.channelBindingsInstalled || + attempt.evidence.remoteXnBound) { + std::cerr << "attempt evidence mismatch\n"; + return 6; + } + if (state.calls.size() != 7 || + state.calls[0].op != TILEXR_CCU_U_OP_SET_MSID_TOKEN || + state.calls[1].op != TILEXR_CCU_U_OP_SET_PFE || + state.calls[2].op != TILEXR_CCU_U_OP_SET_JETTY_CTX || + state.calls[3].op != TILEXR_CCU_U_OP_SET_CHANNEL || + state.calls[4].op != TILEXR_CCU_U_OP_SET_XN || + state.calls[5].op != TILEXR_CCU_U_OP_SET_CKE || + state.calls[6].op != TILEXR_CCU_U_OP_SET_INSTRUCTION) { + std::cerr << "unexpected direct install call sequence\n"; + return 7; + } + if (state.calls[6].offset != attempt.package.repository.missionStartId || + state.calls[6].dataLen != + attempt.package.repository.missionCount * sizeof(TileXRCcuInstr)) { + std::cerr << "repository install range mismatch\n"; + return 8; + } + if (state.calls[0].offset != 0 || state.calls[0].dataLen != 0 || + state.calls[0].arraySize != 0) { + std::cerr << "MSID token install envelope mismatch\n"; + return 10; + } + if (TileXRCcuReleaseRepositoryInstallReceipt( + attempt.repositoryReceipt, memoryOps, &state, &attempt.repositoryReleaseReport) != + TILEXR_SUCCESS || !state.freed) { + std::cerr << "repository release failed: " << attempt.repositoryReleaseReport.message << "\n"; + return 9; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_install_attempt_becomes_submit_ready_with_remote_xn_peer_exchange_proof(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + #include + + using namespace TileXR; + + struct ObservedCall { + uint32_t op = 0; + uint32_t offset = 0; + uint32_t arraySize = 0; + }; + + struct FakeState { + std::vector deviceBytes; + std::vector calls; + bool freed = false; + }; + + int FakeAlloc(uint64_t bytes, void** ptr, void* userData) + { + auto* state = static_cast(userData); + state->deviceBytes.assign(static_cast(bytes), 0); + *ptr = state->deviceBytes.data(); + return 0; + } + + int FakeCopy(void* dst, uint64_t dstBytes, const void* src, uint64_t bytes, void* userData) + { + auto* state = static_cast(userData); + if (dst != state->deviceBytes.data() || dstBytes != state->deviceBytes.size() || + bytes != state->deviceBytes.size()) { + return -1; + } + std::memcpy(dst, src, static_cast(bytes)); + return 0; + } + + int FakeFree(void* ptr, void* userData) + { + auto* state = static_cast(userData); + if (ptr != state->deviceBytes.data()) { + return -1; + } + state->freed = true; + state->deviceBytes.clear(); + return 0; + } + + int FakeCustomChannel( + uint32_t, + const TileXRCcuCustomChannelIn& in, + TileXRCcuCustomChannelOut* out, + void* userData) + { + auto* state = static_cast(userData); + state->calls.push_back({in.op, in.offsetStartIdx, in.data.dataInfo.dataArraySize}); + out->opRet = 0; + return 0; + } + + TileXRCcuLowerLayerInstallPlan MakeCompleteLowerLayer() + { + TileXRCcuLowerLayerInstallPlan lowerLayer; + lowerLayer.msidTokens.push_back({1, 0x45, 0x1234, 0x5678}); + + TileXRCcuPfeInstall pfe; + pfe.dieId = 1; + pfe.pfeOffset = 2; + pfe.ctx.raw[0] = 0xa1; + lowerLayer.pfes.push_back(pfe); + + TileXRCcuJettyInstall jetty; + jetty.dieId = 1; + jetty.startJettyCtxId = 9; + jetty.ctxs.resize(1); + jetty.ctxs[0].raw[0] = 0xb2; + lowerLayer.jettys.push_back(jetty); + + for (uint32_t i = 0; i < 3; ++i) { + TileXRCcuChannelInstall channel; + channel.dieId = 1; + channel.channelId = 2 + i; + channel.ctx.raw[0] = 0xc3 + i; + lowerLayer.channels.push_back(channel); + } + + lowerLayer.xnClears.push_back({1, 1961, 14}); + lowerLayer.ckeClears.push_back({1, 332, 3}); + lowerLayer.remoteXnBindings.push_back({1, 2, 1961, 1975, 332, 0, true, 0, true, true, true}); + lowerLayer.remoteXnBindings.push_back({1, 3, 1962, 1976, 333, 0, true, 0, true, true, true}); + lowerLayer.remoteXnBindings.push_back({1, 4, 1963, 1977, 334, 0, true, 0, true, true, true}); + return lowerLayer; + } + + int main() + { + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.sqeArgCount = TILEXR_CCU_SQE_ARGS_LEN; + options.syncResourceCount = 3; + options.syncInstructionCount = 143; + options.bindingsPerSyncResource = 3; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-direct-install-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.ckeStartId = 332; + options.channelStartId = 2; + TileXRCcuLowerLayerInstallPlan lowerLayer = MakeCompleteLowerLayer(); + options.lowerLayerPlan = &lowerLayer; + + FakeState state; + TileXRCcuDriverAdapter adapter; + TileXRCcuDriverAdapterReport driverReport; + if (adapter.Init(5, FakeCustomChannel, &state, &driverReport) != TILEXR_SUCCESS) { + std::cerr << "adapter init failed\n"; + return 1; + } + TileXRCcuDeviceMemoryOps memoryOps; + memoryOps.alloc = FakeAlloc; + memoryOps.copyHostToDevice = FakeCopy; + memoryOps.free = FakeFree; + options.driverAdapter = &adapter; + options.repositoryMemoryOps = memoryOps; + options.repositoryMemoryUserData = &state; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + if (ret != TILEXR_SUCCESS) { + std::cerr << "direct install attempt should be submit-ready: " + << report.message << "\n"; + return 2; + } + if (!report.pipelineBuilt || !report.installAttempted || !report.installSucceeded || + !report.submitReady || report.missingInstallSurfaceCount != 0 || + report.publicVerifiedInstallSurfaceCount != 6 || + report.submitTaskCount != attempt.package.tasks.size()) { + std::cerr << "submit-ready direct install report mismatch\n"; + return 3; + } + if (!attempt.evidence.repositoryInstalled || !attempt.evidence.missionInstalled || + !attempt.evidence.localXnInstalled || !attempt.evidence.remoteXnBound || + !attempt.evidence.notifyCkeInstalled || !attempt.evidence.channelBindingsInstalled) { + std::cerr << "submit-ready evidence mismatch\n"; + return 4; + } + if (attempt.evidence.remoteXnSource.source.find("ValidateRemoteXnExchangeBindingProof") == + std::string::npos || + attempt.evidence.remoteXnSource.detail.find("syncXn") == std::string::npos) { + std::cerr << "remote XN evidence detail is weak\n"; + return 5; + } + if (state.calls.size() != 10 || + state.calls[0].op != TILEXR_CCU_U_OP_SET_MSID_TOKEN || + state.calls[1].op != TILEXR_CCU_U_OP_SET_PFE || + state.calls[2].op != TILEXR_CCU_U_OP_SET_JETTY_CTX || + state.calls[3].op != TILEXR_CCU_U_OP_SET_CHANNEL || + state.calls[4].op != TILEXR_CCU_U_OP_SET_CHANNEL || + state.calls[5].op != TILEXR_CCU_U_OP_SET_CHANNEL || + state.calls[6].op != TILEXR_CCU_U_OP_SET_XN || + state.calls[6].arraySize != 8 || + state.calls[7].op != TILEXR_CCU_U_OP_SET_XN || + state.calls[7].arraySize != 6 || + state.calls[8].op != TILEXR_CCU_U_OP_SET_CKE || + state.calls[8].arraySize != 3 || + state.calls[9].op != TILEXR_CCU_U_OP_SET_INSTRUCTION) { + std::cerr << "submit-ready call sequence mismatch\n"; + return 6; + } + if (TileXRCcuReleaseRepositoryInstallReceipt( + attempt.repositoryReceipt, memoryOps, &state, &attempt.repositoryReleaseReport) != + TILEXR_SUCCESS || !state.freed) { + std::cerr << "repository release failed: " << attempt.repositoryReleaseReport.message << "\n"; + return 7; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_submit_tasks_runs_submitter_in_order_and_stops_on_failure(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + + using namespace TileXR; + + struct SubmitState { + std::vector starts; + int failOnCall = 0; + }; + + int FakeSubmit(const TileXRCcuTask& task, void* stream, void* userData) + { + if (stream != reinterpret_cast(0x1234)) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + auto* state = static_cast(userData); + state->starts.push_back(task.instStartId); + if (state->failOnCall != 0 && + state->starts.size() == static_cast(state->failOnCall)) { + return TILEXR_ERROR_MKIRT; + } + return TILEXR_SUCCESS; + } + + TileXRCcuTask Task(uint16_t start) + { + TileXRCcuTask task; + task.dieId = 1; + task.missionId = 6; + task.instStartId = start; + task.instCnt = 13; + task.key = 0x059b0f03U; + task.argSize = TILEXR_CCU_SQE_ARGS_LEN; + return task; + } + + int main() + { + std::vector tasks; + tasks.push_back(Task(489)); + tasks.push_back(Task(502)); + + SubmitState okState; + TileXRCcuDirectSubmitReport report; + if (TileXRCcuSubmitPreparedTasks( + tasks, reinterpret_cast(0x1234), FakeSubmit, &okState, &report) != + TILEXR_SUCCESS) { + std::cerr << "submit batch failed: " << report.message << "\n"; + return 1; + } + if (!report.submitted || report.taskCount != 2 || report.submittedTaskCount != 2 || + okState.starts.size() != 2 || okState.starts[0] != 489 || okState.starts[1] != 502 || + report.message.find("submitted") == std::string::npos) { + std::cerr << "submit success report mismatch\n"; + return 2; + } + + SubmitState failState; + failState.failOnCall = 2; + if (TileXRCcuSubmitPreparedTasks( + tasks, reinterpret_cast(0x1234), FakeSubmit, &failState, &report) != + TILEXR_ERROR_MKIRT) { + std::cerr << "submit failure was not propagated\n"; + return 3; + } + if (report.submitted || report.taskCount != 2 || report.submittedTaskCount != 1 || + failState.starts.size() != 2 || + report.message.find("task=1") == std::string::npos || + report.message.find("ret=-2") == std::string::npos || + report.message.find("missionId=6") == std::string::npos || + report.message.find("instStartId=502") == std::string::npos || + report.message.find("key=0x59b0f03") == std::string::npos || + report.message.find("argSize=13") == std::string::npos) { + std::cerr << "submit failure report mismatch: " << report.message << "\n"; + return 4; + } + + if (TileXRCcuSubmitPreparedTasks( + tasks, nullptr, FakeSubmit, &okState, &report) != TILEXR_ERROR_PARA_CHECK_FAIL) { + std::cerr << "null stream accepted\n"; + return 5; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_submit_default_runtime_trace_indexes_final_tasks(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + #include + #include + + #include + + using namespace TileXR; + + namespace { + int g_launchCount = 0; + } + + extern "C" rtError_t rtCCULaunch(rtCcuTaskInfo_t*, rtStream_t) + { + ++g_launchCount; + return RT_ERROR_NONE; + } + + TileXRCcuTask Task(uint16_t start) + { + TileXRCcuTask task; + task.dieId = 1; + task.missionId = 6; + task.timeout = 68; + task.instStartId = start; + task.instCnt = 13; + task.key = 0x059b0f03U; + task.argSize = TILEXR_CCU_SQE_ARGS_LEN; + task.args[0] = 0x1000ULL + start; + return task; + } + + int main() + { + setenv("TILEXR_CCU_DIRECT_TRACE", "1", 1); + + std::vector tasks; + tasks.push_back(Task(489)); + tasks.push_back(Task(502)); + + TileXRCcuDirectSubmitReport report; + const int ret = TileXRCcuSubmitPreparedTasks( + tasks, reinterpret_cast(0x1234), nullptr, nullptr, &report); + if (ret != TILEXR_SUCCESS || !report.submitted || g_launchCount != 2) { + std::cerr << "submit mismatch ret=" << ret + << " submitted=" << report.submitted + << " launchCount=" << g_launchCount + << " message=" << report.message << "\n"; + return 1; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + self.assertIn("TileXRDirectCcuTrace finalRuntimeTask[0] dieId=1 missionId=6 timeout=68 instStartId=489", result.stderr) + self.assertIn("TileXRDirectCcuTrace finalRuntimeTask[1] dieId=1 missionId=6 timeout=68 instStartId=502", result.stderr) + self.assertNotIn("TileXRDirectCcuTrace finalRuntimeTask task=0", result.stderr) + + def test_direct_install_attempt_can_prepare_lower_layer_plan_after_allocation(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + #include + + using namespace TileXR; + + struct FakeState { + std::vector deviceBytes; + std::vector ops; + }; + + struct CallbackState { + uint32_t callCount = 0; + uint32_t syncResourceCount = 0; + }; + + int FakeAlloc(uint64_t bytes, void** ptr, void* userData) + { + auto* state = static_cast(userData); + state->deviceBytes.assign(static_cast(bytes), 0); + *ptr = state->deviceBytes.data(); + return 0; + } + + int FakeCopy(void* dst, uint64_t dstBytes, const void* src, uint64_t bytes, void* userData) + { + auto* state = static_cast(userData); + if (dst != state->deviceBytes.data() || dstBytes != state->deviceBytes.size() || + bytes != state->deviceBytes.size()) { + return -1; + } + std::memcpy(dst, src, static_cast(bytes)); + return 0; + } + + int FakeFree(void*, void*) + { + return 0; + } + + int FakeCustomChannel( + uint32_t, + const TileXRCcuCustomChannelIn& in, + TileXRCcuCustomChannelOut* out, + void* userData) + { + auto* state = static_cast(userData); + state->ops.push_back(in.op); + out->opRet = 0; + return 0; + } + + int PrepareLowerLayerPlan( + const TileXRCcuResourceAllocation& allocation, + TileXRCcuLowerLayerInstallPlan* plan, + TileXRCcuLowerLayerPlanBuilderReport* report, + void* userData) + { + auto* state = static_cast(userData); + ++state->callCount; + state->syncResourceCount = allocation.remoteXn.num; + if (plan == nullptr || report == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *plan = TileXRCcuLowerLayerInstallPlan{}; + plan->msidTokens.push_back({1, 0x45, 0x1234, 0x5678}); + TileXRCcuPfeInstall pfe; + pfe.dieId = 1; + pfe.pfeOffset = allocation.channels.startId; + pfe.ctx.raw[0] = 0xa1; + plan->pfes.push_back(pfe); + TileXRCcuJettyInstall jetty; + jetty.dieId = 1; + jetty.startJettyCtxId = 9; + jetty.ctxs.resize(1); + jetty.ctxs[0].raw[0] = 0xb2; + plan->jettys.push_back(jetty); + TileXRCcuLowerLayerPayloadReport payloadReport; + for (uint32_t i = 0; i < allocation.channels.num; ++i) { + TileXRCcuChannelInstall channel; + channel.dieId = 1; + channel.channelId = allocation.channels.startId + i; + TileXRCcuChannelCtxV1Spec channelSpec; + channelSpec.remoteEid[0] = static_cast(0x10 + i); + channelSpec.tpn = 1; + channelSpec.sourcePfeId = allocation.channels.startId; + channelSpec.startJettyId = 1024; + channelSpec.jettyCount = 1; + channelSpec.dieId = 1; + channelSpec.memoryTokenId = 0x2345; + channelSpec.memoryTokenValue = 0x67890000U; + channelSpec.remoteCcuVa = 0x0000100054800000ULL; + if (TileXRCcuBuildChannelCtxV1(channelSpec, &channel.ctx, &payloadReport) != TILEXR_SUCCESS) { + report->message = payloadReport.message; + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + plan->channels.push_back(channel); + } + plan->xnClears.push_back({1, allocation.localXn.startId, allocation.localXn.num}); + plan->ckeClears.push_back({1, allocation.notifyCke.startId, allocation.notifyCke.num}); + for (uint32_t i = 0; i < allocation.remoteXn.num; ++i) { + plan->remoteXnBindings.push_back({ + 1, + static_cast(allocation.channels.startId + i), + static_cast(allocation.localXn.startId + i), + static_cast(allocation.remoteXn.startId + i), + static_cast(allocation.notifyCke.startId + i), + i, + true, + static_cast(allocation.localWaitCke.startId + i), + true, + true, + true}); + } + report->msidTokenCount = 1; + report->pfeCount = 1; + report->jettyCount = 1; + report->localJettyCtxCount = 1; + report->channelCount = allocation.channels.num; + report->ckeClearCount = 1; + report->message = "callback lower-layer plan prepared"; + return TILEXR_SUCCESS; + } + + int main() + { + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + FakeState fake; + TileXRCcuDriverAdapter adapter; + TileXRCcuDriverAdapterReport driverReport; + if (adapter.Init(5, FakeCustomChannel, &fake, &driverReport) != TILEXR_SUCCESS) { + return 1; + } + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.sqeArgCount = TILEXR_CCU_SQE_ARGS_LEN; + options.syncResourceCount = 3; + options.syncInstructionCount = 143; + options.bindingsPerSyncResource = 3; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-callback-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.gsaStartId = 510; + options.remoteXnStartId = 2361; + options.remoteXnCount = 8; + options.ckeStartId = 332; + options.remoteNotifyCkeStartId = 364; + options.remoteNotifyCkeCount = 8; + options.channelStartId = 2; + options.driverAdapter = &adapter; + options.repositoryMemoryOps = {FakeAlloc, FakeCopy, FakeFree}; + options.repositoryMemoryUserData = &fake; + + CallbackState callback; + options.prepareLowerLayerPlan = PrepareLowerLayerPlan; + options.lowerLayerPlanUserData = &callback; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + if (ret != TILEXR_SUCCESS || !report.submitReady || report.submitTaskCount != 2) { + std::cerr << "callback install should be submit-ready: " << report.message << "\n"; + return 2; + } + if (callback.callCount != 1 || callback.syncResourceCount != 3) { + std::cerr << "lower-layer callback was not called with allocation\n"; + return 3; + } + if (attempt.plan.kernelLocalGsa.startId != 510 || attempt.plan.kernelLocalGsa.num != 1 || + attempt.allocation.localGsa.startId != 510 || attempt.allocation.localGsa.num != 1) { + std::cerr << "GSA resource did not flow into direct CCU producer plan\n"; + return 13; + } + if (attempt.preparedLowerLayerPlan.remoteXnBindings.size() != 3 || + attempt.lowerLayerPlanReport.message.find("callback") == std::string::npos) { + std::cerr << "prepared lower-layer artifacts were not retained\n"; + return 4; + } + if (!attempt.evidence.remoteXnBound || !attempt.evidence.channelBindingsInstalled || + attempt.evidence.remoteXnSource.source.find("ValidateRemoteXnExchangeBindingProof") == + std::string::npos) { + std::cerr << "prepared lower-layer evidence mismatch\n"; + return 5; + } + const uint64_t expectedPackedToken = (1ULL << 52U) | (0x1234ULL << 32U) | 0x5678ULL; + if (attempt.package.plan.taskWindows[0].args.size() != TILEXR_CCU_SQE_ARGS_LEN || + attempt.package.plan.taskWindows[0].args[0] != basic.resourceAddr || + attempt.package.plan.taskWindows[0].args[1] != basic.resourceAddr || + attempt.package.plan.taskWindows[0].args[2] != expectedPackedToken || + attempt.package.plan.taskWindows[0].args[3] != 0x0000100054800000ULL || + attempt.submitTasks[0].args[0] != basic.resourceAddr || + attempt.submitTasks[0].args[1] != basic.resourceAddr || + attempt.submitTasks[0].args[2] != expectedPackedToken || + attempt.submitTasks[0].args[3] != 0x0000100054800000ULL) { + std::cerr << "hcomm-style SQE task args were not generated from lower-layer token state\n"; + return 6; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_install_attempt_passes_split_cke_ranges_to_lower_layer_callback(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + #include + + using namespace TileXR; + + struct FakeState { + std::vector deviceBytes; + }; + + struct CallbackState { + bool sawSplitCke = false; + uint16_t localWaitStart = 0; + uint16_t remoteNotifyStart = 0; + }; + + int FakeAlloc(uint64_t bytes, void** ptr, void* userData) + { + auto* state = static_cast(userData); + state->deviceBytes.assign(static_cast(bytes), 0); + *ptr = state->deviceBytes.data(); + return 0; + } + + int FakeCopy(void* dst, uint64_t dstBytes, const void* src, uint64_t bytes, void* userData) + { + auto* state = static_cast(userData); + if (dst != state->deviceBytes.data() || dstBytes != state->deviceBytes.size() || + bytes != state->deviceBytes.size()) { + return -1; + } + std::memcpy(dst, src, static_cast(bytes)); + return 0; + } + + int FakeFree(void*, void*) + { + return 0; + } + + int FakeCustomChannel( + uint32_t, + const TileXRCcuCustomChannelIn&, + TileXRCcuCustomChannelOut* out, + void*) + { + out->opRet = 0; + return 0; + } + + int PrepareSplitLowerLayerPlan( + const TileXRCcuResourceAllocation& allocation, + TileXRCcuLowerLayerInstallPlan* plan, + TileXRCcuLowerLayerPlanBuilderReport* report, + void* userData) + { + auto* state = static_cast(userData); + state->localWaitStart = allocation.localWaitCke.startId; + state->remoteNotifyStart = allocation.remoteNotifyCke.startId; + state->sawSplitCke = + allocation.localWaitCke.startId == 332 && + allocation.localWaitCke.num == 3 && + allocation.remoteNotifyCke.startId == 364 && + allocation.remoteNotifyCke.num == 3 && + allocation.notifyCke.startId == 364 && + allocation.notifyCke.num == 3; + if (plan == nullptr || report == nullptr || !state->sawSplitCke) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + + *plan = TileXRCcuLowerLayerInstallPlan{}; + plan->msidTokens.push_back({1, 0x45, 0x1234, 0x5678}); + TileXRCcuPfeInstall pfe; + pfe.dieId = 1; + pfe.pfeOffset = allocation.channels.startId; + pfe.ctx.raw[0] = 0xa1; + plan->pfes.push_back(pfe); + TileXRCcuJettyInstall jetty; + jetty.dieId = 1; + jetty.startJettyCtxId = 9; + jetty.ctxs.resize(1); + jetty.ctxs[0].raw[0] = 0xb2; + plan->jettys.push_back(jetty); + for (uint32_t i = 0; i < allocation.channels.num; ++i) { + TileXRCcuChannelInstall channel; + channel.dieId = 1; + channel.channelId = allocation.channels.startId + i; + channel.ctx.raw[0] = 0xc3 + i; + plan->channels.push_back(channel); + } + plan->xnClears.push_back({1, allocation.localXn.startId, allocation.localXn.num}); + plan->ckeClears.push_back({1, allocation.localWaitCke.startId, allocation.localWaitCke.num}); + for (uint32_t i = 0; i < allocation.remoteXn.num; ++i) { + plan->remoteXnBindings.push_back({ + 1, + allocation.channels.startId + i, + static_cast(allocation.localXn.startId + i), + static_cast(allocation.remoteXn.startId + i), + static_cast(allocation.remoteNotifyCke.startId + i), + i, + true, + static_cast(allocation.localWaitCke.startId + i), + true, + true, + true}); + } + report->message = "split CKE lower-layer plan prepared"; + return TILEXR_SUCCESS; + } + + int main() + { + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + FakeState fake; + TileXRCcuDriverAdapter adapter; + TileXRCcuDriverAdapterReport driverReport; + if (adapter.Init(5, FakeCustomChannel, &fake, &driverReport) != TILEXR_SUCCESS) { + return 1; + } + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.syncResourceCount = 3; + options.syncInstructionCount = 143; + options.bindingsPerSyncResource = 3; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-split-cke-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.ckeStartId = 332; + options.localWaitCkeStartId = 332; + options.remoteNotifyCkeStartId = 364; + options.channelStartId = 2; + options.driverAdapter = &adapter; + options.repositoryMemoryOps = {FakeAlloc, FakeCopy, FakeFree}; + options.repositoryMemoryUserData = &fake; + + CallbackState callback; + options.prepareLowerLayerPlan = PrepareSplitLowerLayerPlan; + options.lowerLayerPlanUserData = &callback; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + if (ret != TILEXR_SUCCESS || !report.submitReady) { + std::cerr << "split CKE direct install should be submit-ready: " + << report.message << "\n"; + return 2; + } + if (!callback.sawSplitCke || callback.localWaitStart != 332 || + callback.remoteNotifyStart != 364) { + std::cerr << "split CKE allocation was not passed to callback\n"; + return 3; + } + if (attempt.plan.kernelLocalCke.startId != 332 || + attempt.plan.syncResources[0].localWaitCke != 332 || + attempt.plan.syncResources[0].notifyCke != 364 || + attempt.evidence.remoteXnSource.detail.find("syncXn") == std::string::npos) { + std::cerr << "split CKE plan or evidence mismatch\n"; + return 4; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_install_attempt_rewrites_sync_resources_from_peer_lower_layer_proof(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + #include + + using namespace TileXR; + + struct FakeState { + std::vector deviceBytes; + }; + + int FakeAlloc(uint64_t bytes, void** ptr, void* userData) + { + auto* state = static_cast(userData); + state->deviceBytes.assign(static_cast(bytes), 0); + *ptr = state->deviceBytes.data(); + return 0; + } + + int FakeCopy(void* dst, uint64_t dstBytes, const void* src, uint64_t bytes, void* userData) + { + auto* state = static_cast(userData); + if (dst != state->deviceBytes.data() || dstBytes != state->deviceBytes.size() || + bytes != state->deviceBytes.size()) { + return -1; + } + std::memcpy(dst, src, static_cast(bytes)); + return 0; + } + + int FakeFree(void*, void*) + { + return 0; + } + + int FakeCustomChannel( + uint32_t, + const TileXRCcuCustomChannelIn&, + TileXRCcuCustomChannelOut* out, + void*) + { + out->opRet = 0; + return 0; + } + + int PreparePeerLowerLayerPlan( + const TileXRCcuResourceAllocation& allocation, + TileXRCcuLowerLayerInstallPlan* plan, + TileXRCcuLowerLayerPlanBuilderReport* report, + void*) + { + if (plan == nullptr || report == nullptr) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + *plan = TileXRCcuLowerLayerInstallPlan{}; + plan->msidTokens.push_back({1, 0x45, 0x1234, 0x5678}); + TileXRCcuPfeInstall pfe; + pfe.dieId = 1; + pfe.pfeOffset = allocation.channels.startId; + pfe.ctx.raw[0] = 0xa1; + plan->pfes.push_back(pfe); + TileXRCcuJettyInstall jetty; + jetty.dieId = 1; + jetty.startJettyCtxId = 9; + jetty.ctxs.resize(1); + jetty.ctxs[0].raw[0] = 0xb2; + plan->jettys.push_back(jetty); + for (uint32_t i = 0; i < allocation.channels.num; ++i) { + TileXRCcuChannelInstall channel; + channel.dieId = 1; + channel.channelId = allocation.channels.startId + i; + channel.ctx.raw[0] = 0xc3 + i; + plan->channels.push_back(channel); + } + plan->xnClears.push_back({1, allocation.localXn.startId, allocation.localXn.num}); + plan->ckeClears.push_back({1, allocation.localWaitCke.startId, allocation.localWaitCke.num}); + + for (uint32_t i = 0; i < allocation.remoteXn.num; ++i) { + const uint16_t peerLocalXn = static_cast(0x1a0 + i); + const uint16_t peerNotifyCke = static_cast(0x360 + i); + plan->remoteXnBindings.push_back({ + 1, + allocation.channels.startId + i, + static_cast(allocation.localXn.startId + i), + peerLocalXn, + peerNotifyCke, + i, + true, + static_cast(allocation.localWaitCke.startId + i), + true, + true, + true}); + } + report->message = "peer lower-layer plan prepared"; + return TILEXR_SUCCESS; + } + + int main() + { + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + FakeState fake; + TileXRCcuDriverAdapter adapter; + TileXRCcuDriverAdapterReport driverReport; + if (adapter.Init(5, FakeCustomChannel, &fake, &driverReport) != TILEXR_SUCCESS) { + return 1; + } + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.sqeArgCount = TILEXR_CCU_SQE_ARGS_LEN; + options.syncResourceCount = 2; + options.syncInstructionCount = 9; + options.bindingsPerSyncResource = 1; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-peer-xn-rewrite-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 0x2a0; + options.ckeStartId = 0x220; + options.localWaitCkeStartId = 0x220; + options.remoteNotifyCkeStartId = 0x360; + options.channelStartId = 2; + options.driverAdapter = &adapter; + options.repositoryMemoryOps = {FakeAlloc, FakeCopy, FakeFree}; + options.repositoryMemoryUserData = &fake; + options.prepareLowerLayerPlan = PreparePeerLowerLayerPlan; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + if (ret != TILEXR_SUCCESS || !report.submitReady) { + std::cerr << "peer XN rewrite should make direct install submit-ready: " + << report.message << "\n"; + return 2; + } + if (attempt.plan.syncResources.size() != 2 || + attempt.plan.syncResources[0].remoteXn != 0x1a0 || + attempt.plan.syncResources[1].remoteXn != 0x1a1 || + attempt.plan.syncResources[0].notifyCke != 0x360 || + attempt.plan.syncResources[1].notifyCke != 0x361) { + std::cerr << "producer plan was not rewritten from peer lower-layer proof\n"; + return 3; + } + if (attempt.package.plan.syncResources[0].remoteXn != 0x1a0 || + attempt.package.plan.syncResources[1].remoteXn != 0x1a1 || + attempt.package.plan.syncResources[0].notifyCke != 0x360 || + attempt.package.plan.syncResources[1].notifyCke != 0x361 || + attempt.submitTasks.size() != 2 || + attempt.submitTasks[0].instStartId != attempt.package.repository.missionStartId || + attempt.submitTasks[0].instCnt != 13 || + attempt.submitTasks[1].instStartId != + attempt.package.repository.missionStartId + attempt.package.repository.sqeLoadCount || + attempt.submitTasks[1].instCnt != + attempt.package.repository.missionCount - attempt.package.repository.sqeLoadCount) { + std::cerr << "launch package did not use peer lower-layer proof resources\n"; + return 4; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_install_attempt_release_helper_frees_repository_receipt_once(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + + using namespace TileXR; + + struct FakeState { + int freeCalls = 0; + int payload = 0; + }; + + int FakeFree(void* ptr, void* userData) + { + auto* state = static_cast(userData); + if (ptr != &state->payload) { + return -1; + } + ++state->freeCalls; + return 0; + } + + int main() + { + FakeState state; + TileXRCcuDirectInstallAttempt attempt; + attempt.repositoryMemoryOps.free = FakeFree; + attempt.repositoryMemoryUserData = &state; + attempt.repositoryReceipt.deviceInstructionPtr = &state.payload; + attempt.repositoryReceipt.deviceInstructionAddr = + reinterpret_cast(&state.payload); + attempt.repositoryReceipt.uploaded = true; + attempt.repositoryReceipt.installed = true; + + if (TileXRCcuReleaseDirectInstallAttemptResources(attempt) != TILEXR_SUCCESS) { + std::cerr << "first release failed: " << attempt.repositoryReleaseReport.message << "\n"; + return 1; + } + if (state.freeCalls != 1 || attempt.repositoryReceipt.deviceInstructionPtr != nullptr || + attempt.repositoryReleaseReport.message != "ok") { + std::cerr << "release did not clear repository receipt\n"; + return 2; + } + if (TileXRCcuReleaseDirectInstallAttemptResources(attempt) != TILEXR_SUCCESS || + state.freeCalls != 1) { + std::cerr << "release helper is not idempotent\n"; + return 3; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_install_attempt_trace_dumps_final_peer_resources_and_tasks_when_enabled(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + + #include + #include + + using namespace TileXR; + + int main() + { + setenv("TILEXR_CCU_DIRECT_TRACE", "1", 1); + + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.syncResourceCount = 1; + options.syncInstructionCount = 2; + options.bindingsPerSyncResource = 1; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-direct-trace-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.ckeStartId = 332; + options.channelStartId = 2; + options.offlineOnly = true; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + (void)ret; + if (!report.pipelineBuilt || attempt.package.tasks.empty()) { + std::cerr << "trace setup failed before package build: " << report.message << "\n"; + return 1; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + self.assertIn("TileXRDirectCcuTrace", result.stderr) + self.assertIn("syncResource[0]", result.stderr) + self.assertIn("remoteXn=", result.stderr) + self.assertIn("taskWindow[0]", result.stderr) + self.assertIn("task[0]", result.stderr) + self.assertNotIn("program.sqeLoad[0]", result.stderr) + self.assertIn("program.sync[0]", result.stderr) + + def test_direct_install_attempt_trace_decodes_lower_layer_contexts_when_enabled(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_orchestrator.h" + #include "ccu/tilexr_ccu_lower_layer_plan_builder.h" + + #include + #include + + using namespace TileXR; + + int PrepareLowerLayer( + const TileXRCcuResourceAllocation& allocation, + TileXRCcuLowerLayerInstallPlan* plan, + TileXRCcuLowerLayerPlanBuilderReport* report, + void* userData) + { + auto* basic = static_cast(userData); + + std::array eid; + for (uint32_t i = 0; i < eid.size(); ++i) { + eid[i] = static_cast(0x20 + i); + } + + TileXRCcuLowerLayerPlanSpec spec; + spec.msidToken.dieId = basic->dieId; + spec.msidToken.msId = basic->msId; + spec.msidToken.tokenId = basic->msidToken.tokenId; + spec.msidToken.tokenValue = basic->msidToken.tokenValue; + spec.msidToken.valid = basic->msidToken.valid; + spec.pfe.dieId = basic->dieId; + spec.pfe.pfeOffset = allocation.channels.startId; + spec.pfe.startJettyId = 0x400; + spec.pfe.startLocalJettyCtxId = 0x02; + + TileXRCcuLowerLayerJettySpec jetty; + jetty.dieId = basic->dieId; + jetty.pfeId = allocation.channels.startId; + jetty.startJettyCtxId = spec.pfe.startLocalJettyCtxId; + jetty.doorbellVa = 0x1122334455667788ULL; + jetty.doorbellTokenId = 0xabcdeU; + jetty.doorbellTokenValue = 0x12345678U; + jetty.sqDepth = 8; + jetty.wqeBasicBlockStartId = 0x34; + spec.jettys.push_back(jetty); + + TileXRCcuLowerLayerChannelSpec channel; + channel.dieId = basic->dieId; + channel.channelId = allocation.channels.startId; + channel.remoteEid = eid; + channel.tpn = 0x00876543U; + channel.sourcePfeId = allocation.channels.startId; + channel.startJettyId = 0x400; + channel.jettyCount = 1; + channel.memoryTokenId = 0x23456U; + channel.memoryTokenValue = 0x3456789aU; + channel.remoteCcuVa = + basic->resourceAddr + TILEXR_CCU_V1_XN_RESOURCE_OFFSET + allocation.remoteXn.startId * 8ULL; + spec.channels.push_back(channel); + + spec.xnClear.dieId = basic->dieId; + spec.xnClear.startXnId = allocation.localXn.startId; + spec.xnClear.count = allocation.localXn.num; + spec.xnClear.valid = true; + spec.ckeClear.dieId = basic->dieId; + spec.ckeClear.startCkeId = allocation.notifyCke.startId; + spec.ckeClear.count = allocation.notifyCke.num; + spec.ckeClear.valid = true; + + TileXRCcuRemoteXnBindingProof proof; + proof.dieId = basic->dieId; + proof.channelId = allocation.channels.startId; + proof.localXn = allocation.localXn.startId; + proof.remoteXn = allocation.remoteXn.startId; + proof.notifyCke = allocation.notifyCke.startId; + proof.localWaitCke = allocation.notifyCke.startId; + proof.peerRank = 0; + proof.peerExchangeObserved = true; + proof.endpointRouteVerified = true; + spec.remoteXnBindings.push_back(proof); + + return TileXRCcuBuildLowerLayerInstallPlan(spec, plan, report); + } + + int main() + { + setenv("TILEXR_CCU_DIRECT_TRACE", "1", 1); + + TileXRCcuBasicInfo basic; + basic.dieId = 1; + basic.msId = 0x45; + basic.msidToken.valid = true; + basic.msidToken.tokenId = 0x1234; + basic.msidToken.tokenValue = 0x5678; + basic.missionKey = 0x059b0f03U; + basic.resourceAddr = 0x100000000ULL; + basic.caps.cap0 = (7U << 24) | (11U << 16) | 169U; + basic.caps.cap1 = (61U << 16) | 31U; + basic.caps.cap2 = (63U << 16) | 35U; + basic.caps.cap3 = (127U << 16) | 3U; + basic.caps.cap4 = 15U; + + TileXRCcuDirectInstallOptions options; + options.basicInfo = &basic; + options.syncResourceCount = 1; + options.syncInstructionCount = 2; + options.bindingsPerSyncResource = 1; + options.deviceId = 3; + options.rank = 1; + options.provider = "unit-test-direct-trace-lower-layer-provider"; + options.missionStartId = 6; + options.instructionStartId = 475; + options.xnStartId = 1961; + options.ckeStartId = 332; + options.channelStartId = 2; + options.offlineOnly = true; + options.prepareLowerLayerPlan = PrepareLowerLayer; + options.lowerLayerPlanUserData = &basic; + + TileXRCcuDirectInstallAttempt attempt; + TileXRCcuDirectInstallReport report; + const int ret = TileXRCcuRunDirectInstallAttempt(options, &attempt, &report); + (void)ret; + if (!report.pipelineBuilt || attempt.preparedLowerLayerPlan.channels.empty()) { + std::cerr << "lower-layer trace setup failed: " << report.message << "\n"; + return 1; + } + return 0; + } + ''' + ) + + result = self.compile_and_run(code) + + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + self.assertIn("decoded=PfeCtx", result.stderr) + self.assertIn("decoded=LocalJettyCtx", result.stderr) + self.assertIn("decoded=ChannelCtxV1", result.stderr) + self.assertIn("startTaJettyId=1024", result.stderr) + self.assertIn("startLocalJettyCtxId=2", result.stderr) + self.assertIn("doorbellVa=0x1122334455667788", result.stderr) + self.assertIn("wqeBasicBlockStartId=52", result.stderr) + self.assertIn("sourcePfeId=2", result.stderr) + self.assertIn("remoteCcuVa=", result.stderr) + + def test_direct_orchestrator_is_wired_and_has_no_private_runtime_surface(self): + cmake = COMM_CMAKE.read_text(encoding="utf-8") + header = DIRECT_HEADER.read_text(encoding="utf-8") + source = DIRECT_SOURCE.read_text(encoding="utf-8") + + self.assertIn("ccu/tilexr_ccu_direct_orchestrator.h", cmake) + self.assertIn("ccu/tilexr_ccu_direct_orchestrator.cpp", cmake) + self.assertIn("struct TileXRCcuDirectInstallOptions", header) + self.assertIn("struct TileXRCcuDirectMemoryCopySpec", header) + self.assertIn("struct TileXRCcuDirectInstallAttempt", header) + self.assertIn("struct TileXRCcuDirectInstallReport", header) + self.assertIn("struct TileXRCcuDirectSubmitReport", header) + self.assertIn("TileXRCcuLowerLayerPlanPrepareFn", header) + self.assertIn("localWaitCkeStartId", header) + self.assertIn("gsaStartId", header) + self.assertIn("remoteNotifyCkeStartId", header) + self.assertIn("prepareLowerLayerPlan", header) + self.assertIn("preparedLowerLayerPlan", header) + self.assertIn("lowerLayerPlanReport", header) + self.assertIn("TileXRCcuSubmitPreparedTasks", header) + self.assertIn("TileXRCcuRunDirectInstallAttempt", header) + self.assertIn("TileXRCcuRunDirectMemoryCopyInstallAttempt", header) + self.assertIn("TileXRCcuDecodeBasicInfo", source) + self.assertIn("TileXRCcuBuildResourceSpec", source) + self.assertIn("TileXRCcuResourceAllocator", source) + self.assertIn("TileXRCcuBuildLaunchPackage", source) + self.assertIn("TileXRCcuBuildMemoryCopyProgram", source) + self.assertIn("BuildDirectMemoryCopyLaunchPackage", source) + self.assertIn("TileXRCcuBindLaunchPackageInstallScope", source) + self.assertIn("TileXRCcuBuildInstallManifest", source) + self.assertIn("options.prepareLowerLayerPlan", source) + self.assertIn("TileXRCcuInstallHardware", source) + self.assertIn("TileXRCcuPrepareSubmitTasks", source) + self.assertIn("TileXRCcuSubmitTask", source) + self.assertIn("remote XN install provider is missing", source) + self.assertIn("TraceDecodedInstr", source) + self.assertIn("TraceDecodedPfeCtx", source) + self.assertIn("TraceDecodedLocalJettyCtx", source) + self.assertIn("TraceDecodedChannelCtxV1", source) + self.assertIn("decoded=SyncXn", source) + self.assertIn("decoded=SyncCke", source) + self.assertIn("decoded=LoadImdToXn", source) + self.assertIn("decoded=LoadImdToGSA", source) + self.assertIn("decoded=TransRmtMemToLocMem", source) + self.assertIn("decoded=TransLocMemToRmtMem", source) + self.assertIn("TILEXR_CCU_TRACE_TRANS_RMT_MEM_TO_LOC_MEM_HEADER", source) + self.assertIn("TILEXR_CCU_TRACE_TRANS_LOC_MEM_TO_RMT_MEM_HEADER", source) + + combined = header + "\n" + source + for needle in PRIVATE_CCU_PRODUCER_NEEDLES: + with self.subTest(needle=needle): + self.assertNotIn(needle, combined) + + def test_direct_install_options_default_to_lower_layer_first(self): + header = DIRECT_HEADER.read_text(encoding="utf-8") + source = DIRECT_SOURCE.read_text(encoding="utf-8") + + self.assertIn( + "TileXRCcuInstallOrder installOrder = " + "TileXRCcuInstallOrder::InstallLowerLayerFirst", + header, + ) + self.assertIn("installRequest.installOrder = options.installOrder", source) + + +if __name__ == "__main__": + unittest.main() diff --git a/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py b/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py new file mode 100644 index 00000000..36c2068c --- /dev/null +++ b/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py @@ -0,0 +1,2121 @@ +#!/usr/bin/env python3 +# +# Copyright (c) 2026 TileXR Project +# + +import shutil +import subprocess +import tempfile +import textwrap +import unittest +from pathlib import Path + + +REPO_ROOT = Path(__file__).resolve().parents[2] +LOADER_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_hccp_loader.h" +LOADER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_hccp_loader.cpp" +TYPES_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_hccp_types.h" +ABI_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_abi_constants.h" +DIRECT_RUNTIME_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_direct_runtime.h" +DIRECT_RUNTIME_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_direct_runtime.cpp" +DRIVER_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.h" +DRIVER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp" + + +class TileXRCcuRaCustomChannelLoaderTest(unittest.TestCase): + def setUp(self): + lifecycle_hardening_tests = { + "test_direct_runtime_ra_init_failure_report_keeps_phy_and_hdc_context", + "test_direct_runtime_tlv_init_failure_report_keeps_ra_phy_and_hdc_context", + "test_direct_runtime_uses_optional_tilexr_endpoint_route_provider_before_env", + } + runtime_source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8") + if (self._testMethodName in lifecycle_hardening_tests and + "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX" not in runtime_source): + self.skipTest("direct runtime lifecycle hardening is introduced by a later stack layer") + + def compile_and_run(self, code: str, extra_sources=None, extra_link_flags=None, extra_env=None): + compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++") + if compiler is None: + self.skipTest("no local C++ compiler found") + extra_sources = extra_sources or [] + extra_link_flags = extra_link_flags or [] + with tempfile.TemporaryDirectory() as temp_dir: + temp_path = Path(temp_dir) + test_cpp = temp_path / "hccp_loader_test.cpp" + test_bin = temp_path / "hccp_loader_test" + test_cpp.write_text(code, encoding="utf-8") + subprocess.run( + [ + compiler, + "-std=c++14", + "-I", + str(REPO_ROOT / "src" / "include"), + "-I", + str(REPO_ROOT / "src" / "comm"), + str(test_cpp), + *[str(source) for source in extra_sources], + "-o", + str(test_bin), + *extra_link_flags, + ], + cwd=REPO_ROOT, + check=True, + text=True, + capture_output=True, + ) + env = None if extra_env is None else {**extra_env} + return subprocess.run( + [str(test_bin)], + cwd=REPO_ROOT, + check=False, + text=True, + capture_output=True, + env=env, + ) + + def test_ccu_hccp_loader_exposes_ra_custom_channel_for_tilexr_ccu(self): + header = LOADER_HEADER.read_text(encoding="utf-8") + source = LOADER_SOURCE.read_text(encoding="utf-8") + + self.assertIn("TileXRCcuRaCustomChannelFunc", TYPES_HEADER.read_text(encoding="utf-8")) + self.assertIn("RaCustomChannel", header) + self.assertIn("RaCustomChannel = nullptr", header) + self.assertIn('LoadSymbol(raHandle_, RaCustomChannel, "RaCustomChannel", "ra_custom_channel")', source) + self.assertIn("RaCustomChannel = nullptr", source) + + def test_ccu_hccp_loader_exposes_runtime_phy_id_mapping_for_ccu_runtime(self): + header = LOADER_HEADER.read_text(encoding="utf-8") + source = LOADER_SOURCE.read_text(encoding="utf-8") + + self.assertIn("TileXRCcuRtGetDevicePhyIdByIndexFunc", TYPES_HEADER.read_text(encoding="utf-8")) + self.assertIn("RtGetDevicePhyIdByIndex", header) + self.assertIn("RtGetDevicePhyIdByIndex = nullptr", header) + self.assertIn('dlopen("libruntime.so", RTLD_NOW)', source) + self.assertIn( + 'LoadOptionalSymbol(runtimeHandle_, RtGetDevicePhyIdByIndex, "rtGetDevicePhyIdByIndex", nullptr)', + source, + ) + self.assertIn("ResolveDevicePhyId", header) + self.assertIn("ResolveDevicePhyId", source) + + def test_ccu_hccp_loader_exposes_public_ra_ctx_resource_window_symbols(self): + header = LOADER_HEADER.read_text(encoding="utf-8") + source = LOADER_SOURCE.read_text(encoding="utf-8") + types = TYPES_HEADER.read_text(encoding="utf-8") + + for needle in [ + "TileXRCcuRaGetDevEidInfoNumFunc", + "TileXRCcuRaGetDevEidInfoListFunc", + "TileXRCcuRaCtxInitFunc", + "TileXRCcuRaCtxDeinitFunc", + "TileXRCcuRaCtxTokenIdAllocFunc", + "TileXRCcuRaCtxTokenIdFreeFunc", + "TileXRCcuRaCtxLmemRegisterFunc", + "TileXRCcuRaCtxLmemUnregisterFunc", + "TileXRCcuRaGetSecRandomFunc", + "TileXRCcuHccpMrRegInfo", + ]: + with self.subTest(needle=needle): + self.assertIn(needle, types) + + for needle in [ + "RaGetDevEidInfoNum", + "RaGetDevEidInfoList", + "RaCtxInit", + "RaCtxDeinit", + "RaCtxTokenIdAlloc", + "RaCtxTokenIdFree", + "RaCtxLmemRegister", + "RaCtxLmemUnregister", + "RaGetSecRandom", + ]: + with self.subTest(needle=needle): + self.assertIn(needle, header) + self.assertIn(needle, source) + + def test_ccu_hccp_loader_exposes_public_ra_ctx_endpoint_route_symbols(self): + header = LOADER_HEADER.read_text(encoding="utf-8") + source = LOADER_SOURCE.read_text(encoding="utf-8") + types = TYPES_HEADER.read_text(encoding="utf-8") + + for needle in [ + "TileXRCcuRaCtxCqCreateFunc", + "TileXRCcuRaCtxCqDestroyFunc", + "TileXRCcuRaCtxQpCreateFunc", + "TileXRCcuRaCtxQpDestroyFunc", + "TileXRCcuRaCtxQpImportFunc", + "TileXRCcuRaCtxQpUnimportFunc", + "TileXRCcuRaCtxQpBindFunc", + "TileXRCcuRaCtxQpUnbindFunc", + "TileXRCcuRaGetTpInfoListAsyncFunc", + "TileXRCcuRaGetAsyncReqResultFunc", + "TileXRCcuHccpQpCreateAttr", + "TileXRCcuHccpQpImportInfo", + "TileXRCcuHccpGetTpCfg", + ]: + with self.subTest(needle=needle): + self.assertIn(needle, types) + + for needle in [ + "RaCtxCqCreate", + "RaCtxCqDestroy", + "RaCtxQpCreate", + "RaCtxQpDestroy", + "RaCtxQpImport", + "RaCtxQpUnimport", + "RaCtxQpBind", + "RaCtxQpUnbind", + "RaGetTpInfoListAsync", + "RaGetAsyncReqResult", + ]: + with self.subTest(needle=needle): + self.assertIn(needle, header) + self.assertIn(needle, source) + + def test_ccu_hccp_loader_fails_closed_when_runtime_phy_id_mapping_fails(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + int FakeRtGetDevicePhyIdByIndex(uint32_t logicDevId, uint32_t* phyId) + { + if (logicDevId != 7 || phyId == nullptr) { + return -99; + } + *phyId = 0xbeefU; + return -1234; + } + + int main() + { + TileXRCcuHccpLoader loader; + loader.loaded_ = true; + loader.RtGetDevicePhyIdByIndex = FakeRtGetDevicePhyIdByIndex; + + uint32_t phyId = 0; + TileXRCcuHccpLoaderReport report; + const int ret = loader.ResolveDevicePhyId(7, &phyId, &report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "phy id mapping failure silently succeeded\n"; + return 1; + } + if (phyId == 7 || phyId == 0xbeefU) { + std::cerr << "failed mapping leaked fallback phy id: " << phyId << "\n"; + return 2; + } + if (report.message.find("rtGetDevicePhyIdByIndex failed") == std::string::npos) { + std::cerr << "missing phy id failure message: " << report.message << "\n"; + return 3; + } + if (report.logicDevId != 7 || report.runtimePhyIdRet != -1234) { + std::cerr << "phy id report mismatch\n"; + return 4; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_initializes_and_deinitializes_ra_hdc_session(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + TileXRCcuRaInitConfig g_lastInitConfig {}; + TileXRCcuRaInitConfig g_lastDeinitConfig {}; + int g_initCalls = 0; + int g_deinitCalls = 0; + int g_openCalls = 0; + int g_closeCalls = 0; + + int FakeRtOpenNetService(const TileXRCcuRtNetServiceOpenArgs*) + { + ++g_openCalls; + return 0; + } + + int FakeRtCloseNetService() + { + ++g_closeCalls; + return 0; + } + + int FakeRaInit(TileXRCcuRaInitConfig* config) + { + if (config == nullptr) { + return -9; + } + g_lastInitConfig = *config; + ++g_initCalls; + return 0; + } + + int FakeRaDeinit(TileXRCcuRaInitConfig* config) + { + if (config == nullptr) { + return -8; + } + g_lastDeinitConfig = *config; + ++g_deinitCalls; + return 0; + } + + int main() + { + TileXRCcuHccpLoader first; + first.loaded_ = true; + first.RtOpenNetService = FakeRtOpenNetService; + first.RtCloseNetService = FakeRtCloseNetService; + first.RaInit = FakeRaInit; + first.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoader second; + second.loaded_ = true; + second.RtOpenNetService = FakeRtOpenNetService; + second.RtCloseNetService = FakeRtCloseNetService; + second.RaInit = FakeRaInit; + second.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoaderReport report; + if (first.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report) != TILEXR_SUCCESS) { + std::cerr << "first init failed: " << report.message << "\n"; + return 1; + } + if (second.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report) != TILEXR_SUCCESS) { + std::cerr << "second init failed: " << report.message << "\n"; + return 2; + } + if (g_initCalls != 1 || !report.raInitialized || report.raInitRefCount != 2) { + std::cerr << "RA init refcount mismatch\n"; + return 3; + } + if (g_openCalls != 1 || report.netServiceRefCount != 2) { + std::cerr << "net service refcount mismatch\n"; + return 8; + } + if (g_lastInitConfig.phyId != 3 || + g_lastInitConfig.nicPosition != TILEXR_CCU_NETWORK_OFFLINE || + g_lastInitConfig.hdcType != TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2 || + !g_lastInitConfig.enableHdcAsync) { + std::cerr << "RA init config mismatch\n"; + return 4; + } + + second.Unload(); + if (g_deinitCalls != 0) { + std::cerr << "RA deinit happened before last reference\n"; + return 5; + } + if (g_closeCalls != 0) { + std::cerr << "runtime net service closed before last reference\n"; + return 9; + } + first.Unload(); + if (g_deinitCalls != 1) { + std::cerr << "RA deinit count mismatch: " << g_deinitCalls << "\n"; + return 6; + } + if (g_closeCalls != 1) { + std::cerr << "runtime net service close count mismatch: " << g_closeCalls << "\n"; + return 10; + } + if (g_lastDeinitConfig.phyId != 3 || + g_lastDeinitConfig.nicPosition != TILEXR_CCU_NETWORK_OFFLINE || + g_lastDeinitConfig.hdcType != TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2) { + std::cerr << "RA deinit config mismatch\n"; + return 7; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_opens_runtime_net_service_before_ra_init(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + #include + #include + + using namespace TileXR; + + std::string g_openArg; + int g_openCalls = 0; + int g_closeCalls = 0; + int g_initCalls = 0; + int g_order = 0; + int g_openOrder = 0; + int g_initOrder = 0; + + int FakeRtOpenNetService(const TileXRCcuRtNetServiceOpenArgs* args) + { + ++g_openCalls; + g_openOrder = ++g_order; + if (args == nullptr || args->extParamList == nullptr || args->extParamCnt != 1 || + args->extParamList[0].paramInfo == nullptr) { + return -11; + } + g_openArg.assign(args->extParamList[0].paramInfo, args->extParamList[0].paramLen); + return 0; + } + + int FakeRtCloseNetService() + { + ++g_closeCalls; + return 0; + } + + int FakeRaInit(TileXRCcuRaInitConfig*) + { + ++g_initCalls; + g_initOrder = ++g_order; + return 0; + } + + int FakeRaDeinit(TileXRCcuRaInitConfig*) + { + return 0; + } + + int main() + { + TileXRCcuHccpLoader loader; + loader.loaded_ = true; + loader.RtOpenNetService = FakeRtOpenNetService; + loader.RtCloseNetService = FakeRtCloseNetService; + loader.RaInit = FakeRaInit; + loader.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoaderReport report; + if (loader.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report) != TILEXR_SUCCESS) { + std::cerr << "init failed: " << report.message << "\n"; + return 1; + } + if (g_openCalls != 1 || g_initCalls != 1 || g_openArg != "--hdcType=18") { + std::cerr << "open/init mismatch open=" << g_openCalls + << " init=" << g_initCalls + << " arg=" << g_openArg << "\n"; + return 2; + } + if (g_openOrder == 0 || g_initOrder == 0 || g_openOrder > g_initOrder) { + std::cerr << "rtOpenNetService did not precede RaInit\n"; + return 3; + } + loader.Unload(); + if (g_closeCalls != 1) { + std::cerr << "close count mismatch: " << g_closeCalls << "\n"; + return 4; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_refcounts_runtime_net_service_by_hdc_type(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + int g_openCalls = 0; + int g_closeCalls = 0; + int g_initCalls = 0; + int g_deinitCalls = 0; + + int FakeRtOpenNetService(const TileXRCcuRtNetServiceOpenArgs*) + { + ++g_openCalls; + return 0; + } + + int FakeRtCloseNetService() + { + ++g_closeCalls; + return 0; + } + + int FakeRaInit(TileXRCcuRaInitConfig*) + { + ++g_initCalls; + return 0; + } + + int FakeRaDeinit(TileXRCcuRaInitConfig*) + { + ++g_deinitCalls; + return 0; + } + + int main() + { + TileXRCcuHccpLoader first; + first.loaded_ = true; + first.RtOpenNetService = FakeRtOpenNetService; + first.RtCloseNetService = FakeRtCloseNetService; + first.RaInit = FakeRaInit; + first.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoader second; + second.loaded_ = true; + second.RtOpenNetService = FakeRtOpenNetService; + second.RtCloseNetService = FakeRtCloseNetService; + second.RaInit = FakeRaInit; + second.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoaderReport report; + if (first.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report) != TILEXR_SUCCESS || + second.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report) != TILEXR_SUCCESS) { + std::cerr << "init failed: " << report.message << "\n"; + return 1; + } + if (g_openCalls != 1 || g_initCalls != 1 || report.raInitRefCount != 2) { + std::cerr << "open/init ref mismatch open=" << g_openCalls + << " init=" << g_initCalls + << " refs=" << report.raInitRefCount << "\n"; + return 2; + } + second.Unload(); + if (g_closeCalls != 0 || g_deinitCalls != 0) { + std::cerr << "closed/deinit before final reference\n"; + return 3; + } + first.Unload(); + if (g_closeCalls != 1 || g_deinitCalls != 1) { + std::cerr << "final close/deinit mismatch close=" << g_closeCalls + << " deinit=" << g_deinitCalls << "\n"; + return 4; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_fails_before_ra_init_when_net_service_open_fails(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + int g_openCalls = 0; + int g_initCalls = 0; + + int FakeRtOpenNetService(const TileXRCcuRtNetServiceOpenArgs*) + { + ++g_openCalls; + return 128000; + } + + int FakeRtCloseNetService() + { + return 0; + } + + int FakeRaInit(TileXRCcuRaInitConfig*) + { + ++g_initCalls; + return 0; + } + + int FakeRaDeinit(TileXRCcuRaInitConfig*) + { + return 0; + } + + int main() + { + TileXRCcuHccpLoader loader; + loader.loaded_ = true; + loader.RtOpenNetService = FakeRtOpenNetService; + loader.RtCloseNetService = FakeRtCloseNetService; + loader.RaInit = FakeRaInit; + loader.RaDeinit = FakeRaDeinit; + + TileXRCcuHccpLoaderReport report; + const int ret = loader.InitRaHdc(3, TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2, true, &report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "open failure unexpectedly succeeded\n"; + return 1; + } + if (g_openCalls != 1 || g_initCalls != 0) { + std::cerr << "open failure call mismatch open=" << g_openCalls + << " init=" << g_initCalls << "\n"; + return 2; + } + if (report.message.find("rtOpenNetService failed ret=128000") == std::string::npos || + report.message.find("--hdcType=18") == std::string::npos) { + std::cerr << "missing open failure diagnostic: " << report.message << "\n"; + return 3; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_initializes_and_releases_ccu_tlv_session(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + TileXRCcuTlvInitInfo g_lastInitInfo {}; + uint32_t g_lastRequestModule = 0; + uint32_t g_lastRequestType = 0; + uint32_t g_lastReleaseType = 0; + int g_initCalls = 0; + int g_requestCalls = 0; + int g_deinitCalls = 0; + int g_order = 0; + int g_initOrder = 0; + int g_requestOrder = 0; + int g_deinitOrder = 0; + + int FakeRaTlvInit(TileXRCcuTlvInitInfo* initInfo, uint32_t* bufferSize, void** tlvHandle) + { + ++g_initCalls; + g_initOrder = ++g_order; + if (initInfo == nullptr || bufferSize == nullptr || tlvHandle == nullptr) { + return -1; + } + g_lastInitInfo = *initInfo; + *bufferSize = 4096; + *tlvHandle = reinterpret_cast(0x12345678ULL); + return 0; + } + + int FakeRaTlvRequest(void* tlvHandle, uint32_t moduleType, TileXRCcuTlvMsg* sendMsg, TileXRCcuTlvMsg* recvMsg) + { + ++g_requestCalls; + g_requestOrder = ++g_order; + if (tlvHandle != reinterpret_cast(0x12345678ULL) || sendMsg == nullptr || recvMsg == nullptr) { + return -2; + } + g_lastRequestModule = moduleType; + if (sendMsg->type == TILEXR_CCU_TLV_MSG_TYPE_CCU_UNINIT) { + g_lastReleaseType = sendMsg->type; + } else { + g_lastRequestType = sendMsg->type; + } + return 0; + } + + int FakeRaTlvDeinit(void* tlvHandle) + { + ++g_deinitCalls; + g_deinitOrder = ++g_order; + return tlvHandle == reinterpret_cast(0x12345678ULL) ? 0 : -3; + } + + int main() + { + TileXRCcuHccpLoader loader; + loader.loaded_ = true; + loader.RaTlvInit = FakeRaTlvInit; + loader.RaTlvRequest = FakeRaTlvRequest; + loader.RaTlvDeinit = FakeRaTlvDeinit; + + TileXRCcuHccpLoaderReport report; + if (loader.InitCcuTlv(3, &report) != TILEXR_SUCCESS) { + std::cerr << "tlv init failed: " << report.message << "\n"; + return 1; + } + if (!report.ccuTlvInitialized || report.ccuTlvRefCount != 1 || + report.ccuTlvBufferSize != 4096) { + std::cerr << "tlv report mismatch\n"; + return 2; + } + if (g_initCalls != 1 || g_requestCalls != 1 || g_deinitCalls != 0) { + std::cerr << "tlv init call count mismatch\n"; + return 3; + } + if (g_lastInitInfo.version != TILEXR_CCU_TLV_VERSION || + g_lastInitInfo.phyId != 3 || + g_lastInitInfo.nicPosition != TILEXR_CCU_NETWORK_OFFLINE || + g_lastRequestModule != TILEXR_CCU_TLV_MODULE_TYPE_CCU || + g_lastRequestType != TILEXR_CCU_TLV_MSG_TYPE_CCU_INIT) { + std::cerr << "tlv init/request envelope mismatch\n"; + return 4; + } + if (g_initOrder == 0 || g_requestOrder == 0 || g_initOrder > g_requestOrder) { + std::cerr << "tlv request did not follow init\n"; + return 5; + } + + loader.Unload(); + if (g_requestCalls != 2 || g_lastReleaseType != TILEXR_CCU_TLV_MSG_TYPE_CCU_UNINIT || + g_deinitCalls != 1) { + std::cerr << "tlv release mismatch requestCalls=" << g_requestCalls + << " releaseType=" << g_lastReleaseType + << " deinitCalls=" << g_deinitCalls << "\n"; + return 6; + } + if (g_deinitOrder == 0 || g_requestOrder > g_deinitOrder) { + std::cerr << "tlv deinit happened before release request\n"; + return 7; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_refcounts_ccu_tlv_by_device_phy_id(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + int g_initCalls = 0; + int g_requestCalls = 0; + int g_deinitCalls = 0; + + int FakeRaTlvInit(TileXRCcuTlvInitInfo*, uint32_t* bufferSize, void** tlvHandle) + { + ++g_initCalls; + *bufferSize = 1024; + *tlvHandle = reinterpret_cast(0x1000ULL); + return 0; + } + + int FakeRaTlvRequest(void*, uint32_t, TileXRCcuTlvMsg*, TileXRCcuTlvMsg*) + { + ++g_requestCalls; + return 0; + } + + int FakeRaTlvDeinit(void*) + { + ++g_deinitCalls; + return 0; + } + + int main() + { + TileXRCcuHccpLoader first; + first.loaded_ = true; + first.RaTlvInit = FakeRaTlvInit; + first.RaTlvRequest = FakeRaTlvRequest; + first.RaTlvDeinit = FakeRaTlvDeinit; + + TileXRCcuHccpLoader second; + second.loaded_ = true; + second.RaTlvInit = FakeRaTlvInit; + second.RaTlvRequest = FakeRaTlvRequest; + second.RaTlvDeinit = FakeRaTlvDeinit; + + TileXRCcuHccpLoaderReport report; + if (first.InitCcuTlv(3, &report) != TILEXR_SUCCESS || + second.InitCcuTlv(3, &report) != TILEXR_SUCCESS) { + std::cerr << "tlv init failed: " << report.message << "\n"; + return 1; + } + if (g_initCalls != 1 || g_requestCalls != 1 || report.ccuTlvRefCount != 2) { + std::cerr << "tlv refcount mismatch init=" << g_initCalls + << " request=" << g_requestCalls + << " refs=" << report.ccuTlvRefCount << "\n"; + return 2; + } + second.Unload(); + if (g_deinitCalls != 0 || g_requestCalls != 1) { + std::cerr << "tlv released before final reference\n"; + return 3; + } + first.Unload(); + if (g_deinitCalls != 1 || g_requestCalls != 2) { + std::cerr << "tlv final release mismatch deinit=" << g_deinitCalls + << " request=" << g_requestCalls << "\n"; + return 4; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_hccp_loader_rolls_back_tlv_handle_when_ccu_init_request_fails(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_hccp_loader.h" + #undef private + + #include + + using namespace TileXR; + + int g_initCalls = 0; + int g_requestCalls = 0; + int g_deinitCalls = 0; + + int FakeRaTlvInit(TileXRCcuTlvInitInfo*, uint32_t* bufferSize, void** tlvHandle) + { + ++g_initCalls; + *bufferSize = 1024; + *tlvHandle = reinterpret_cast(0x1000ULL); + return 0; + } + + int FakeRaTlvRequest(void*, uint32_t moduleType, TileXRCcuTlvMsg* sendMsg, TileXRCcuTlvMsg*) + { + ++g_requestCalls; + if (moduleType != TILEXR_CCU_TLV_MODULE_TYPE_CCU || + sendMsg == nullptr || sendMsg->type != TILEXR_CCU_TLV_MSG_TYPE_CCU_INIT) { + return -77; + } + return 128308; + } + + int FakeRaTlvDeinit(void*) + { + ++g_deinitCalls; + return 0; + } + + int main() + { + TileXRCcuHccpLoader loader; + loader.loaded_ = true; + loader.RaTlvInit = FakeRaTlvInit; + loader.RaTlvRequest = FakeRaTlvRequest; + loader.RaTlvDeinit = FakeRaTlvDeinit; + + TileXRCcuHccpLoaderReport report; + const int ret = loader.InitCcuTlv(3, &report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "tlv request failure unexpectedly succeeded\n"; + return 1; + } + if (g_initCalls != 1 || g_requestCalls != 1 || g_deinitCalls != 1) { + std::cerr << "tlv rollback mismatch init=" << g_initCalls + << " request=" << g_requestCalls + << " deinit=" << g_deinitCalls << "\n"; + return 2; + } + if (report.ccuTlvInitialized || report.raTlvRequestRet != 128308 || + report.message.find("RaTlvRequest CCU_INIT failed ret=128308") == std::string::npos) { + std::cerr << "weak tlv failure diagnostic: " << report.message << "\n"; + return 3; + } + return 0; + } + ''' + ) + result = self.compile_and_run(code, extra_sources=[LOADER_SOURCE], extra_link_flags=["-ldl"]) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_ccu_runtime_initializes_ra_hdc_after_resolving_device_phy_id(self): + header = DIRECT_RUNTIME_HEADER.read_text(encoding="utf-8") + source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8") + types = TYPES_HEADER.read_text(encoding="utf-8") + loader_header = LOADER_HEADER.read_text(encoding="utf-8") + loader_source = LOADER_SOURCE.read_text(encoding="utf-8") + + self.assertIn("TileXRCcuRaInitConfig", types) + self.assertIn("TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2", types) + self.assertIn("TileXRCcuRaInitFunc", types) + self.assertIn("TileXRCcuRaDeinitFunc", types) + self.assertIn("RaInit", loader_header) + self.assertIn("RaDeinit", loader_header) + self.assertIn("InitRaHdc", loader_header) + self.assertIn('LoadSymbol(raHandle_, RaInit, "RaInit", nullptr)', loader_source) + self.assertIn('LoadSymbol(raHandle_, RaDeinit, "RaDeinit", nullptr)', loader_source) + self.assertIn("InitCcuTlv", loader_header) + self.assertIn('LoadOptionalSymbol(raHandle_, RaTlvInit, "RaTlvInit", nullptr)', loader_source) + self.assertIn('LoadOptionalSymbol(raHandle_, RaTlvRequest, "RaTlvRequest", nullptr)', loader_source) + self.assertIn('LoadOptionalSymbol(raHandle_, RaTlvDeinit, "RaTlvDeinit", nullptr)', loader_source) + self.assertIn("loader_.InitRaHdc", source) + self.assertIn("loader_.InitRaHdc(devicePhyId_, hdcType, true, &raReport)", source) + self.assertIn("loader_.InitCcuTlv(devicePhyId_, &tlvReport)", source) + self.assertIn("SelectDirectCcuHdcType", source) + self.assertIn("TILEXR_CCU_DIRECT_HDC_TYPE", source) + self.assertIn("raInitialized", header) + self.assertIn("ccuTlvInitialized", header) + + def test_direct_runtime_ra_init_failure_report_keeps_phy_and_hdc_context(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_runtime.h" + + #include + #include + + using namespace TileXR; + + int main() + { + setenv("TILEXR_CCU_DIRECT_HDC_TYPE", "6", 1); + + TileXRCcuDirectRuntime runtime; + TileXRCcuDirectRuntimeOptions options; + options.devId = 7; + options.rank = 0; + options.rankSize = 2; + TileXRCcuDirectRuntimeReport report; + const int ret = runtime.Init(options, &report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "runtime init unexpectedly succeeded\n"; + return 1; + } + if (report.logicDevId != 7 || report.devicePhyId != 0xbeefU || + report.hdcType != TILEXR_CCU_HDC_SERVICE_TYPE_RDMA || + report.raInitialized) { + std::cerr << "report lost init context logic=" << report.logicDevId + << " phy=" << report.devicePhyId + << " hdc=" << report.hdcType + << " ra=" << report.raInitialized << "\n"; + return 2; + } + if (report.message.find("RaInit failed") == std::string::npos) { + std::cerr << "missing RA failure diagnostic: " << report.message << "\n"; + return 3; + } + return 0; + } + ''' + ) + libra_code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_hccp_types.h" + + extern "C" int RaCustomChannel( + TileXR::TileXRCcuRaInfo, + TileXR::TileXRCcuCustomChannelIn*, + TileXR::TileXRCcuCustomChannelOut*) + { + return 0; + } + + extern "C" int RaInit(TileXR::TileXRCcuRaInitConfig*) + { + return -55; + } + + extern "C" int RaDeinit(TileXR::TileXRCcuRaInitConfig*) + { + return 0; + } + ''' + ) + runtime_code = textwrap.dedent( + r''' + #include + + extern "C" int rtGetDevicePhyIdByIndex(uint32_t logicDevId, uint32_t* phyId) + { + if (logicDevId != 7 || phyId == nullptr) { + return -1; + } + *phyId = 0xbeefU; + return 0; + } + + struct TileXRCcuRtProcExtParam { + const char *paramInfo; + uint64_t paramLen; + }; + + struct TileXRCcuRtNetServiceOpenArgs { + TileXRCcuRtProcExtParam *extParamList; + uint64_t extParamCnt; + }; + + extern "C" int rtOpenNetService(const TileXRCcuRtNetServiceOpenArgs*) + { + return 0; + } + + extern "C" int rtCloseNetService() + { + return 0; + } + ''' + ) + + compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++") + if compiler is None: + self.skipTest("no local C++ compiler found") + with tempfile.TemporaryDirectory() as temp_dir: + temp_path = Path(temp_dir) + libra_cpp = temp_path / "fake_libra.cpp" + runtime_cpp = temp_path / "fake_libruntime.cpp" + libra_so = temp_path / "libra.so" + runtime_so = temp_path / "libruntime.so" + libra_cpp.write_text(libra_code, encoding="utf-8") + runtime_cpp.write_text(runtime_code, encoding="utf-8") + common_flags = [ + compiler, + "-std=c++14", + "-shared", + "-fPIC", + "-I", + str(REPO_ROOT / "src" / "include"), + "-I", + str(REPO_ROOT / "src" / "comm"), + ] + subprocess.run( + [*common_flags, str(libra_cpp), "-o", str(libra_so)], + cwd=REPO_ROOT, + check=True, + text=True, + capture_output=True, + ) + subprocess.run( + [*common_flags, str(runtime_cpp), "-o", str(runtime_so)], + cwd=REPO_ROOT, + check=True, + text=True, + capture_output=True, + ) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + DRIVER_SOURCE, + ], + extra_link_flags=["-ldl"], + extra_env={"LD_LIBRARY_PATH": str(temp_path)}, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_runtime_tlv_init_failure_report_keeps_ra_phy_and_hdc_context(self): + code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_direct_runtime.h" + + #include + + using namespace TileXR; + + int main() + { + TileXRCcuDirectRuntime runtime; + TileXRCcuDirectRuntimeOptions options; + options.devId = 7; + options.rank = 0; + options.rankSize = 2; + TileXRCcuDirectRuntimeReport report; + const int ret = runtime.Init(options, &report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "runtime init unexpectedly succeeded\n"; + return 1; + } + if (report.logicDevId != 7 || report.devicePhyId != 0xbeefU || + report.hdcType != TILEXR_CCU_HDC_SERVICE_TYPE_RDMA_V2 || + !report.raInitialized || report.ccuTlvInitialized) { + std::cerr << "report lost TLV failure context logic=" << report.logicDevId + << " phy=" << report.devicePhyId + << " hdc=" << report.hdcType + << " ra=" << report.raInitialized + << " tlv=" << report.ccuTlvInitialized << "\n"; + return 2; + } + if (report.message.find("RaTlvRequest CCU_INIT failed ret=128308") == std::string::npos) { + std::cerr << "missing TLV failure diagnostic: " << report.message << "\n"; + return 3; + } + return 0; + } + ''' + ) + libra_code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_hccp_types.h" + + extern "C" int RaCustomChannel( + TileXR::TileXRCcuRaInfo, + TileXR::TileXRCcuCustomChannelIn*, + TileXR::TileXRCcuCustomChannelOut*) + { + return 0; + } + + extern "C" int RaInit(TileXR::TileXRCcuRaInitConfig*) + { + return 0; + } + + extern "C" int RaDeinit(TileXR::TileXRCcuRaInitConfig*) + { + return 0; + } + + extern "C" int RaTlvInit(TileXR::TileXRCcuTlvInitInfo*, uint32_t* bufferSize, void** tlvHandle) + { + *bufferSize = 1024; + *tlvHandle = reinterpret_cast(0x1000ULL); + return 0; + } + + extern "C" int RaTlvRequest( + void*, + uint32_t, + TileXR::TileXRCcuTlvMsg* sendMsg, + TileXR::TileXRCcuTlvMsg*) + { + return sendMsg != nullptr && sendMsg->type == TileXR::TILEXR_CCU_TLV_MSG_TYPE_CCU_INIT ? + 128308 : 0; + } + + extern "C" int RaTlvDeinit(void*) + { + return 0; + } + ''' + ) + runtime_code = textwrap.dedent( + r''' + #include + + extern "C" int rtGetDevicePhyIdByIndex(uint32_t logicDevId, uint32_t* phyId) + { + if (logicDevId != 7 || phyId == nullptr) { + return -1; + } + *phyId = 0xbeefU; + return 0; + } + + struct TileXRCcuRtProcExtParam { + const char *paramInfo; + uint64_t paramLen; + }; + + struct TileXRCcuRtNetServiceOpenArgs { + TileXRCcuRtProcExtParam *extParamList; + uint64_t extParamCnt; + }; + + extern "C" int rtOpenNetService(const TileXRCcuRtNetServiceOpenArgs*) + { + return 0; + } + + extern "C" int rtCloseNetService() + { + return 0; + } + ''' + ) + + compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++") + if compiler is None: + self.skipTest("no local C++ compiler found") + with tempfile.TemporaryDirectory() as temp_dir: + temp_path = Path(temp_dir) + libra_cpp = temp_path / "fake_libra.cpp" + runtime_cpp = temp_path / "fake_libruntime.cpp" + libra_so = temp_path / "libra.so" + runtime_so = temp_path / "libruntime.so" + libra_cpp.write_text(libra_code, encoding="utf-8") + runtime_cpp.write_text(runtime_code, encoding="utf-8") + common_flags = [ + compiler, + "-std=c++14", + "-shared", + "-fPIC", + "-I", + str(REPO_ROOT / "src" / "include"), + "-I", + str(REPO_ROOT / "src" / "comm"), + ] + subprocess.run( + [*common_flags, str(libra_cpp), "-o", str(libra_so)], + cwd=REPO_ROOT, + check=True, + text=True, + capture_output=True, + ) + subprocess.run( + [*common_flags, str(runtime_cpp), "-o", str(runtime_so)], + cwd=REPO_ROOT, + check=True, + text=True, + capture_output=True, + ) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + DRIVER_SOURCE, + ], + extra_link_flags=["-ldl"], + extra_env={"LD_LIBRARY_PATH": str(temp_path)}, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_ccu_driver_adapter_keeps_dynamic_loading_in_ccu_hccp_loader(self): + combined = DRIVER_HEADER.read_text(encoding="utf-8") + "\n" + DRIVER_SOURCE.read_text(encoding="utf-8") + + for needle in [ + "dlopen", + "dlsym", + "libra.so", + "libhcomm", + "libhccl_v2", + "#include + #include + + using namespace TileXR; + + int main() + { + TileXRCcuHccpLoader loader; + TileXRCcuHccpLoaderReport report; + if (loader.LoadEndpointRouteProviderFromEnv(&report) != TILEXR_SUCCESS) { + std::cerr << "load provider failed: " << report.message << "\n"; + return 1; + } + if (loader.CollectLocalEndpointRoute == nullptr) { + std::cerr << "endpoint route provider function missing\n"; + return 2; + } + + TileXRCcuEndpointRouteProviderResourceWindow window; + window.addr = 0x10000000ULL; + window.bytes = 0x2000; + window.tokenId = 0x1234; + window.rawTokenId = 0x2234; + window.tokenValue = 0; + + TileXRCcuEndpointRouteProviderRoute route; + if (loader.CollectLocalEndpointRoute(0x55, &window, &route) != TILEXR_SUCCESS) { + std::cerr << "provider call failed\n"; + return 3; + } + if (!route.endpointRouteVerified || route.remoteEid[0] != 0xaa || + route.tpn != 0x10203 || route.doorbellVa != 0x1122334455667788ULL || + route.doorbellTokenId != 0x3456 || route.doorbellTokenValue != 0 || + route.sqDepth != 64) { + std::cerr << "route mismatch\n"; + return 4; + } + if (!report.endpointRouteProviderLoaded) { + std::cerr << "provider report did not record loaded state\n"; + return 5; + } + return 0; + } + ''' + ) + plugin_code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_hccp_loader.h" + #include "tilexr_types.h" + + #include + + using namespace TileXR; + + extern "C" int TileXRCcuCollectLocalEndpointRoute( + uint32_t devicePhyId, + const TileXRCcuEndpointRouteProviderResourceWindow* window, + TileXRCcuEndpointRouteProviderRoute* route) + { + if (devicePhyId != 0x55 || window == nullptr || route == nullptr || + window->addr != 0x10000000ULL || window->tokenId != 0x1234) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + std::memset(route, 0, sizeof(*route)); + route->remoteEid[0] = 0xaa; + route->tpn = 0x10203; + route->doorbellVa = 0x1122334455667788ULL; + route->doorbellTokenId = 0x3456; + route->doorbellTokenValue = 0; + route->sqDepth = 64; + route->endpointRouteVerified = true; + return TILEXR_SUCCESS; + } + ''' + ) + + compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++") + if compiler is None: + self.skipTest("no local C++ compiler found") + with tempfile.TemporaryDirectory() as temp_dir: + temp_path = Path(temp_dir) + plugin_cpp = temp_path / "tilexr_ccu_endpoint_route_provider.cpp" + plugin_so = temp_path / "tilexr_ccu_endpoint_route_provider.so" + plugin_cpp.write_text(plugin_code, encoding="utf-8") + plugin_build = subprocess.run( + [ + compiler, + "-std=c++14", + "-shared", + "-fPIC", + "-I", + str(REPO_ROOT / "src" / "include"), + "-I", + str(REPO_ROOT / "src" / "comm"), + str(plugin_cpp), + "-o", + str(plugin_so), + ], + cwd=REPO_ROOT, + check=False, + text=True, + capture_output=True, + ) + self.assertEqual(0, plugin_build.returncode, plugin_build.stdout + plugin_build.stderr) + result = self.compile_and_run( + code, + extra_sources=[LOADER_SOURCE], + extra_link_flags=["-ldl"], + extra_env={"TILEXR_CCU_ENDPOINT_ROUTE_PROVIDER": str(plugin_so)}, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_runtime_uses_optional_tilexr_endpoint_route_provider_before_env(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_direct_runtime.h" + #undef private + + #include + #include + + using namespace TileXR; + + int EchoLocalExchangeAsPeer(const void* sendBuf, size_t sendBytes, void* recvBuf, void*) + { + if (sendBuf == nullptr || sendBytes != sizeof(TileXRCcuResourceWindowExchange)) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + const auto* local = static_cast(sendBuf); + auto* out = static_cast(recvBuf); + std::memset(out, 0, sizeof(TileXRCcuResourceWindowExchange) * 2); + out[0] = *local; + return TILEXR_SUCCESS; + } + + int main() + { + TileXRCcuDirectRuntime runtime; + runtime.initialized_ = true; + runtime.resourceWindowRegistered_ = true; + runtime.devicePhyId_ = 0x55; + runtime.options_.rank = 1; + runtime.options_.rankSize = 2; + runtime.options_.allGather = EchoLocalExchangeAsPeer; + runtime.localResourceWindow_.addr = 0x10000000ULL; + runtime.localResourceWindow_.bytes = 0x2000; + runtime.localResourceWindow_.tokenId = 0x1234; + runtime.localResourceWindow_.rawTokenId = 0x2234; + runtime.localResourceWindow_.tokenValue = 0; + + TileXRCcuDirectRuntimeReport report; + if (runtime.RefreshLocalVerifiedEndpointRoute(&report) != TILEXR_SUCCESS) { + std::cerr << "provider route was not collected: " << report.message << "\n"; + return 1; + } + + std::vector buffers; + if (runtime.ExportRemoteCcuRmaBuffers(&buffers) != TILEXR_SUCCESS) { + std::cerr << "remote export failed\n"; + return 2; + } + if (buffers.size() != 1 || !buffers[0].endpointRouteVerified || + buffers[0].remoteEid[0] != 0xaa || + buffers[0].tpn != 0x10203 || + buffers[0].doorbellVa != 0x1122334455667788ULL || + buffers[0].doorbellTokenId != 0x3456 || + buffers[0].doorbellTokenValue != 0 || + buffers[0].sqDepth != 64) { + std::cerr << "provider route was not exported\n"; + return 3; + } + return 0; + } + ''' + ) + plugin_code = textwrap.dedent( + r''' + #include "ccu/tilexr_ccu_hccp_loader.h" + #include "tilexr_types.h" + + #include + + using namespace TileXR; + + extern "C" int TileXRCcuCollectLocalEndpointRoute( + uint32_t devicePhyId, + const TileXRCcuEndpointRouteProviderResourceWindow* window, + TileXRCcuEndpointRouteProviderRoute* route) + { + if (devicePhyId != 0x55 || window == nullptr || route == nullptr || + window->addr != 0x10000000ULL || window->tokenId != 0x1234) { + return TILEXR_ERROR_PARA_CHECK_FAIL; + } + std::memset(route, 0, sizeof(*route)); + route->remoteEid[0] = 0xaa; + route->tpn = 0x10203; + route->doorbellVa = 0x1122334455667788ULL; + route->doorbellTokenId = 0x3456; + route->doorbellTokenValue = 0; + route->sqDepth = 64; + route->endpointRouteVerified = true; + return TILEXR_SUCCESS; + } + ''' + ) + + compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++") + if compiler is None: + self.skipTest("no local C++ compiler found") + with tempfile.TemporaryDirectory() as temp_dir: + temp_path = Path(temp_dir) + plugin_cpp = temp_path / "tilexr_ccu_endpoint_route_provider.cpp" + plugin_so = temp_path / "tilexr_ccu_endpoint_route_provider.so" + plugin_cpp.write_text(plugin_code, encoding="utf-8") + plugin_build = subprocess.run( + [ + compiler, + "-std=c++14", + "-shared", + "-fPIC", + "-I", + str(REPO_ROOT / "src" / "include"), + "-I", + str(REPO_ROOT / "src" / "comm"), + str(plugin_cpp), + "-o", + str(plugin_so), + ], + cwd=REPO_ROOT, + check=False, + text=True, + capture_output=True, + ) + self.assertEqual(0, plugin_build.returncode, plugin_build.stdout + plugin_build.stderr) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_specs.cpp", + ], + extra_link_flags=["-ldl"], + extra_env={"TILEXR_CCU_ENDPOINT_ROUTE_PROVIDER": str(plugin_so)}, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_runtime_collects_ra_ctx_loop_endpoint_route_before_env_fallback(self): + source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8") + + self.assertIn("TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE", source) + self.assertIn("ra_ctx_loop", source) + self.assertIn("CollectLocalEndpointRouteWithRaCtx(&route)", source) + + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_direct_runtime.h" + #undef private + + #include + #include + + using namespace TileXR; + + constexpr uint32_t kRawDbTokenId = 0x345600U; + constexpr uint32_t kTokenValue = 0x7799U; + constexpr uint64_t kTpHandle = 0xabcddcbaULL; + constexpr uint64_t kCcuResourceBase = 0x10000000ULL; + constexpr uint64_t kHcommWqeBasicBlockOffset = 0x1000000ULL; + constexpr uint64_t kHcommSqBufferSize = 256ULL * 1024ULL; + constexpr uint32_t kDirectLoopJettyCtxId = 0; + constexpr uint64_t kExpectedSqVa = + kCcuResourceBase + kHcommWqeBasicBlockOffset + + static_cast(kDirectLoopJettyCtxId) * kHcommSqBufferSize; + + int g_cqCreateCalls = 0; + int g_qpCreateCalls = 0; + int g_tpInfoCalls = 0; + int g_reqResultCalls = 0; + int g_qpImportCalls = 0; + int g_qpBindCalls = 0; + int g_qpUnbindCalls = 0; + int g_qpUnimportCalls = 0; + int g_qpDestroyCalls = 0; + int g_cqDestroyCalls = 0; + TileXRCcuHccpQpCreateAttr g_lastQpAttr {}; + TileXRCcuHccpGetTpCfg g_lastTpCfg {}; + TileXRCcuHccpQpImportInfo g_lastImportInfo {}; + + int FakeCqCreate(void* ctx, TileXRCcuHccpCqInfo* info, void** cqHandle) + { + if (ctx != reinterpret_cast(0x1000) || info == nullptr || cqHandle == nullptr) { + return -1; + } + if (info->in.depth != 64 || info->in.ub.mode != 2) { + return -2; + } + ++g_cqCreateCalls; + *cqHandle = reinterpret_cast(0x2000); + return 0; + } + + int FakeCqDestroy(void*, void* cqHandle) + { + if (cqHandle == reinterpret_cast(0x2000)) { + ++g_cqDestroyCalls; + } + return 0; + } + + int FakeQpCreate( + void* ctx, + TileXRCcuHccpQpCreateAttr* attr, + TileXRCcuHccpQpCreateInfo* info, + void** qpHandle) + { + if (ctx != reinterpret_cast(0x1000) || attr == nullptr || info == nullptr || + qpHandle == nullptr) { + return -3; + } + g_lastQpAttr = *attr; + if (attr->scqHandle != reinterpret_cast(0x2000) || + attr->rcqHandle != reinterpret_cast(0x2000) || + attr->sqDepth != 8 || + attr->rqDepth != TILEXR_CCU_HCCP_RQ_DEPTH_DEFAULT || + attr->transportMode != TILEXR_CCU_HCCP_TRANSPORT_MODE_RM || + attr->ub.mode != static_cast(TILEXR_CCU_HCCP_JETTY_MODE_CCU) || + attr->ub.jettyId != 1024 || + attr->ub.tokenIdHandle != reinterpret_cast(0x3000) || + attr->ub.tokenValue != kTokenValue || + attr->ub.flag.bs.shareJfr != 1 || + attr->ub.jfsFlag.bs.errorSuspend != 1 || + attr->ub.extMode.cstmFlag.bs.sqCstm != 1 || + attr->ub.extMode.sqebbNum != 8 || + attr->ub.extMode.sq.buffVa != kExpectedSqVa || + attr->ub.extMode.sq.buffSize != 8 * 4 * 64) { + return -4; + } + for (uint32_t i = 0; i < TILEXR_CCU_HCCP_QP_KEY_BYTES; ++i) { + info->key.value[i] = static_cast(0x80U + i); + } + info->key.size = TILEXR_CCU_HCCP_QP_KEY_BYTES; + info->ub.dbAddr = 0x1122334455667788ULL; + info->ub.dbTokenId = kRawDbTokenId; + ++g_qpCreateCalls; + *qpHandle = reinterpret_cast(0x4000); + return 0; + } + + int FakeQpDestroy(void* qpHandle) + { + if (qpHandle == reinterpret_cast(0x4000)) { + ++g_qpDestroyCalls; + } + return 0; + } + + int FakeGetTpInfoListAsync( + void* ctx, + TileXRCcuHccpGetTpCfg* cfg, + TileXRCcuHccpTpInfo infoList[], + uint32_t* num, + void** reqHandle) + { + if (ctx != reinterpret_cast(0x1000) || cfg == nullptr || infoList == nullptr || + num == nullptr || reqHandle == nullptr || *num != 1) { + return -5; + } + g_lastTpCfg = *cfg; + if (cfg->flag.bs.rtp != 1 || cfg->transMode != TILEXR_CCU_HCCP_TRANSPORT_MODE_RM || + std::memcmp(cfg->localEid.raw, cfg->peerEid.raw, TILEXR_CCU_EID_BYTES) != 0) { + return -6; + } + infoList[0].tpHandle = kTpHandle; + ++g_tpInfoCalls; + *reqHandle = reinterpret_cast(0x5000); + return 0; + } + + int FakeGetAsyncReqResult(void* reqHandle, int* reqResult) + { + if (reqHandle != reinterpret_cast(0x5000) || reqResult == nullptr) { + return -7; + } + *reqResult = 0; + ++g_reqResultCalls; + return 0; + } + + int FakeQpImport(void* ctx, TileXRCcuHccpQpImportInfo* info, void** remoteQpHandle) + { + if (ctx != reinterpret_cast(0x1000) || info == nullptr || remoteQpHandle == nullptr) { + return -8; + } + g_lastImportInfo = *info; + if (info->in.key.size != TILEXR_CCU_HCCP_QP_KEY_BYTES || + info->in.key.value[0] != 0x80 || + info->in.ub.mode != TILEXR_CCU_HCCP_JETTY_IMPORT_MODE_EXP || + info->in.ub.tokenValue != kTokenValue || + info->in.ub.policy != TILEXR_CCU_HCCP_JETTY_GRP_POLICY_RR || + info->in.ub.type != TILEXR_CCU_HCCP_TARGET_TYPE_JETTY || + info->in.ub.flag.bs.tokenPolicy != TILEXR_CCU_HCCP_TOKEN_POLICY_PLAIN_TEXT || + info->in.ub.expImportCfg.tpHandle != kTpHandle || + info->in.ub.expImportCfg.peerTpHandle != kTpHandle || + info->in.ub.expImportCfg.txPsn == 0 || + info->in.ub.expImportCfg.rxPsn == 0 || + info->in.ub.tpType != TILEXR_CCU_HCCP_TP_TYPE_RTP) { + return -9; + } + info->out.ub.tpn = 0; + ++g_qpImportCalls; + *remoteQpHandle = reinterpret_cast(0x6000); + return 0; + } + + int FakeQpUnimport(void*, void* remoteQpHandle) + { + if (remoteQpHandle == reinterpret_cast(0x6000)) { + ++g_qpUnimportCalls; + } + return 0; + } + + int FakeQpBind(void* qpHandle, void* remoteQpHandle) + { + if (qpHandle != reinterpret_cast(0x4000) || + remoteQpHandle != reinterpret_cast(0x6000)) { + return -10; + } + ++g_qpBindCalls; + return -11; + } + + int FakeQpUnbind(void* qpHandle) + { + if (qpHandle == reinterpret_cast(0x4000)) { + ++g_qpUnbindCalls; + } + return 0; + } + + int main() + { + TileXRCcuDirectRuntime runtime; + runtime.initialized_ = true; + runtime.resourceWindowRegistered_ = true; + runtime.devicePhyId_ = 0x55; + runtime.options_.rank = 0; + runtime.options_.rankSize = 1; + runtime.localResourceWindow_.addr = kCcuResourceBase; + runtime.localResourceWindow_.bytes = TILEXR_CCU_RESOURCE_WINDOW_BYTES; + runtime.localResourceWindow_.tokenId = 0x3456; + runtime.localResourceWindow_.rawTokenId = 0x345600; + runtime.localResourceWindow_.tokenValue = kTokenValue; + runtime.localResourceWindow_.raCtxHandle = reinterpret_cast(0x1000); + runtime.localResourceWindow_.tokenIdHandle = reinterpret_cast(0x3000); + runtime.localResourceWindow_.raCtxRegistered = true; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + runtime.localResourceWindow_.eid[i] = static_cast(0xa0U + i); + } + + runtime.loader_.RaCtxCqCreate = FakeCqCreate; + runtime.loader_.RaCtxCqDestroy = FakeCqDestroy; + runtime.loader_.RaCtxQpCreate = FakeQpCreate; + runtime.loader_.RaCtxQpDestroy = FakeQpDestroy; + runtime.loader_.RaGetTpInfoListAsync = FakeGetTpInfoListAsync; + runtime.loader_.RaGetAsyncReqResult = FakeGetAsyncReqResult; + runtime.loader_.RaCtxQpImport = FakeQpImport; + runtime.loader_.RaCtxQpUnimport = FakeQpUnimport; + runtime.loader_.RaCtxQpBind = FakeQpBind; + runtime.loader_.RaCtxQpUnbind = FakeQpUnbind; + + TileXRCcuDirectRuntimeReport report; + if (runtime.RefreshLocalVerifiedEndpointRoute(&report) != TILEXR_SUCCESS) { + std::cerr << "ra ctx loop route collection failed: " << report.message << "\n"; + return 1; + } + if (!runtime.localVerifiedEndpointRouteValid_) { + std::cerr << "verified route flag not set\n"; + return 2; + } + const auto& route = runtime.localVerifiedEndpointRoute_; + if (!route.endpointRouteVerified || + route.remoteEid[0] != 0xa0 || + route.tpn != 0 || + route.doorbellVa != 0x1122334455667788ULL || + route.doorbellTokenId != (kRawDbTokenId >> 8) || + route.doorbellTokenValue != kTokenValue || + route.sqDepth != 8) { + std::cerr << "verified route mismatch\n"; + return 3; + } + if (g_cqCreateCalls != 1 || g_qpCreateCalls != 1 || g_tpInfoCalls != 1 || + g_reqResultCalls != 1 || g_qpImportCalls != 1 || g_qpBindCalls != 0) { + std::cerr << "unexpected RA call counts\n"; + return 4; + } + + runtime.Shutdown(); + if (g_qpUnbindCalls != 0 || g_qpUnimportCalls != 1 || + g_qpDestroyCalls != 1 || g_cqDestroyCalls != 1) { + std::cerr << "endpoint route handles were not released\n"; + return 5; + } + return 0; + } + ''' + ) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_specs.cpp", + ], + extra_link_flags=["-ldl"], + extra_env={ + "TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE": "ra_ctx_loop", + }, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_runtime_retries_ra_ctx_endpoint_route_after_async_result_failure(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_direct_runtime.h" + #undef private + + #include + #include + + using namespace TileXR; + + constexpr uint32_t kRawDbTokenId = 0x456700U; + constexpr uint32_t kTokenValue = 0x8811U; + constexpr uint64_t kCcuResourceBase = 0x11000000ULL; + + int g_cqCreateCalls = 0; + int g_qpCreateCalls = 0; + int g_tpInfoCalls = 0; + int g_reqResultCalls = 0; + int g_qpImportCalls = 0; + int g_qpDestroyCalls = 0; + int g_cqDestroyCalls = 0; + + int FakeCqCreate(void*, TileXRCcuHccpCqInfo*, void** cqHandle) + { + if (cqHandle == nullptr) { + return -1; + } + ++g_cqCreateCalls; + *cqHandle = reinterpret_cast(0x2000 + g_cqCreateCalls); + return 0; + } + + int FakeCqDestroy(void*, void* cqHandle) + { + if (cqHandle != nullptr) { + ++g_cqDestroyCalls; + } + return 0; + } + + int FakeQpCreate( + void*, + TileXRCcuHccpQpCreateAttr*, + TileXRCcuHccpQpCreateInfo* info, + void** qpHandle) + { + if (info == nullptr || qpHandle == nullptr) { + return -2; + } + ++g_qpCreateCalls; + for (uint32_t i = 0; i < TILEXR_CCU_HCCP_QP_KEY_BYTES; ++i) { + info->key.value[i] = static_cast(0x40U + i); + } + info->key.size = TILEXR_CCU_HCCP_QP_KEY_BYTES; + info->ub.dbAddr = 0x1222333444555666ULL; + info->ub.dbTokenId = kRawDbTokenId; + *qpHandle = reinterpret_cast(0x4000 + g_qpCreateCalls); + return 0; + } + + int FakeQpDestroy(void* qpHandle) + { + if (qpHandle != nullptr) { + ++g_qpDestroyCalls; + } + return 0; + } + + int FakeGetTpInfoListAsync( + void*, + TileXRCcuHccpGetTpCfg*, + TileXRCcuHccpTpInfo infoList[], + uint32_t* num, + void** reqHandle) + { + if (infoList == nullptr || num == nullptr || reqHandle == nullptr || *num != 1) { + return -3; + } + ++g_tpInfoCalls; + infoList[0].tpHandle = 0xabcd0000ULL + g_tpInfoCalls; + *reqHandle = reinterpret_cast(0x5000 + g_tpInfoCalls); + return 0; + } + + int FakeGetAsyncReqResult(void*, int* reqResult) + { + if (reqResult == nullptr) { + return -4; + } + ++g_reqResultCalls; + *reqResult = (g_reqResultCalls <= 3) ? 528101 : 0; + return 0; + } + + int FakeQpImport(void*, TileXRCcuHccpQpImportInfo* info, void** remoteQpHandle) + { + if (info == nullptr || remoteQpHandle == nullptr) { + return -5; + } + ++g_qpImportCalls; + info->out.ub.tpn = 0x7000 + g_qpImportCalls; + *remoteQpHandle = reinterpret_cast(0x6000 + g_qpImportCalls); + return 0; + } + + int FakeQpUnimport(void*, void*) + { + return 0; + } + + int main() + { + TileXRCcuDirectRuntime runtime; + runtime.initialized_ = true; + runtime.resourceWindowRegistered_ = true; + runtime.devicePhyId_ = 0x55; + runtime.options_.rank = 0; + runtime.options_.rankSize = 1; + runtime.localResourceWindow_.addr = kCcuResourceBase; + runtime.localResourceWindow_.bytes = TILEXR_CCU_RESOURCE_WINDOW_BYTES; + runtime.localResourceWindow_.tokenId = 0x4567; + runtime.localResourceWindow_.rawTokenId = kRawDbTokenId; + runtime.localResourceWindow_.tokenValue = kTokenValue; + runtime.localResourceWindow_.raCtxHandle = reinterpret_cast(0x1000); + runtime.localResourceWindow_.tokenIdHandle = reinterpret_cast(0x3000); + runtime.localResourceWindow_.raCtxRegistered = true; + for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) { + runtime.localResourceWindow_.eid[i] = static_cast(0xb0U + i); + } + + runtime.loader_.RaCtxCqCreate = FakeCqCreate; + runtime.loader_.RaCtxCqDestroy = FakeCqDestroy; + runtime.loader_.RaCtxQpCreate = FakeQpCreate; + runtime.loader_.RaCtxQpDestroy = FakeQpDestroy; + runtime.loader_.RaGetTpInfoListAsync = FakeGetTpInfoListAsync; + runtime.loader_.RaGetAsyncReqResult = FakeGetAsyncReqResult; + runtime.loader_.RaCtxQpImport = FakeQpImport; + runtime.loader_.RaCtxQpUnimport = FakeQpUnimport; + + TileXRCcuDirectRuntimeReport report; + const int ret = runtime.RefreshLocalVerifiedEndpointRoute(&report); + if (ret != TILEXR_SUCCESS) { + std::cerr << "ra ctx endpoint route retry failed: " << report.message + << " ret=" << ret << "\n"; + return 1; + } + if (!runtime.localVerifiedEndpointRouteValid_) { + std::cerr << "verified endpoint route flag not set after retry\n"; + return 2; + } + const auto& route = runtime.localVerifiedEndpointRoute_; + if (!route.endpointRouteVerified || + route.remoteEid[0] != 0xb0 || + route.doorbellVa != 0x1222333444555666ULL || + route.doorbellTokenId != (kRawDbTokenId >> 8) || + route.doorbellTokenValue != kTokenValue || + route.sqDepth != 8) { + std::cerr << "verified endpoint route mismatch after retry\n"; + return 3; + } + if (g_cqCreateCalls != 4 || g_qpCreateCalls != 4 || g_tpInfoCalls != 4 || + g_reqResultCalls != 4 || g_qpImportCalls != 1) { + std::cerr << "unexpected retry call counts" + << " cqCreate=" << g_cqCreateCalls + << " qpCreate=" << g_qpCreateCalls + << " tpInfo=" << g_tpInfoCalls + << " reqResult=" << g_reqResultCalls + << " qpImport=" << g_qpImportCalls << "\n"; + return 4; + } + if (g_qpDestroyCalls != 3 || g_cqDestroyCalls != 3) { + std::cerr << "failed attempt was not cleaned up before retry" + << " qpDestroy=" << g_qpDestroyCalls + << " cqDestroy=" << g_cqDestroyCalls << "\n"; + return 5; + } + + runtime.Shutdown(); + if (g_qpDestroyCalls != 4 || g_cqDestroyCalls != 4) { + std::cerr << "successful retry handles were not released on shutdown" + << " qpDestroy=" << g_qpDestroyCalls + << " cqDestroy=" << g_cqDestroyCalls << "\n"; + return 6; + } + return 0; + } + ''' + ) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_specs.cpp", + ], + extra_link_flags=["-ldl"], + extra_env={ + "TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE": "ra_ctx_loop", + }, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + def test_direct_runtime_fails_closed_when_configured_endpoint_route_provider_cannot_load(self): + code = textwrap.dedent( + r''' + #define private public + #include "ccu/tilexr_ccu_direct_runtime.h" + #undef private + + #include + + using namespace TileXR; + + int main() + { + TileXRCcuDirectRuntime runtime; + runtime.initialized_ = true; + runtime.resourceWindowRegistered_ = true; + runtime.devicePhyId_ = 0x55; + runtime.options_.rank = 1; + runtime.options_.rankSize = 2; + runtime.localResourceWindow_.addr = 0x10000000ULL; + runtime.localResourceWindow_.bytes = 0x2000; + runtime.localResourceWindow_.tokenId = 0x1234; + runtime.localResourceWindow_.rawTokenId = 0x2234; + runtime.localResourceWindow_.tokenValue = 0; + + TileXRCcuDirectRuntimeReport report; + const int ret = runtime.RefreshLocalVerifiedEndpointRoute(&report); + if (ret == TILEXR_SUCCESS) { + std::cerr << "configured bad provider fell back to env route\n"; + return 1; + } + if (runtime.localVerifiedEndpointRouteValid_) { + std::cerr << "bad provider left a verified route\n"; + return 2; + } + if (report.message.find("failed to load direct CCU endpoint route provider") == std::string::npos) { + std::cerr << "unexpected failure message: " << report.message << "\n"; + return 3; + } + return 0; + } + ''' + ) + result = self.compile_and_run( + code, + extra_sources=[ + DIRECT_RUNTIME_SOURCE, + LOADER_SOURCE, + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp", + REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_specs.cpp", + ], + extra_link_flags=["-ldl"], + extra_env={ + "TILEXR_CCU_ENDPOINT_ROUTE_PROVIDER": "/tmp/tilexr_missing_endpoint_route_provider.so", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_EID_RANK1": "aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_TPN_RANK1": "0x10203", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_VA_RANK1": "0x1122334455667788", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_TOKEN_ID_RANK1": "0x3456", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_DOORBELL_TOKEN_VALUE_RANK1": "0", + "TILEXR_CCU_DIRECT_LOCAL_ENDPOINT_SQ_DEPTH_RANK1": "64", + }, + ) + + self.assertEqual("", result.stderr) + self.assertEqual(0, result.returncode, result.stdout + result.stderr) + + +if __name__ == "__main__": + unittest.main()