diff --git a/docs/CCU_DIRECT_RUNTIME.md b/docs/CCU_DIRECT_RUNTIME.md
new file mode 100644
index 00000000..ef984504
--- /dev/null
+++ b/docs/CCU_DIRECT_RUNTIME.md
@@ -0,0 +1,503 @@
+# TileXR Direct CCU 技术说明
+
+本文记录 `direct-ccu-rebased` 分支中 TileXR 自研 Direct CCU 路径的技术原理、运行流程、执行方法和测试方法。它面向继续开发和问题定位，不作为外部稳定 API 文档。
+
+## 1. 当前边界
+
+Direct CCU 当前属于 `src/comm/ccu` 下的内部后端能力，由 `TileXRComm` 按需持有 `TileXRCcuBackend`。它的目标是在不链接 hcomm/HCCL 私有 CCU producer 的前提下，由 TileXR host 侧完成 CCU repository、mission、lower-layer 资源和 submit task 的准备，并最终通过 CANN runtime 的 `rtCCULaunch` 下发任务。
+
+当前需要明确的边界如下：
+
+- 生产代码不得链接或包含 hcomm/HCCL 私有 CCU producer 接口。
+- 对外安装头文件当前不暴露 `TileXRDirectCcu*`、`PrepareDirectCcu`、`SubmitPrepared` 等 Direct CCU C API。
+- `TileXRCcuCollectivePlanner::Supports()` 目前仍返回 `false`，通用 collective/alltoall 尚未接入为正式 TileXR collective 后端。
+- 已维护的数据面验证路径是 direct CCU P2P copy probe，基于 CCU memory-copy microcode，不走 alltoall。
+- 默认非 P2P barrier/smoke 只能作为安装和提交链路诊断，不能替代 P2P copy 的数据正确性证明。
+
+## 2. 总体架构
+
+核心对象关系：
+
+```text
+TileXRComm
+  |
+  +-- TileXRCcuBackend
+        |
+        +-- TileXRCcuRuntimeSession
+        |     |
+        |     +-- TileXRCcuDirectRuntime
+        |
+        +-- TileXRCcuCollectivePlanner
+        |
+        +-- TileXRCcuExecutor
+```
+
+主要模块职责：
+
+- `tilexr_ccu_backend.*`：Direct CCU 后端门面，封装 runtime session、planner 和 executor。
+- `tilexr_ccu_runtime_session.*`：管理 rank、rankSize、device、socket/thread allgather 和 direct runtime 可用状态。
+- `tilexr_ccu_direct_runtime.*`：动态加载 HCCP/RA/runtime 符号，初始化底层 RA/HDC/CCU TLV，注册 resource window，导出本端和远端 transport 信息。
+- `tilexr_ccu_resource_allocator.*`：按 driver basic info 解出的资源窗口分配 mission、repository instruction、local/remote XN、GSA、CKE、channel 等资源。
+- `tilexr_ccu_collective_planner.*`：把 runtime session、资源分配、lower-layer plan、repository 安装和 prepared submit task 串起来。
+- `tilexr_ccu_direct_orchestrator.*`：执行完整 direct install attempt，包括资源规格解码、资源分配、lower-layer plan、launch package、manifest、hardware install 和 submit task 生成。
+- `tilexr_ccu_install_provider.*`：定义硬件安装需求、证据和校验，执行 repository/lower-layer/mission 等安装步骤。
+- `tilexr_ccu_repository.*`：构造 CCU instruction repository image，并负责把 repository image 安装到设备侧。
+- `tilexr_ccu_barrier_program.*`、`tilexr_ccu_memory_program.*`、`tilexr_ccu_microcode.*`：生成 CCU barrier 和 memory-copy 指令。
+- `tilexr_ccu_runtime.*`：把 `TileXRCcuTask` 映射为 runtime task，并调用 `rtCCULaunch`。
+
+## 3. 初始化与生命周期
+
+`TileXRComm::Init()` 和 `TileXRComm::InitThread()` 结束前会调用 `InitCcuBackendIfEnabled()`。是否启用由环境变量控制：
+
+```bash
+export TILEXR_ENABLE_CCU_BACKEND=1
+```
+
+启用后流程如下：
+
+1. `TileXRComm::InitCcuBackend()` 创建 `TileXRCcuBackend`。
+2. `TileXRCcuBackend::Init()` 重置 planner，并调用 `TileXRCcuRuntimeSession::Init()`。
+3. `TileXRCcuRuntimeSession::Init()` 对单 rank communicator 直接跳过 direct runtime。
+4. 多 rank 时创建 `TileXRCcuDirectRuntime`，传入 rank、rankSize、device 和 allgather 回调。
+5. `TileXRCcuDirectRuntime::Init()` 动态加载底层符号，解析逻辑 device 到物理 device，初始化 RA/HDC 和 CCU TLV。
+6. runtime 初始化成功后刷新 direct CCU basic info，缓存 die、resource address、mission/resource 范围等基础信息。
+
+失败策略是降级而不是让 communicator 初始化失败：
+
+- runtime 初始化失败会记录 warning，并将 direct CCU 后端关闭。
+- 失败状态按 `devId` 记录，避免一个 device 初始化失败污染同进程其他 device。
+- 后续同 device 再初始化会直接跳过，并保留明确的 unavailable message。
+
+`TileXRComm::Destroy()` 会关闭 CCU backend。`TileXRCcuBackend::Shutdown()` 会先重置 planner，再关闭 runtime session；runtime session 会释放 direct runtime、basic info 缓存和 allgather 轮次状态。
+
+## 4. Direct Runtime 原理
+
+`TileXRCcuDirectRuntime` 是 TileXR 和底层 driver/RA/HCCP/runtime 交互的边界，承担四类工作。
+
+### 4.1 符号与设备初始化
+
+runtime 动态解析所需符号，不在 `tile-comm` 链接期引入 hcomm/HCCL 私有依赖。初始化阶段会：
+
+- 加载 HCCP/RA 相关入口。
+- 选择 direct CCU HDC 类型。
+- 解析 `logicDevId -> devicePhyId`。
+- 初始化 RA/HDC。
+- 初始化 CCU TLV。
+- 创建 driver adapter，用于 basic info 查询和后续 install。
+
+如果任一关键步骤失败，runtime 返回不可用状态，并由 session 记录 device-scoped failure。
+
+### 4.2 Resource Window 注册
+
+CCU lower-layer/repository 安装需要一块可被 peer 识别的 resource window。当前推荐使用 RA ctx 模式：
+
+```bash
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE=ra_ctx
+```
+
+RA ctx 模式的关键步骤：
+
+1. 通过 `RaGetDevEidInfoNum` / `RaGetDevEidInfoList` 获取可用 EID。
+2. 通过 `TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX` 或 rank 级变量选择 EID。
+3. 调用 `RaCtxInit` 创建 RA ctx。
+4. 通过 `RaCtxTokenIdAlloc` 分配 token id。
+5. 对 CCU resource address 做页对齐后调用 `RaCtxLmemRegister`。
+6. 导出 `addr/bytes/tokenId/rawTokenId/tokenValue/eid/raCtxHandle`。
+
+常用 EID 选择变量：
+
+```bash
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX=3
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_RANK0=3
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_RANK1=3
+```
+
+### 4.3 Peer 信息交换
+
+runtime session 为 direct runtime 提供 allgather：
+
+- 有 socket exchange 时使用 `TileXRSockExchange::AllGather()`。
+- thread mode 或测试模式下没有 socket exchange 时，使用进程内 thread allgather。
+
+allgather 用于交换：
+
+- 本端 resource window token 和 EID。
+- endpoint route 信息。
+- remote XN/CKE/channel 绑定所需的 peer 资源窗口。
+- P2P copy 端点中的 source/destination address 和 process token。
+
+thread allgather 带超时和失败中止逻辑，避免单 rank runtime init 失败后其他 rank 永久等待。
+
+### 4.4 Endpoint Route 与 Remote Buffer
+
+`ExportRemoteCcuRmaBuffers()` 会把本端 resource window 信息 allgather 到所有 rank，并为每个 peer 生成 remote CCU buffer 信息。可用 RA ctx endpoint route 时，会进一步：
+
+- 查询本端到 peer EID 的 TP handle。
+- 交换 TP handle。
+- 导入 peer QP。
+- 使用导入得到的 TPN 和本端 doorbell token 构造 channel route。
+
+这里存在一个重要细节：channel 中使用的 remote EID 采用 hcomm 兼容语义，导入后会使用反向 EID 表示。
+
+## 5. 资源模型与所有权
+
+Direct CCU 执行前需要把硬件资源划分为明确的窗口。资源来自 `TileXRCcuBasicInfo`，再由 `TileXRCcuDecodeBasicInfo()` 和 `TileXRCcuBuildResourceSpec()` 转换为 `TileXRCcuResourceSpec`。
+
+主要资源：
+
+- `mission`：mission id 和 mission key。
+- `repository`：CCU instruction repository 中的指令槽。
+- `localXn`：本端 CCU XN 资源。
+- `localGsa`：本端 GSA 资源，P2P memory copy 用于装载地址、token 和长度。
+- `remoteXn`：绑定 peer 侧 XN 的本地表示。
+- `notifyCke`、`localWaitCke`、`remoteNotifyCke`：同步/完成通知使用的 CKE。
+- `channels`：指向 peer resource window/endpoint route 的 channel 绑定。
+
+`TileXRCcuResourceAllocator` 使用 receipt 记录每次分配。当前释放策略是严格 LIFO，用来保证 cursor 回退不会和仍在使用的资源重叠。这个设计的目的不是做复杂资源池，而是在多次 prepare/launch 场景中先保证资源窗口所有权清晰、可诊断、可回收。
+
+## 6. Lower-layer Install 原理
+
+lower-layer install 负责把 CCU 执行需要的资源上下文安装到 driver 可识别的位置。主要输入来自：
+
+- basic info 解码出的资源窗口。
+- allocator 分配出的 mission/XN/CKE/channel range。
+- 本端 resource window token。
+- peer resource window token。
+- verified endpoint route。
+
+准备流程：
+
+1. `PrepareDirectCcuLowerLayerTemplateFromAllocation()` 先注册本端 resource window。
+2. 导出本端 token 后，如果已经有 verified endpoint route 就配置到 runtime；否则尝试自动采集。
+3. `ExportRemoteCcuRmaBuffers()` 交换 peer token 和 endpoint 信息。
+4. `ExchangeDirectCcuRemoteNotifyCke()` 交换 peer 本地 XN/CKE/channel 分配结果。
+5. `TileXRCcuBuildLowerLayerTransportTemplate()` 构造 transport snapshot。
+6. `TileXRCcuBuildLowerLayerInstallPlanFromTransportSnapshot()` 生成安装计划。
+
+安装计划覆盖的 surface 包括 local XN、remote XN binding、notify CKE、channel binding、repository 和 mission/key。`tilexr_ccu_install_provider` 会为这些 surface 建立 manifest 和 install evidence。只有 install evidence 与 launch package fingerprint、device、rank、provider 和资源范围匹配时，后续 prepared submit task 才被认为可提交。
+
+## 7. Repository、Mission 与 Launch Package
+
+Direct CCU 的 host 下发不是直接下发一段普通 kernel，而是下发 CCU mission 和 instruction repository。
+
+完整 install attempt 的关键步骤在 `RunDirectInstallAttemptImpl()` 中：
+
+1. 校验 basic info、provider 和 repository install 输入。
+2. 解码 basic info，生成 resource spec。
+3. 分配 mission、repository、XN、GSA、CKE、channel 等资源。
+4. 如是 P2P memory copy，调整 local XN/GSA/remote XN 资源布局。
+5. 准备 lower-layer install plan。
+6. 将 producer plan 与 lower-layer proof 对齐，确保 sync resource 使用真实 peer 资源。
+7. 构造 CCU program 和 repository image。
+8. 绑定 launch package 的 device/rank/provider scope。
+9. 构造 install manifest。
+10. 执行 `TileXRCcuInstallHardware()`，安装 repository/lower-layer/mission。
+11. 根据 install evidence 生成 `submitTasks`。
+
+`TileXRCcuTask` 是最终提交给 runtime 的任务描述，关键字段包括：
+
+- `dieId`
+- `missionId`
+- `key`
+- `instStartId`
+- `instCnt`
+- `timeout`
+- `argSize`
+- `args[]`
+
+提交时 `TileXRCcuSubmitPreparedTasks()` 会逐个任务调用 `TileXRCcuSubmitTaskWithReport()`，后者在 `tilexr_ccu_runtime.cpp` 中映射到 `rtCCULaunch()`。如果中途某个 task 失败，submit report 会记录已提交数量和失败 task 的 mission/key/instruction/args 诊断信息。
+
+## 8. P2P CCU Copy 执行原理
+
+当前硬件数据面验证使用 direct CCU P2P copy，不走 alltoall。它验证的是 CCU 指令真实从 peer device memory 读写数据，而不是只验证 host marker 或 IPC 同步。
+
+### 8.1 端点准备
+
+每个 rank 分配两块 device buffer：
+
+- `source`：写入 rank 相关的数据 pattern。
+- `destination`：初始化为固定填充值。
+
+然后通过 `rtUbDevQueryInfo(QUERY_PROCESS_TOKEN, ...)` 查询 source/destination 的 process token，打包为 CCU memory token。所有 rank 通过 session allgather 交换端点：
+
+```text
+rank -> {sourceAddr, sourceToken, destinationAddr, destinationToken, bytes}
+```
+
+### 8.2 方向语义
+
+`TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_DIRECTION` 支持两种方向：
+
+- `remote_to_local`：active rank 从 peer source 拷贝到本 rank destination，active rank 校验本地 destination。
+- `local_to_remote`：active rank 从本 rank source 拷贝到 peer destination，inactive rank 等待 done gate 后校验本地 destination。
+
+只有 active rank 真正 submit direct CCU task；inactive rank 会打印 `p2pCcuCopy skipped`，但仍通过 done gate 参与同步，必要时校验自己的 destination。
+
+### 8.3 Memory-copy Microcode
+
+P2P copy 会构造一个 memory copy program，典型使用 7 条 CCU 指令：
+
+- 把 local/remote address、token、length 装入 GSA/XN。
+- 根据方向生成 `TransRmtMemToLocMem` 或 `TransLocMemToRmtMem`。
+- 使用 CKE 完成通知，host 侧随后 `aclrtSynchronizeStream()` 等待 stream 完成。
+
+通过 `TILEXR_CCU_DIRECT_TRACE=1` 可在日志中看到解码后的 CCU 指令，例如：
+
+```text
+decoded=TransRmtMemToLocMem
+decoded=TransLocMemToRmtMem
+```
+
+## 9. 执行方法
+
+### 9.1 构建
+
+在 NPU 服务器上执行：
+
+```bash
+cd /home/tileXR
+source scripts/common_env.sh
+cmake --build build_ccu_direct --target tile-comm -j4
+```
+
+如果 build 目录不同，需要同步修改 `TILEXR_TILE_COMM_LIB`。
+
+### 9.2 两卡 P2P Copy Smoke
+
+以下命令是当前推荐的 direct CCU 数据面 smoke。设备号按实际健康设备调整。
+
+```bash
+cd /home/tileXR
+source scripts/common_env.sh
+
+export TILEXR_TILE_COMM_LIB=/home/tileXR/build_ccu_direct/src/comm/libtile-comm.so
+export TILEXR_CCU_SMOKE_DEVICES=3,2
+export TILEXR_RUN_CCU_DIRECT_SMOKE_PROBE=1
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY=1
+export TILEXR_CCU_DIRECT_SMOKE_EXPECT_P2P_CCU_COPY=1
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_BYTES=64
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_ACTIVE_RANK=0
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_RANK0=3
+export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_RANK1=3
+export TILEXR_CCU_DIRECT_SMOKE_READY_TIMEOUT_MS=180000
+export TILEXR_CCU_DIRECT_SMOKE_SUBMIT=1
+export TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_ON_PREPARE_FAILURE=0
+export TILEXR_CCU_SMOKE_TIMEOUT=180
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_DIRECTION=remote_to_local
+
+timeout 420s bash tests/ccu/run_tilexr_ccu_direct_smoke.sh
+```
+
+反方向验证：
+
+```bash
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_DIRECTION=local_to_remote
+timeout 420s bash tests/ccu/run_tilexr_ccu_direct_smoke.sh
+```
+
+预期关键日志：
+
+```text
+tilexr_ccu_direct_smoke prepare ret=0 ... installSucceeded=1 ... submitReady=1
+tilexr_ccu_direct_smoke submit ret=0 ... submitted=1
+tilexr_ccu_direct_smoke p2pCcuCopy ... mismatches=0 ... passed=1
+tilexr_ccu_direct_smoke_runner success
+```
+
+### 9.3 覆盖 `TileXRComm` 自动初始化路径
+
+默认 P2P smoke 可以只用 direct CCU internal init。若要覆盖 `TileXRComm::Init()` 中的 backend auto-init 路径，增加：
+
+```bash
+export TILEXR_ENABLE_CCU_BACKEND=1
+export TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT=0
+timeout 420s bash tests/ccu/run_tilexr_ccu_direct_smoke.sh
+```
+
+### 9.4 Dry-run 参数检查
+
+不触碰 ACL/NPU，只检查 runner 推导出的 repository/task 参数：
+
+```bash
+export TILEXR_RUN_CCU_DIRECT_SMOKE_PROBE=1
+export TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY=1
+export TILEXR_CCU_DIRECT_SMOKE_DRY_RUN=1
+bash tests/ccu/run_tilexr_ccu_direct_smoke.sh
+```
+
+## 10. 测试矩阵
+
+### 10.1 本地单元测试
+
+推荐先跑 CCU 相关 Python 测试：
+
+```bash
+python3 -m unittest \
+  tests.ccu.test_tilexr_ccu_resource_allocator \
+  tests.ccu.test_tilexr_ccu_lower_layer_plan_builder \
+  tests.ccu.test_tilexr_ccu_backend_boundary \
+  tests.ccu.test_tilexr_ccu_direct_orchestrator \
+  tests.ccu.test_tilexr_ccu_direct_smoke_probe \
+  tests.ccu.test_tilexr_ccu_direct_smoke_runner \
+  tests.ccu.test_tilexr_ccu_public_comm_api
+```
+
+覆盖更完整的 CCU suite：
+
+```bash
+python3 -m unittest discover tests/ccu
+```
+
+### 10.2 依赖边界检查
+
+每次修改 `src/comm/ccu` 后应确认 `tile-comm` 没有引入 hcomm/HCCL 私有 CCU 依赖：
+
+```bash
+source scripts/common_env.sh
+bash tests/ccu/check_tile_comm_no_hcomm_deps.sh build/src/comm/libtile-comm.so
+```
+
+### 10.3 硬件 Smoke
+
+硬件 smoke 由 `tests/ccu/run_tilexr_ccu_direct_smoke.sh` 启动两个 rank 进程。runner 默认是安全的，必须设置：
+
+```bash
+export TILEXR_RUN_CCU_DIRECT_SMOKE_PROBE=1
+```
+
+runner 会：
+
+- 编译 `tests/ccu/ccu_tilexr_direct_smoke_probe.cpp`。
+- 检查 `npu-smi info` 是否可完成。
+- 默认拒绝 busy/unhealthy 设备。
+- 分别启动 rank0/rank1。
+- 为 prepare、install、submit、p2p result 做日志断言。
+- 用 `timeout` 包住 rank 进程，避免测试卡死。
+
+如确认设备健康但 `npu-smi` health 字段不是 OK，可显式放开 unhealthy 检查：
+
+```bash
+export TILEXR_CCU_SMOKE_ALLOW_UNHEALTHY_NPU=1
+```
+
+不要在没有明确授权时设置：
+
+```bash
+export TILEXR_CCU_SMOKE_ALLOW_BUSY_NPU=1
+```
+
+### 10.4 前序验证记录
+
+在 950 验证服务器的前序复测中，当前分支曾通过以下检查：
+
+- `cmake --build build_ccu_direct --target tile-comm -j4`
+- 143 个 CCU 单元测试通过，1 个 skip。
+- 两卡 P2P copy smoke 在 `remote_to_local` 和 `local_to_remote` 两个方向通过。
+- `TILEXR_ENABLE_CCU_BACKEND=1` 且 `TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT=0` 的自动初始化路径通过 P2P copy smoke。
+
+注意：默认非 P2P smoke 曾暴露 repository install 旧参数问题，因此不能把默认非 P2P smoke 作为当前推荐验证路径。
+
+## 11. 常见问题定位
+
+### 11.1 runtime 初始化失败
+
+看日志中的：
+
+```text
+TileXR direct CCU runtime init failed
+logicDevId
+devicePhyId
+hdcType
+raInitialized
+ccuTlvInitialized
+message
+```
+
+如果同一 device 后续被跳过，说明 device-scoped unavailable 状态已经记录。需要先解决首次失败原因，或换健康 device 重试。
+
+### 11.2 prepare 卡住或 allgather 超时
+
+重点检查：
+
+- rank0/rank1 是否都启动。
+- `TILEXR_COMM_ID`、端口和 rankSize 是否一致。
+- 某个 rank 是否先因 runtime unavailable 退出。
+- thread mode 下 uid 是否一致。
+- `TILEXR_CCU_DIRECT_SMOKE_READY_TIMEOUT_MS` 是否过短。
+
+### 11.3 resource window 注册失败
+
+重点检查：
+
+- 是否设置 `TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE=ra_ctx`。
+- EID index 是否在 `RaGetDevEidInfoList` 返回列表内。
+- 是否选择了错误 device 或 busy device。
+- resource address 是否来自当前 device 的 CCU basic info。
+
+可打开：
+
+```bash
+export TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE=1
+```
+
+查看 EID、TPN、doorbell token、QP import 等信息。
+
+### 11.4 repository 或 lower-layer install 失败
+
+重点看 prepare report：
+
+- `installAttempted`
+- `installSucceeded`
+- `requiredInstallSurfaceCount`
+- `publicVerifiedInstallSurfaceCount`
+- `missingInstallSurfaceCount`
+- `message`
+
+如果缺失 surface，通常说明 lower-layer plan 没有拿到完整 peer route/token，或 resource range 与 manifest 不匹配。
+
+### 11.5 submit 失败
+
+submit 失败时 report 会打印 task 关键字段：
+
+```text
+missionId
+key
+instStartId
+instCnt
+argSize
+args[]
+rtRet
+```
+
+定位顺序：
+
+1. 确认 prepare 阶段 `submitReady=1`。
+2. 确认 repository mission window 覆盖 task 的 `instStartId/instCnt`。
+3. 确认 mission id/key 与 install manifest 匹配。
+4. 确认 stream 非空。
+5. 打开 `TILEXR_CCU_DIRECT_TRACE=1` 查看最终 task 和 microcode 解码。
+
+### 11.6 P2P copy 不匹配
+
+重点区分方向：
+
+- `remote_to_local`：active rank 校验本地 destination。
+- `local_to_remote`：inactive rank 等待 active rank done 后校验本地 destination。
+
+常见原因：
+
+- active rank 设置错误。
+- direction 与期望校验 rank 不一致。
+- process token 查询失败。
+- endpoint route 不完整导致 channel 指向错误 peer。
+- submit 成功但 stream synchronize 失败。
+
+## 12. 后续扩展建议
+
+当前代码已经具备 direct runtime、resource allocator、lower-layer install、repository install、mission/task submit 和 P2P copy 数据面验证基础。后续要做泛化 alltoall/collective，建议按以下顺序推进：
+
+1. 先把 collective request 到 producer plan 的映射补齐，而不是绕过 `TileXRCcuCollectivePlanner::Supports()`。
+2. 复用现有 resource allocator 和 lower-layer evidence，不新增无证据的 env override。
+3. 用多 task submit 的单元测试覆盖 launch package 批量下发。
+4. 在两卡 P2P copy 稳定后，再扩展到 4 卡 N-to-N 数据流。
+5. alltoall 数据正确性必须以 device buffer 内容校验为准，不能用 host marker 或单纯 submit 成功替代。
diff --git a/src/comm/CMakeLists.txt b/src/comm/CMakeLists.txt
index e0e6c35d..52aa2d98 100644
--- a/src/comm/CMakeLists.txt
+++ b/src/comm/CMakeLists.txt
@@ -145,6 +145,8 @@ set(TILEXR_SOURCE_FILE tilexr_comm.cpp
         ccu/tilexr_ccu_resource_allocator.cpp
         ccu/tilexr_ccu_runtime.h
         ccu/tilexr_ccu_runtime.cpp
+        ccu/tilexr_ccu_signal_wait_program.h
+        ccu/tilexr_ccu_signal_wait_program.cpp
         ccu/tilexr_ccu_runtime_session.h
         ccu/tilexr_ccu_runtime_session.cpp
         ccu/tilexr_ccu_executor.h
diff --git a/src/comm/ccu/tilexr_ccu_backend.cpp b/src/comm/ccu/tilexr_ccu_backend.cpp
index 3c3afbde..27efa405 100644
--- a/src/comm/ccu/tilexr_ccu_backend.cpp
+++ b/src/comm/ccu/tilexr_ccu_backend.cpp
@@ -20,6 +20,11 @@ class TileXRCcuBackend::Impl {
     bool Supports(const TileXRCcuCollectiveRequest &request) const;
     int PrepareCollective(const TileXRCcuCollectiveRequest &request, TileXRCcuCollectivePlan *plan);
     int SubmitCollective(const TileXRCcuCollectivePlan &plan, aclrtStream stream);
+    int PrepareSignalWait(const TileXRCcuSignalWaitRequest &request, TileXRCcuSignalWaitPlan *plan);
+    int SubmitSignalWait(
+        const TileXRCcuSignalWaitPlan &plan,
+        aclrtStream stream,
+        TileXRCcuDirectSubmitReport *report);
 #ifdef TILEXR_CCU_TESTING
     bool RuntimeInitializedForTest() const;
 #endif
@@ -105,6 +110,38 @@ int TileXRCcuBackend::Impl::SubmitCollective(const TileXRCcuCollectivePlan &plan
     return executor_->Submit(*runtimeSession_, plan, stream);
 }
 
+int TileXRCcuBackend::Impl::PrepareSignalWait(
+    const TileXRCcuSignalWaitRequest &request,
+    TileXRCcuSignalWaitPlan *plan)
+{
+    if (plan == nullptr) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    *plan = TileXRCcuSignalWaitPlan {};
+    if (runtimeSession_ == nullptr || planner_ == nullptr) {
+        return TILEXR_ERROR_INTERNAL;
+    }
+    TileXRCcuDirectInstallReport report;
+    return planner_->PrepareSignalWait(*runtimeSession_, request, plan, &report);
+}
+
+int TileXRCcuBackend::Impl::SubmitSignalWait(
+    const TileXRCcuSignalWaitPlan &plan,
+    aclrtStream stream,
+    TileXRCcuDirectSubmitReport *report)
+{
+    if (report != nullptr) {
+        *report = TileXRCcuDirectSubmitReport {};
+    }
+    if (!plan.ready || plan.submitTasks.empty()) {
+        if (report != nullptr) {
+            report->message = "direct CCU signal/wait plan is not ready";
+        }
+        return TILEXR_ERROR_NOT_INITIALIZED;
+    }
+    return TileXRCcuSubmitPreparedTasks(plan.submitTasks, stream, nullptr, nullptr, report);
+}
+
 #ifdef TILEXR_CCU_TESTING
 bool TileXRCcuBackend::Impl::RuntimeInitializedForTest() const
 {
@@ -165,6 +202,27 @@ int TileXRCcuBackend::SubmitCollective(const TileXRCcuCollectivePlan &plan, aclr
     return impl_->SubmitCollective(plan, stream);
 }
 
+int TileXRCcuBackend::PrepareSignalWait(
+    const TileXRCcuSignalWaitRequest &request,
+    TileXRCcuSignalWaitPlan *plan)
+{
+    if (impl_ == nullptr) {
+        return TILEXR_ERROR_INTERNAL;
+    }
+    return impl_->PrepareSignalWait(request, plan);
+}
+
+int TileXRCcuBackend::SubmitSignalWait(
+    const TileXRCcuSignalWaitPlan &plan,
+    aclrtStream stream,
+    TileXRCcuDirectSubmitReport *report)
+{
+    if (impl_ == nullptr) {
+        return TILEXR_ERROR_INTERNAL;
+    }
+    return impl_->SubmitSignalWait(plan, stream, report);
+}
+
 #ifdef TILEXR_CCU_TESTING
 bool TileXRCcuBackend::RuntimeInitializedForTest() const
 {
diff --git a/src/comm/ccu/tilexr_ccu_backend.h b/src/comm/ccu/tilexr_ccu_backend.h
index 8177e211..65ebfd82 100644
--- a/src/comm/ccu/tilexr_ccu_backend.h
+++ b/src/comm/ccu/tilexr_ccu_backend.h
@@ -8,7 +8,9 @@
 #include <cstdint>
 #include <memory>
 #include <string>
+#include <vector>
 
+#include "ccu/tilexr_ccu_direct_orchestrator.h"
 #include "acl/acl_base.h"
 #include "tilexr_types.h"
 
@@ -42,6 +44,38 @@ struct TileXRCcuCollectivePlan {
     bool ready = false;
 };
 
+enum class TileXRCcuSignalWaitRole {
+    Signal = 0,
+    Wait = 1,
+    SignalAndWait = 2,
+};
+
+struct TileXRCcuSignalWaitRequest {
+    int peerRank = -1;
+    TileXRCcuSignalWaitRole role = TileXRCcuSignalWaitRole::Signal;
+    uint32_t syncInstructionCount = 0;
+    uint16_t missionStartId = 0;
+    uint16_t instructionStartId = 0;
+    uint16_t missionInstructionStartId = 0;
+    uint16_t xnStartId = 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;
+    uint16_t timeout = 0;
+    std::string provider;
+};
+
+struct TileXRCcuSignalWaitPlan {
+    bool ready = false;
+    TileXRCcuDirectInstallAttempt attempt;
+    std::vector<TileXRCcuTask> submitTasks;
+};
+
 class TileXRCcuBackend {
 public:
     TileXRCcuBackend();
@@ -56,6 +90,11 @@ class TileXRCcuBackend {
     bool Supports(const TileXRCcuCollectiveRequest &request) const;
     int PrepareCollective(const TileXRCcuCollectiveRequest &request, TileXRCcuCollectivePlan *plan);
     int SubmitCollective(const TileXRCcuCollectivePlan &plan, aclrtStream stream);
+    int PrepareSignalWait(const TileXRCcuSignalWaitRequest &request, TileXRCcuSignalWaitPlan *plan);
+    int SubmitSignalWait(
+        const TileXRCcuSignalWaitPlan &plan,
+        aclrtStream stream,
+        TileXRCcuDirectSubmitReport *report);
 #ifdef TILEXR_CCU_TESTING
     bool RuntimeInitializedForTest() const;
 #endif
diff --git a/src/comm/ccu/tilexr_ccu_barrier_program.cpp b/src/comm/ccu/tilexr_ccu_barrier_program.cpp
index 319caeea..73b3a783 100644
--- a/src/comm/ccu/tilexr_ccu_barrier_program.cpp
+++ b/src/comm/ccu/tilexr_ccu_barrier_program.cpp
@@ -221,7 +221,7 @@ int TileXRCcuBuildBarrierProgram(
             TileXRCcuSyncCkeSpec post;
             post.remoteCke = spec.remoteNotifyCke;
             post.localCke = spec.sourceCke;
-            post.localCkeMask = spec.remoteNotifyMask;
+            post.localCkeMask = spec.sourceCkeMask;
             post.channelId = spec.channelId;
             if (TileXRCcuEncodeSyncCke(post, &instr) != TILEXR_SUCCESS) {
                 return Fail(program, report, "failed to encode CCU barrier SyncCKE post instruction");
diff --git a/src/comm/ccu/tilexr_ccu_collective_planner.cpp b/src/comm/ccu/tilexr_ccu_collective_planner.cpp
index 311eec0d..7991522b 100644
--- a/src/comm/ccu/tilexr_ccu_collective_planner.cpp
+++ b/src/comm/ccu/tilexr_ccu_collective_planner.cpp
@@ -13,6 +13,7 @@
 #include <algorithm>
 #include <cstdlib>
 #include <cstring>
+#include <iostream>
 
 #ifdef TILEXR_CCU_TESTING
 #include "runtime/dev.h"
@@ -25,6 +26,9 @@ namespace TileXR {
 #ifdef TILEXR_CCU_TESTING
 constexpr uint32_t TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT = 7U;
 #endif
+constexpr uint32_t TILEXR_CCU_DIRECT_SIGNAL_INSTRUCTION_COUNT = 5U;
+constexpr uint32_t TILEXR_CCU_DIRECT_WAIT_INSTRUCTION_COUNT = 5U;
+constexpr uint32_t TILEXR_CCU_DIRECT_SIGNAL_WAIT_INSTRUCTION_COUNT = 6U;
 
 uint8_t SelectDirectCcuInstallDieId()
 {
@@ -73,17 +77,98 @@ uint16_t SelectDirectCcuChannelBoundRemoteXnId(
         SelectDirectCcuChannelBoundRemoteXnOffset(peerLocalIndex, syncIndex, peerRouteCount));
 }
 
+TileXRCcuSignalWaitProgramRole ToDirectSignalWaitProgramRole(TileXRCcuSignalWaitRole role)
+{
+    if (role == TileXRCcuSignalWaitRole::Wait) {
+        return TileXRCcuSignalWaitProgramRole::Wait;
+    }
+    if (role == TileXRCcuSignalWaitRole::SignalAndWait) {
+        return TileXRCcuSignalWaitProgramRole::SignalAndWait;
+    }
+    return TileXRCcuSignalWaitProgramRole::Signal;
+}
+
+uint32_t DefaultSignalWaitInstructionCount(TileXRCcuSignalWaitRole role)
+{
+    if (role == TileXRCcuSignalWaitRole::Wait) {
+        return TILEXR_CCU_DIRECT_WAIT_INSTRUCTION_COUNT;
+    }
+    if (role == TileXRCcuSignalWaitRole::SignalAndWait) {
+        return TILEXR_CCU_DIRECT_SIGNAL_WAIT_INSTRUCTION_COUNT;
+    }
+    return TILEXR_CCU_DIRECT_SIGNAL_INSTRUCTION_COUNT;
+}
+
 #ifdef TILEXR_CCU_TESTING
 struct DirectCcuMemoryCopyEndpoint {
     uint64_t sourceAddr = 0;
     uint64_t sourceToken = 0;
+    TileXRCcuRemoteMemoryBufferImportRequest sourceRemoteImport {};
     uint64_t destinationAddr = 0;
     uint64_t destinationToken = 0;
+    TileXRCcuRemoteMemoryBufferImportRequest destinationRemoteImport {};
     uint64_t bytes = 0;
     uint32_t rank = 0;
     uint32_t valid = 0;
 };
 
+bool TraceDirectCcuMemoryCopy()
+{
+    const char *trace = std::getenv("TILEXR_CCU_DIRECT_TRACE");
+    return trace != nullptr && trace[0] != '\0' && trace[0] != '0';
+}
+
+bool UseCcuResourceWindowForMemoryCopy()
+{
+    const char *value = std::getenv("TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_RESOURCE_WINDOW");
+    return value != nullptr && value[0] != '\0' && value[0] != '0';
+}
+
+void TraceDirectCcuProcessMemoryToken(
+    uint64_t addr,
+    uint64_t bytes,
+    uint32_t rawTokenId,
+    uint32_t tokenId,
+    uint32_t tokenValue,
+    uint64_t packedToken)
+{
+    if (!TraceDirectCcuMemoryCopy()) {
+        return;
+    }
+    std::cerr << "TileXRDirectCcuTrace memoryToken"
+              << " addr=0x" << std::hex << addr
+              << " bytes=0x" << bytes
+              << " rawTokenId=0x" << rawTokenId
+              << " tokenId=0x" << tokenId
+              << " tokenValue=0x" << tokenValue
+              << " packed=0x" << packedToken
+              << std::dec << std::endl;
+}
+
+void TraceDirectCcuRegisteredMemoryToken(
+    const char *label,
+    uint64_t addr,
+    uint64_t bytes,
+    const TileXRCcuRegisteredMemoryBufferInfo &info,
+    uint64_t packedToken)
+{
+    if (!TraceDirectCcuMemoryCopy()) {
+        return;
+    }
+    std::cerr << "TileXRDirectCcuTrace registeredMemoryToken"
+              << " label=" << (label == nullptr ? "" : label)
+              << " addr=0x" << std::hex << addr
+              << " bytes=0x" << bytes
+              << " alignedAddr=0x" << info.alignedAddr
+              << " alignedBytes=0x" << info.alignedBytes
+              << " targetSegVa=0x" << info.targetSegVa
+              << " rawTokenId=0x" << info.rawTokenId
+              << " tokenId=0x" << info.tokenId
+              << " tokenValue=0x" << info.tokenValue
+              << " packed=0x" << packedToken
+              << std::dec << std::endl;
+}
+
 int QueryDirectCcuProcessMemoryToken(uint64_t addr, uint64_t bytes, uint64_t *packedToken)
 {
     if (addr == 0 || bytes == 0 || packedToken == nullptr) {
@@ -97,13 +182,14 @@ int QueryDirectCcuProcessMemoryToken(uint64_t addr, uint64_t bytes, uint64_t *pa
     if (ret != RT_ERROR_NONE) {
         return TILEXR_ERROR_MKIRT;
     }
-    constexpr uint32_t tokenIdRightShift = 8U;
-    const uint32_t tokenId = info.tokenId >> tokenIdRightShift;
+    const uint32_t tokenId = info.tokenId;
     *packedToken = TileXRCcuPackMemoryToken(tokenId, info.tokenValue, true);
+    TraceDirectCcuProcessMemoryToken(addr, bytes, info.tokenId, tokenId, info.tokenValue, *packedToken);
     return *packedToken == 0 ? TILEXR_ERROR_NOT_FOUND : TILEXR_SUCCESS;
 }
 
 int BuildDirectCcuLocalMemoryCopyEndpoint(
+    TileXRCcuRuntimeSession &session,
     uint32_t rank,
     uint64_t sourceAddr,
     uint64_t destinationAddr,
@@ -116,16 +202,53 @@ int BuildDirectCcuLocalMemoryCopyEndpoint(
     *endpoint = DirectCcuMemoryCopyEndpoint {};
     endpoint->rank = rank;
     endpoint->bytes = bytes;
-    endpoint->sourceAddr = sourceAddr;
-    endpoint->destinationAddr = destinationAddr;
-    int ret = QueryDirectCcuProcessMemoryToken(sourceAddr, bytes, &endpoint->sourceToken);
+    TileXRCcuRegisteredMemoryBufferInfo sourceInfo;
+    int ret = session.RegisterMemoryBuffer(sourceAddr, bytes, &sourceInfo);
     if (ret != TILEXR_SUCCESS) {
         return ret;
     }
-    ret = QueryDirectCcuProcessMemoryToken(destinationAddr, bytes, &endpoint->destinationToken);
+    if (sourceInfo.targetSegVa == 0) {
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+    endpoint->sourceAddr = sourceInfo.addr;
+    endpoint->sourceToken = TileXRCcuPackMemoryToken(sourceInfo.tokenId, sourceInfo.tokenValue, true);
+    endpoint->sourceRemoteImport.addr = sourceInfo.addr;
+    endpoint->sourceRemoteImport.bytes = sourceInfo.bytes;
+    endpoint->sourceRemoteImport.alignedAddr = sourceInfo.alignedAddr;
+    endpoint->sourceRemoteImport.offset = sourceInfo.addr - sourceInfo.alignedAddr;
+    endpoint->sourceRemoteImport.tokenId = sourceInfo.tokenId;
+    endpoint->sourceRemoteImport.rawTokenId = sourceInfo.rawTokenId;
+    endpoint->sourceRemoteImport.tokenValue = sourceInfo.tokenValue;
+    endpoint->sourceRemoteImport.key = sourceInfo.key;
+    endpoint->sourceRemoteImport.valid = true;
+    TraceDirectCcuRegisteredMemoryToken("source", sourceAddr, bytes, sourceInfo, endpoint->sourceToken);
+
+    TileXRCcuRegisteredMemoryBufferInfo destinationInfo;
+    ret = session.RegisterMemoryBuffer(destinationAddr, bytes, &destinationInfo);
     if (ret != TILEXR_SUCCESS) {
         return ret;
     }
+    if (destinationInfo.targetSegVa == 0) {
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+    endpoint->destinationAddr = destinationInfo.addr;
+    endpoint->destinationToken =
+        TileXRCcuPackMemoryToken(destinationInfo.tokenId, destinationInfo.tokenValue, true);
+    endpoint->destinationRemoteImport.addr = destinationInfo.addr;
+    endpoint->destinationRemoteImport.bytes = destinationInfo.bytes;
+    endpoint->destinationRemoteImport.alignedAddr = destinationInfo.alignedAddr;
+    endpoint->destinationRemoteImport.offset = destinationInfo.addr - destinationInfo.alignedAddr;
+    endpoint->destinationRemoteImport.tokenId = destinationInfo.tokenId;
+    endpoint->destinationRemoteImport.rawTokenId = destinationInfo.rawTokenId;
+    endpoint->destinationRemoteImport.tokenValue = destinationInfo.tokenValue;
+    endpoint->destinationRemoteImport.key = destinationInfo.key;
+    endpoint->destinationRemoteImport.valid = true;
+    TraceDirectCcuRegisteredMemoryToken(
+        "destination",
+        destinationAddr,
+        bytes,
+        destinationInfo,
+        endpoint->destinationToken);
     endpoint->valid = 1;
     return TILEXR_SUCCESS;
 }
@@ -332,6 +455,9 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuLowerLayerTemplateFromAllocation
         directCcuLowerLayerPlanStatus_ = ret;
         return directCcuLowerLayerPlanStatus_;
     }
+#ifdef TILEXR_CCU_TESTING
+    ApplyDirectCcuRemoteRouteMemoryOverride(&remoteCcuBuffers);
+#endif
 
     TileXRCcuLowerLayerTransportSnapshot templateSnapshot;
     TileXRCcuLowerLayerPlanBuilderReport report;
@@ -539,6 +665,48 @@ int TileXRCcuCollectivePlanner::ExchangeDirectCcuRemoteNotifyCke(
     return TILEXR_SUCCESS;
 }
 
+#ifdef TILEXR_CCU_TESTING
+void TileXRCcuCollectivePlanner::SetDirectCcuRemoteRouteMemoryOverride(
+    uint32_t peerRank,
+    uint64_t remoteCcuVa,
+    uint32_t memoryTokenId,
+    uint32_t rawMemoryTokenId,
+    uint32_t memoryTokenValue)
+{
+    directCcuRemoteRouteMemoryOverride_ = TileXRCcuRemoteCcuBufferInfo {};
+    directCcuRemoteRouteMemoryOverride_.peerRank = peerRank;
+    directCcuRemoteRouteMemoryOverride_.remoteCcuVa = remoteCcuVa;
+    directCcuRemoteRouteMemoryOverride_.memoryTokenId = memoryTokenId;
+    directCcuRemoteRouteMemoryOverride_.rawMemoryTokenId = rawMemoryTokenId;
+    directCcuRemoteRouteMemoryOverride_.memoryTokenValue = memoryTokenValue;
+    directCcuRemoteRouteMemoryOverrideValid_ =
+        remoteCcuVa != 0 && memoryTokenId != 0 && memoryTokenValue != 0;
+}
+
+void TileXRCcuCollectivePlanner::ClearDirectCcuRemoteRouteMemoryOverride()
+{
+    directCcuRemoteRouteMemoryOverride_ = TileXRCcuRemoteCcuBufferInfo {};
+    directCcuRemoteRouteMemoryOverrideValid_ = false;
+}
+
+void TileXRCcuCollectivePlanner::ApplyDirectCcuRemoteRouteMemoryOverride(
+    std::vector<TileXRCcuRemoteCcuBufferInfo> *remoteCcuBuffers) const
+{
+    if (!directCcuRemoteRouteMemoryOverrideValid_ || remoteCcuBuffers == nullptr) {
+        return;
+    }
+    for (auto &remoteCcuBuffer : *remoteCcuBuffers) {
+        if (remoteCcuBuffer.peerRank != directCcuRemoteRouteMemoryOverride_.peerRank) {
+            continue;
+        }
+        remoteCcuBuffer.remoteCcuVa = directCcuRemoteRouteMemoryOverride_.remoteCcuVa;
+        remoteCcuBuffer.memoryTokenId = directCcuRemoteRouteMemoryOverride_.memoryTokenId;
+        remoteCcuBuffer.rawMemoryTokenId = directCcuRemoteRouteMemoryOverride_.rawMemoryTokenId;
+        remoteCcuBuffer.memoryTokenValue = directCcuRemoteRouteMemoryOverride_.memoryTokenValue;
+    }
+}
+#endif
+
 int TileXRCcuCollectivePlanner::PrepareDirectCcuLowerLayerPlanCallback(
     const TileXRCcuResourceAllocation &allocation,
     TileXRCcuLowerLayerInstallPlan *plan,
@@ -569,12 +737,11 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuInstallAttempt(
         }
         return TILEXR_ERROR_NOT_INITIALIZED;
     }
-    const std::string processUnavailableMessage =
-        TileXRCcuRuntimeSession::ProcessDirectCcuRuntimeUnavailableMessage();
-    if (!processUnavailableMessage.empty()) {
+    const std::string unavailableMessage = session.DirectCcuRuntimeUnavailableMessage();
+    if (!unavailableMessage.empty()) {
         if (report != nullptr) {
             *report = TileXRCcuDirectInstallReport{};
-            report->message = processUnavailableMessage;
+            report->message = unavailableMessage;
         }
         return TILEXR_ERROR_NOT_FOUND;
     }
@@ -627,6 +794,121 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuInstallAttempt(
     return TileXRCcuRunDirectInstallAttempt(next, attempt, report);
 }
 
+int TileXRCcuCollectivePlanner::PrepareSignalWait(
+    TileXRCcuRuntimeSession &session,
+    const TileXRCcuSignalWaitRequest &request,
+    TileXRCcuSignalWaitPlan *plan,
+    TileXRCcuDirectInstallReport *report)
+{
+    if (plan == nullptr) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = "missing direct CCU signal/wait plan output";
+        }
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    *plan = TileXRCcuSignalWaitPlan {};
+    if (!session.Available()) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = "TileXRCcuBackend is not initialized for direct CCU signal/wait";
+        }
+        return TILEXR_ERROR_NOT_INITIALIZED;
+    }
+    const int rank = session.Rank();
+    const int rankSize = session.RankSize();
+    if (rankSize != 2 || request.peerRank < 0 || request.peerRank >= rankSize || request.peerRank == rank) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = "direct CCU signal/wait requires two ranks and a valid peer rank";
+        }
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    const std::string unavailableMessage = session.DirectCcuRuntimeUnavailableMessage();
+    if (!unavailableMessage.empty()) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = unavailableMessage;
+        }
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+
+    const uint8_t installDieId = SelectDirectCcuInstallDieId();
+    const TileXRCcuBasicInfo *basicInfo = session.GetDirectCcuBasicInfo();
+    if (basicInfo == nullptr || basicInfo->dieId != installDieId) {
+        const int ret = session.RefreshDirectCcuBasicInfo(installDieId);
+        if (ret != TILEXR_SUCCESS) {
+            if (report != nullptr) {
+                *report = TileXRCcuDirectInstallReport {};
+                report->message = session.GetDirectCcuBasicInfoReport().message;
+            }
+            return ret;
+        }
+        basicInfo = session.GetDirectCcuBasicInfo();
+    }
+    if (basicInfo == nullptr || !session.Available()) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = "direct CCU runtime is unavailable for signal/wait";
+        }
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+
+    TileXRCcuDriverAdapter adapter;
+    TileXRCcuDriverAdapterReport adapterReport;
+    int ret = session.CreateDriverAdapter(&adapter, &adapterReport);
+    if (ret != TILEXR_SUCCESS) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = adapterReport.message;
+        }
+        return ret;
+    }
+
+    LowerLayerPlanCallbackContext callbackContext {this, &session};
+    TileXRCcuDirectInstallOptions options;
+    options.basicInfo = basicInfo;
+    options.sqeArgCount = 0;
+    options.syncResourceCount = 1;
+    options.syncInstructionCount = request.syncInstructionCount == 0 ?
+        DefaultSignalWaitInstructionCount(request.role) :
+        request.syncInstructionCount;
+    options.bindingsPerSyncResource = 1;
+    options.missionStartId = request.missionStartId;
+    options.instructionStartId = request.instructionStartId;
+    options.missionInstructionStartId = request.missionInstructionStartId;
+    options.xnStartId = request.xnStartId;
+    options.remoteXnStartId = request.remoteXnStartId;
+    options.remoteXnCount = request.remoteXnCount;
+    options.ckeStartId = request.ckeStartId;
+    options.channelStartId = request.channelStartId;
+    options.localWaitCkeStartId = request.localWaitCkeStartId;
+    options.localWaitCkeCount = request.localWaitCkeCount;
+    options.remoteNotifyCkeStartId = request.remoteNotifyCkeStartId;
+    options.remoteNotifyCkeCount = request.remoteNotifyCkeCount;
+    options.taskTimeout = request.timeout;
+    options.deviceId = static_cast<uint32_t>(session.DevId());
+    options.rank = static_cast<uint32_t>(rank);
+    options.provider = request.provider.empty() ? "tilexr-comm-direct-ccu-signal-wait" : request.provider;
+    options.offlineOnly = false;
+    options.driverAdapter = &adapter;
+    options.repositoryMemoryOps = TileXRCcuMakeRepositoryDeviceMemoryOps(options.repositoryMemoryAllocMode);
+    options.repositoryMemoryUserData = nullptr;
+    options.prepareLowerLayerPlan = &TileXRCcuCollectivePlanner::PrepareDirectCcuLowerLayerPlanCallback;
+    options.lowerLayerPlanUserData = &callbackContext;
+
+    TileXRCcuDirectSignalWaitSpec signalWait;
+    signalWait.role = ToDirectSignalWaitProgramRole(request.role);
+    ret = TileXRCcuRunDirectSignalWaitInstallAttempt(options, signalWait, &plan->attempt, report);
+    if (ret != TILEXR_SUCCESS) {
+        *plan = TileXRCcuSignalWaitPlan {};
+        return ret;
+    }
+    plan->submitTasks = plan->attempt.submitTasks;
+    plan->ready = !plan->submitTasks.empty();
+    return plan->ready ? TILEXR_SUCCESS : TILEXR_ERROR_NOT_FOUND;
+}
+
 #ifdef TILEXR_CCU_TESTING
 int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt(
     TileXRCcuRuntimeSession &session,
@@ -656,12 +938,11 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt(
         }
         return TILEXR_ERROR_PARA_CHECK_FAIL;
     }
-    const std::string processUnavailableMessage =
-        TileXRCcuRuntimeSession::ProcessDirectCcuRuntimeUnavailableMessage();
-    if (!processUnavailableMessage.empty()) {
+    const std::string unavailableMessage = session.DirectCcuRuntimeUnavailableMessage();
+    if (!unavailableMessage.empty()) {
         if (report != nullptr) {
             *report = TileXRCcuDirectInstallReport {};
-            report->message = processUnavailableMessage;
+            report->message = unavailableMessage;
         }
         return TILEXR_ERROR_NOT_FOUND;
     }
@@ -686,8 +967,18 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt(
         return TILEXR_ERROR_NOT_FOUND;
     }
 
+    int ret = session.RegisterCcuResourceRmaBuffer(basicInfo->resourceAddr);
+    if (ret != TILEXR_SUCCESS) {
+        if (report != nullptr) {
+            *report = TileXRCcuDirectInstallReport {};
+            report->message = "failed to register direct CCU resource window before memory copy buffer registration";
+        }
+        return ret;
+    }
+
     DirectCcuMemoryCopyEndpoint localEndpoint;
-    int ret = BuildDirectCcuLocalMemoryCopyEndpoint(
+    ret = BuildDirectCcuLocalMemoryCopyEndpoint(
+        session,
         static_cast<uint32_t>(rank),
         localSourceAddr,
         localDestinationAddr,
@@ -722,19 +1013,91 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt(
         return TILEXR_ERROR_PARA_CHECK_FAIL;
     }
 
+    TileXRCcuLocalResourceWindowInfo localResourceWindow;
+    std::vector<TileXRCcuLocalResourceWindowInfo> allResourceWindows;
+    const bool useResourceWindow = UseCcuResourceWindowForMemoryCopy();
+    if (useResourceWindow) {
+        ret = session.ExportLocalCcuRmaBuffer(&localResourceWindow);
+        if (ret != TILEXR_SUCCESS || localResourceWindow.addr == 0 || localResourceWindow.tokenId == 0) {
+            if (report != nullptr) {
+                *report = TileXRCcuDirectInstallReport {};
+                report->message = "failed to export direct CCU resource window for memory copy";
+            }
+            return ret == TILEXR_SUCCESS ? TILEXR_ERROR_NOT_FOUND : ret;
+        }
+        allResourceWindows.resize(static_cast<size_t>(rankSize));
+        ret = session.AllGather(
+            &localResourceWindow,
+            sizeof(localResourceWindow),
+            allResourceWindows.data());
+        if (ret != TILEXR_SUCCESS) {
+            if (report != nullptr) {
+                *report = TileXRCcuDirectInstallReport {};
+                report->message = "failed to exchange direct CCU resource windows for memory copy";
+            }
+            return ret;
+        }
+    }
+
     TileXRCcuDirectMemoryCopySpec memoryCopy;
     memoryCopy.direction = direction;
     memoryCopy.lengthBytes = bytes;
+    TileXRCcuImportedRemoteMemoryBufferInfo importedRemoteBuffer;
+    TileXRCcuRemoteMemoryBufferImportRequest remoteImportRequest;
     if (direction == TileXRCcuMemoryCopyDirection::RemoteToLocal) {
+        remoteImportRequest = peerEndpoint.sourceRemoteImport;
+        ret = session.ImportRemoteMemoryBuffer(remoteImportRequest, &importedRemoteBuffer);
+        if (ret != TILEXR_SUCCESS) {
+            if (report != nullptr) {
+                *report = TileXRCcuDirectInstallReport {};
+                report->message = "failed to import direct CCU memory copy remote source buffer";
+            }
+            return ret;
+        }
         memoryCopy.localAddr = localEndpoint.destinationAddr;
         memoryCopy.localToken = localEndpoint.destinationToken;
-        memoryCopy.remoteAddr = peerEndpoint.sourceAddr;
-        memoryCopy.remoteToken = peerEndpoint.sourceToken;
+        if (useResourceWindow) {
+            const auto &peerWindow = allResourceWindows[peerRank];
+            memoryCopy.remoteAddr = peerWindow.addr;
+            memoryCopy.remoteToken = TileXRCcuPackMemoryToken(peerWindow.tokenId, peerWindow.tokenValue, true);
+        } else {
+            memoryCopy.remoteAddr = remoteImportRequest.addr;
+            memoryCopy.remoteToken =
+                TileXRCcuPackMemoryToken(remoteImportRequest.tokenId, remoteImportRequest.tokenValue, true);
+        }
     } else {
+        remoteImportRequest = peerEndpoint.destinationRemoteImport;
+        ret = session.ImportRemoteMemoryBuffer(remoteImportRequest, &importedRemoteBuffer);
+        if (ret != TILEXR_SUCCESS) {
+            if (report != nullptr) {
+                *report = TileXRCcuDirectInstallReport {};
+                report->message = "failed to import direct CCU memory copy remote destination buffer";
+            }
+            return ret;
+        }
         memoryCopy.localAddr = localEndpoint.sourceAddr;
         memoryCopy.localToken = localEndpoint.sourceToken;
-        memoryCopy.remoteAddr = peerEndpoint.destinationAddr;
-        memoryCopy.remoteToken = peerEndpoint.destinationToken;
+        if (useResourceWindow) {
+            const auto &peerWindow = allResourceWindows[peerRank];
+            memoryCopy.remoteAddr = peerWindow.addr;
+            memoryCopy.remoteToken = TileXRCcuPackMemoryToken(peerWindow.tokenId, peerWindow.tokenValue, true);
+        } else {
+            memoryCopy.remoteAddr = remoteImportRequest.addr;
+            memoryCopy.remoteToken =
+                TileXRCcuPackMemoryToken(remoteImportRequest.tokenId, remoteImportRequest.tokenValue, true);
+        }
+    }
+    if (TraceDirectCcuMemoryCopy()) {
+        std::cerr << "TileXRDirectCcuTrace memoryCopyEndpoint"
+                  << " rank=" << rank
+                  << " peerRank=" << peerRank
+                  << " direction=" << static_cast<int>(direction)
+                  << " localAddr=0x" << std::hex << memoryCopy.localAddr
+                  << " localToken=0x" << memoryCopy.localToken
+                  << " remoteAddr=0x" << memoryCopy.remoteAddr
+                  << " remoteToken=0x" << memoryCopy.remoteToken
+                  << " length=0x" << memoryCopy.lengthBytes
+                  << std::dec << std::endl;
     }
 
     TileXRCcuDriverAdapter adapter;
@@ -768,7 +1131,17 @@ int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt(
         next.provider = "tilexr-comm-direct-ccu-memory-copy";
     }
 
-    return TileXRCcuRunDirectMemoryCopyInstallAttempt(next, memoryCopy, attempt, report);
+    if (!useResourceWindow) {
+        SetDirectCcuRemoteRouteMemoryOverride(
+            peerRank,
+            importedRemoteBuffer.targetSegVa,
+            remoteImportRequest.tokenId,
+            remoteImportRequest.rawTokenId,
+            remoteImportRequest.tokenValue);
+    }
+    ret = TileXRCcuRunDirectMemoryCopyInstallAttempt(next, memoryCopy, attempt, report);
+    ClearDirectCcuRemoteRouteMemoryOverride();
+    return ret;
 }
 #endif
 
diff --git a/src/comm/ccu/tilexr_ccu_collective_planner.h b/src/comm/ccu/tilexr_ccu_collective_planner.h
index 27888e5f..9c512711 100644
--- a/src/comm/ccu/tilexr_ccu_collective_planner.h
+++ b/src/comm/ccu/tilexr_ccu_collective_planner.h
@@ -47,6 +47,11 @@ class TileXRCcuCollectivePlanner {
         const TileXRCcuDirectInstallOptions &options,
         TileXRCcuDirectInstallAttempt *attempt,
         TileXRCcuDirectInstallReport *report);
+    int PrepareSignalWait(
+        TileXRCcuRuntimeSession &session,
+        const TileXRCcuSignalWaitRequest &request,
+        TileXRCcuSignalWaitPlan *plan,
+        TileXRCcuDirectInstallReport *report);
 #ifdef TILEXR_CCU_TESTING
     int PrepareDirectCcuMemoryCopyInstallAttempt(
         TileXRCcuRuntimeSession &session,
@@ -82,6 +87,17 @@ class TileXRCcuCollectivePlanner {
         const TileXRCcuResourceAllocation &allocation,
         std::vector<TileXRCcuRemoteCcuBufferInfo> *remoteCcuBuffers,
         TileXRCcuLowerLayerPlanBuilderReport *report);
+#ifdef TILEXR_CCU_TESTING
+    void SetDirectCcuRemoteRouteMemoryOverride(
+        uint32_t peerRank,
+        uint64_t remoteCcuVa,
+        uint32_t memoryTokenId,
+        uint32_t rawMemoryTokenId,
+        uint32_t memoryTokenValue);
+    void ClearDirectCcuRemoteRouteMemoryOverride();
+    void ApplyDirectCcuRemoteRouteMemoryOverride(
+        std::vector<TileXRCcuRemoteCcuBufferInfo> *remoteCcuBuffers) const;
+#endif
     static int PrepareDirectCcuLowerLayerPlanCallback(
         const TileXRCcuResourceAllocation &allocation,
         TileXRCcuLowerLayerInstallPlan *plan,
@@ -98,6 +114,10 @@ class TileXRCcuCollectivePlanner {
     std::vector<TileXRCcuLowerLayerTransportRoute> directCcuVerifiedEndpointRoutes_ = {};
     TileXRCcuLowerLayerTransportRoute directCcuLocalVerifiedEndpointRoute_ = {};
     bool directCcuLocalVerifiedEndpointRouteValid_ = false;
+#ifdef TILEXR_CCU_TESTING
+    TileXRCcuRemoteCcuBufferInfo directCcuRemoteRouteMemoryOverride_ = {};
+    bool directCcuRemoteRouteMemoryOverrideValid_ = false;
+#endif
 };
 
 } // namespace TileXR
diff --git a/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp b/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp
index f3d8aef5..3f0fc29c 100644
--- a/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp
+++ b/src/comm/ccu/tilexr_ccu_direct_orchestrator.cpp
@@ -11,6 +11,7 @@
 #include <cstdlib>
 #include <iomanip>
 #include <iostream>
+#include <limits>
 #include <sstream>
 
 namespace TileXR {
@@ -34,6 +35,9 @@ 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;
+constexpr uint32_t TILEXR_CCU_DIRECT_SIGNAL_INSTRUCTION_COUNT = 5U;
+constexpr uint32_t TILEXR_CCU_DIRECT_WAIT_INSTRUCTION_COUNT = 5U;
+constexpr uint32_t TILEXR_CCU_DIRECT_SIGNAL_WAIT_INSTRUCTION_COUNT = 6U;
 
 void ResetReport(TileXRCcuDirectInstallReport* report)
 {
@@ -784,6 +788,34 @@ int BuildDirectMemoryCopyLaunchPackage(
     copySpec.channelId = resource.channelId;
     copySpec.completionCke = resource.localWaitCke == 0 ? resource.notifyCke : resource.localWaitCke;
     copySpec.completionMask = resource.localWaitMask == 0 ? 1U : resource.localWaitMask;
+    if (DirectTraceEnabled()) {
+        const auto tracePackedToken = [](uint64_t token, const char* label) {
+            std::cerr << " " << label << "=0x" << std::hex << token
+                      << " " << label << "Valid=" << ((token >> TILEXR_CCU_PACKED_TOKEN_VALID_SHIFT) & 0x1ULL)
+                      << " " << label << "Id=0x"
+                      << ((token >> TILEXR_CCU_PACKED_TOKEN_ID_SHIFT) & TILEXR_CCU_PACKED_TOKEN_ID_MASK)
+                      << " " << label << "Value=0x"
+                      << (token & TILEXR_CCU_PACKED_TOKEN_VALUE_MASK)
+                      << std::dec;
+        };
+        std::cerr << "TileXRDirectCcuTrace memoryCopySpec"
+                  << " direction=" << static_cast<int>(copySpec.direction)
+                  << " localGsa=" << copySpec.localGsa
+                  << " localXn=" << copySpec.localXn
+                  << " remoteGsa=" << copySpec.remoteGsa
+                  << " remoteXn=" << copySpec.remoteXn
+                  << " lengthXn=" << copySpec.lengthXn
+                  << " channelId=" << copySpec.channelId
+                  << " completionCke=" << copySpec.completionCke
+                  << " completionMask=" << copySpec.completionMask
+                  << " localAddr=0x" << std::hex << copySpec.localAddr
+                  << " remoteAddr=0x" << copySpec.remoteAddr
+                  << " lengthBytes=0x" << copySpec.lengthBytes
+                  << std::dec;
+        tracePackedToken(copySpec.localToken, "localToken");
+        tracePackedToken(copySpec.remoteToken, "remoteToken");
+        std::cerr << std::endl;
+    }
 
     TileXRCcuProgram program;
     TileXRCcuMemoryProgramReport memoryReport;
@@ -821,6 +853,96 @@ int BuildDirectMemoryCopyLaunchPackage(
     return TILEXR_SUCCESS;
 }
 
+uint32_t SignalWaitInstructionCount(TileXRCcuSignalWaitProgramRole role)
+{
+    if (role == TileXRCcuSignalWaitProgramRole::Wait) {
+        return TILEXR_CCU_DIRECT_WAIT_INSTRUCTION_COUNT;
+    }
+    if (role == TileXRCcuSignalWaitProgramRole::Signal) {
+        return TILEXR_CCU_DIRECT_SIGNAL_INSTRUCTION_COUNT;
+    }
+    return TILEXR_CCU_DIRECT_SIGNAL_WAIT_INSTRUCTION_COUNT;
+}
+
+TileXRCcuBarrierMode SignalWaitBarrierMode(TileXRCcuSignalWaitProgramRole role)
+{
+    if (role == TileXRCcuSignalWaitProgramRole::Wait) {
+        return TileXRCcuBarrierMode::SyncCkePostOnly;
+    }
+    return role == TileXRCcuSignalWaitProgramRole::Signal ?
+        TileXRCcuBarrierMode::SyncCkePostOnly :
+        TileXRCcuBarrierMode::SyncCke;
+}
+
+int BuildDirectSignalWaitLaunchPackage(
+    const TileXRCcuDirectSignalWaitSpec& signalWait,
+    TileXRCcuDirectInstallAttempt* attempt,
+    TileXRCcuDirectInstallReport* report)
+{
+    if (attempt == nullptr || attempt->plan.syncResources.empty()) {
+        if (report != nullptr) {
+            report->message = "missing direct CCU signal/wait producer resources";
+        }
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    const TileXRCcuSyncResource& resource = attempt->plan.syncResources[0];
+
+    TileXRCcuSignalWaitProgramSpec spec;
+    spec.role = signalWait.role;
+    spec.channelId = resource.channelId;
+    spec.remoteXn = resource.remoteXn;
+    spec.localXn = resource.localXn;
+    spec.localGsa = attempt->plan.kernelLocalGsa.num == 0 ? 0 : attempt->plan.kernelLocalGsa.startId;
+    spec.remoteNotifyCke = resource.notifyCke;
+    spec.remoteNotifyMask = resource.remoteNotifyMask == 0 ? 1U : resource.remoteNotifyMask;
+    spec.localWaitCke = resource.localWaitCke == 0 ? resource.notifyCke : resource.localWaitCke;
+    spec.localWaitMask = resource.localWaitMask == 0 ? 1U : resource.localWaitMask;
+    spec.sourceCke = resource.sourceCke;
+    spec.sourceCkeMask = spec.remoteNotifyMask;
+
+    TileXRCcuProgram program;
+    TileXRCcuBarrierProgramReport signalWaitReport;
+    if (TileXRCcuBuildSignalWaitProgram(spec, &program.sync, &signalWaitReport) != TILEXR_SUCCESS) {
+        if (report != nullptr) {
+            report->message = signalWaitReport.message;
+        }
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    if (program.sync.empty() || program.sync.size() > std::numeric_limits<uint16_t>::max()) {
+        if (report != nullptr) {
+            report->message = "invalid direct CCU signal/wait instruction count";
+        }
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    attempt->plan.taskWindows[0].instCnt = static_cast<uint16_t>(program.sync.size());
+
+    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<TileXRCcuTask> 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) {
@@ -866,6 +988,19 @@ void TraceFinalRuntimeTask(size_t taskIndex, const TileXRCcuTask& task)
     std::cerr << "\n";
 }
 
+void ApplyTaskTimeoutOverride(uint16_t taskTimeout, TileXRCcuDirectInstallAttempt* attempt)
+{
+    if (taskTimeout == 0 || attempt == nullptr) {
+        return;
+    }
+    for (auto& task : attempt->package.tasks) {
+        task.timeout = taskTimeout;
+    }
+    for (auto& task : attempt->submitTasks) {
+        task.timeout = taskTimeout;
+    }
+}
+
 std::string FormatSubmitTaskFailure(
     size_t taskIndex,
     int ret,
@@ -983,6 +1118,7 @@ int TileXRCcuSubmitPreparedTasks(
 int RunDirectInstallAttemptImpl(
     const TileXRCcuDirectInstallOptions& options,
     const TileXRCcuDirectMemoryCopySpec* memoryCopy,
+    const TileXRCcuDirectSignalWaitSpec* signalWait,
     TileXRCcuDirectInstallAttempt* attempt,
     TileXRCcuDirectInstallReport* report)
 {
@@ -1027,13 +1163,18 @@ int RunDirectInstallAttemptImpl(
     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<uint32_t>(options.syncInstructionCount, TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT);
+    attempt->resourceRequest.sqeArgCount = memoryCopy == nullptr && signalWait == nullptr ? options.sqeArgCount : 0U;
+    attempt->resourceRequest.syncResourceCount =
+        memoryCopy == nullptr && signalWait == nullptr ? options.syncResourceCount : 1U;
+    attempt->resourceRequest.syncInstructionCount =
+        memoryCopy != nullptr ?
+        std::max<uint32_t>(options.syncInstructionCount, TILEXR_CCU_DIRECT_MEMORY_COPY_INSTRUCTION_COUNT) :
+        signalWait != nullptr ?
+        std::max<uint32_t>(options.syncInstructionCount, SignalWaitInstructionCount(signalWait->role)) :
+        options.syncInstructionCount;
     attempt->resourceRequest.bindingsPerSyncResource = options.bindingsPerSyncResource;
-    attempt->resourceRequest.barrierMode = options.barrierMode;
+    attempt->resourceRequest.barrierMode =
+        signalWait == nullptr ? options.barrierMode : SignalWaitBarrierMode(signalWait->role);
 
     TileXRCcuResourceAllocator allocator;
     if (allocator.Init(attempt->resourceSpec) != TILEXR_SUCCESS) {
@@ -1061,6 +1202,8 @@ int RunDirectInstallAttemptImpl(
                     report->message);
         }
     }
+    attempt->plan.barrierMode =
+        signalWait == nullptr ? attempt->plan.barrierMode : SignalWaitBarrierMode(signalWait->role);
 
     ret = PrepareLowerLayerPlanIfNeeded(options, attempt, report);
     if (ret != TILEXR_SUCCESS) {
@@ -1089,19 +1232,24 @@ int RunDirectInstallAttemptImpl(
     }
 
     TileXRCcuLaunchPackageReport packageReport;
-    ret = memoryCopy == nullptr ?
-        TileXRCcuBuildLaunchPackage(attempt->plan, &attempt->package, &packageReport) :
-        BuildDirectMemoryCopyLaunchPackage(*memoryCopy, attempt, report);
+    ret = memoryCopy != nullptr ?
+        BuildDirectMemoryCopyLaunchPackage(*memoryCopy, attempt, report) :
+        signalWait != nullptr ?
+        BuildDirectSignalWaitLaunchPackage(*signalWait, attempt, report) :
+        TileXRCcuBuildLaunchPackage(attempt->plan, &attempt->package, &packageReport);
     if (ret != TILEXR_SUCCESS) {
         return Fail(
             attempt,
             report,
-            memoryCopy == nullptr ? packageReport.message :
+            memoryCopy == nullptr && signalWait == nullptr ? packageReport.message :
                 (report == nullptr || report->message.empty() ?
-                    "failed to build direct CCU memory copy launch package" :
+                    (memoryCopy != nullptr ?
+                        "failed to build direct CCU memory copy launch package" :
+                        "failed to build direct CCU signal/wait launch package") :
                     report->message));
     }
 
+    ApplyTaskTimeoutOverride(options.taskTimeout, attempt);
     TraceDirectInstallAttempt(*attempt);
 
     ret = TileXRCcuBindLaunchPackageInstallScope(
@@ -1145,6 +1293,7 @@ int RunDirectInstallAttemptImpl(
         attempt->evidence,
         &attempt->submitTasks,
         &attempt->providerReport);
+    ApplyTaskTimeoutOverride(options.taskTimeout, attempt);
     if (submitRet == TILEXR_SUCCESS) {
         return ReturnWithAttemptStatus(TILEXR_SUCCESS, attempt, report);
     }
@@ -1159,7 +1308,7 @@ int TileXRCcuRunDirectInstallAttempt(
     TileXRCcuDirectInstallAttempt* attempt,
     TileXRCcuDirectInstallReport* report)
 {
-    return RunDirectInstallAttemptImpl(options, nullptr, attempt, report);
+    return RunDirectInstallAttemptImpl(options, nullptr, nullptr, attempt, report);
 }
 
 int TileXRCcuRunDirectMemoryCopyInstallAttempt(
@@ -1178,7 +1327,16 @@ int TileXRCcuRunDirectMemoryCopyInstallAttempt(
         }
         return TILEXR_ERROR_PARA_CHECK_FAIL;
     }
-    return RunDirectInstallAttemptImpl(options, &memoryCopy, attempt, report);
+    return RunDirectInstallAttemptImpl(options, &memoryCopy, nullptr, attempt, report);
+}
+
+int TileXRCcuRunDirectSignalWaitInstallAttempt(
+    const TileXRCcuDirectInstallOptions& options,
+    const TileXRCcuDirectSignalWaitSpec& signalWait,
+    TileXRCcuDirectInstallAttempt* attempt,
+    TileXRCcuDirectInstallReport* report)
+{
+    return RunDirectInstallAttemptImpl(options, nullptr, &signalWait, attempt, report);
 }
 
 int TileXRCcuReleaseDirectInstallAttemptResources(TileXRCcuDirectInstallAttempt& attempt)
diff --git a/src/comm/ccu/tilexr_ccu_direct_orchestrator.h b/src/comm/ccu/tilexr_ccu_direct_orchestrator.h
index 60530cf7..7a7271cf 100644
--- a/src/comm/ccu/tilexr_ccu_direct_orchestrator.h
+++ b/src/comm/ccu/tilexr_ccu_direct_orchestrator.h
@@ -9,6 +9,7 @@
 #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_signal_wait_program.h"
 #include "ccu/tilexr_ccu_specs.h"
 
 #include <cstdint>
@@ -43,6 +44,7 @@ struct TileXRCcuDirectInstallOptions {
     uint16_t remoteNotifyCkeStartId = 0;
     uint16_t remoteNotifyCkeCount = 0;
     TileXRCcuBarrierMode barrierMode = TileXRCcuBarrierMode::SyncXn;
+    uint16_t taskTimeout = 0;
     uint32_t deviceId = 0;
     uint32_t rank = 0;
     std::string provider;
@@ -67,6 +69,10 @@ struct TileXRCcuDirectMemoryCopySpec {
     uint64_t lengthBytes = 0;
 };
 
+struct TileXRCcuDirectSignalWaitSpec {
+    TileXRCcuSignalWaitProgramRole role = TileXRCcuSignalWaitProgramRole::Signal;
+};
+
 struct TileXRCcuDirectInstallAttempt {
     TileXRCcuSpecInfo specInfo;
     TileXRCcuResourceSpec resourceSpec;
@@ -123,6 +129,12 @@ int TileXRCcuRunDirectMemoryCopyInstallAttempt(
     TileXRCcuDirectInstallAttempt* attempt,
     TileXRCcuDirectInstallReport* report);
 
+int TileXRCcuRunDirectSignalWaitInstallAttempt(
+    const TileXRCcuDirectInstallOptions& options,
+    const TileXRCcuDirectSignalWaitSpec& signalWait,
+    TileXRCcuDirectInstallAttempt* attempt,
+    TileXRCcuDirectInstallReport* report);
+
 int TileXRCcuReleaseDirectInstallAttemptResources(TileXRCcuDirectInstallAttempt& attempt);
 
 int TileXRCcuSubmitPreparedTasks(
diff --git a/src/comm/ccu/tilexr_ccu_direct_runtime.cpp b/src/comm/ccu/tilexr_ccu_direct_runtime.cpp
index c1291d26..26b91e71 100644
--- a/src/comm/ccu/tilexr_ccu_direct_runtime.cpp
+++ b/src/comm/ccu/tilexr_ccu_direct_runtime.cpp
@@ -42,6 +42,8 @@ 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_RESOURCE_WINDOW_EID_INDEX_ENV =
+    "TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX";
 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 =
@@ -52,6 +54,7 @@ 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 uint8_t TILEXR_CCU_DIRECT_DEFAULT_DIE_ID = 0;
 constexpr uint64_t TILEXR_CCU_UB_MEM_PAGE_SIZE = 4096ULL;
 constexpr uint32_t TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT = 8;
 
@@ -172,6 +175,15 @@ bool HasRankedEnv(const char* base, int rank)
     return SelectRankedEnv(base, rank) != nullptr;
 }
 
+std::array<uint8_t, TILEXR_CCU_EID_BYTES> CopyRawEid(const TileXRCcuHccpEid& eid)
+{
+    std::array<uint8_t, TILEXR_CCU_EID_BYTES> copied {};
+    for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) {
+        copied[i] = eid.raw[i];
+    }
+    return copied;
+}
+
 bool IsRaCtxResourceWindowRegistrationMode()
 {
     const char* value = std::getenv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE_ENV);
@@ -201,6 +213,63 @@ void TraceEndpointRouteStep(const std::string& message)
     }
 }
 
+void TraceTaskKillCleanup(uint8_t dieId, int ret, const TileXRCcuDriverAdapterReport& report)
+{
+    if (!TraceEndpointRoute()) {
+        return;
+    }
+    std::cerr << "TileXRDirectCcuTrace taskKillCleanup"
+              << " dieId=" << static_cast<uint32_t>(dieId)
+              << " ret=" << ret
+              << " opcode=" << report.opcode
+              << " driverRet=" << report.driverRet
+              << " opRet=" << report.opRet
+              << " message=\"" << report.message << "\""
+              << std::endl;
+}
+
+void TraceRaCtxEidInfos(const std::vector<TileXRCcuHccpDevEidInfo>& eidInfos)
+{
+    if (!TraceEndpointRoute()) {
+        return;
+    }
+    for (size_t i = 0; i < eidInfos.size(); ++i) {
+        std::cerr << "TileXRDirectCcuTrace endpointRoute raCtxEidInfo"
+                  << " ordinal=" << i
+                  << " eidIndex=" << eidInfos[i].eidIndex
+                  << " funcId=" << eidInfos[i].funcId
+                  << " eid=" << FormatEndpointEid(CopyRawEid(eidInfos[i].eid))
+                  << std::endl;
+    }
+}
+
+bool SelectRaCtxResourceWindowEidInfo(
+    int rank,
+    const std::vector<TileXRCcuHccpDevEidInfo>& eidInfos,
+    TileXRCcuHccpDevEidInfo* selectedEid)
+{
+    if (eidInfos.empty() || selectedEid == nullptr) {
+        return false;
+    }
+    TraceRaCtxEidInfos(eidInfos);
+    const char* configured = SelectRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_ENV, rank);
+    if (configured == nullptr) {
+        *selectedEid = eidInfos[0];
+        return true;
+    }
+    uint64_t configuredIndex = 0;
+    if (!ParseUnsignedEnv(configured, &configuredIndex) || configuredIndex > 0xffffffffULL) {
+        return false;
+    }
+    for (const auto& eidInfo : eidInfos) {
+        if (eidInfo.eidIndex == static_cast<uint32_t>(configuredIndex)) {
+            *selectedEid = eidInfo;
+            return true;
+        }
+    }
+    return false;
+}
+
 bool IsRaCtxLoopEndpointRouteCollectionMode()
 {
     const char* value = std::getenv(TILEXR_CCU_DIRECT_ENDPOINT_ROUTE_COLLECTION_MODE_ENV);
@@ -229,6 +298,12 @@ bool HasRaCtxResourceWindowSymbols(const TileXRCcuHccpLoader& loader)
         loader.RaCtxLmemUnregister != nullptr;
 }
 
+bool HasRaCtxRemoteMemoryImportSymbols(const TileXRCcuHccpLoader& loader)
+{
+    return loader.RaCtxRmemImport != nullptr &&
+        loader.RaCtxRmemUnimport != nullptr;
+}
+
 bool HasRaCtxEndpointRouteSymbols(const TileXRCcuHccpLoader& loader)
 {
     return loader.RaCtxCqCreate != nullptr &&
@@ -509,6 +584,13 @@ int TileXRCcuDirectRuntime::Init(
     }
 
     initialized_ = true;
+    TileXRCcuDriverAdapter adapter;
+    TileXRCcuDriverAdapterReport adapterReport;
+    int cleanupRet = CreateDriverAdapter(&adapter, &adapterReport);
+    if (cleanupRet == TILEXR_SUCCESS) {
+        cleanupRet = adapter.CleanTaskKillState(TILEXR_CCU_DIRECT_DEFAULT_DIE_ID, &adapterReport);
+    }
+    TraceTaskKillCleanup(TILEXR_CCU_DIRECT_DEFAULT_DIE_ID, cleanupRet, adapterReport);
     if (report != nullptr) {
         report->initialized = true;
         report->raInitialized = true;
@@ -537,6 +619,7 @@ void TileXRCcuDirectRuntime::Shutdown()
     endpointQpHandle_ = nullptr;
     endpointRemoteQpHandle_ = nullptr;
     endpointPeerRemoteQpHandles_.clear();
+    importedRemoteMemoryBuffers_.clear();
     endpointQpKey_ = TileXRCcuHccpQpKey{};
     endpointQpKeyValid_ = false;
     endpointRouteBound_ = false;
@@ -637,6 +720,11 @@ int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBuffer(uint64_t resourceAddr)
     if (!IsAvailable()) {
         return TILEXR_ERROR_NOT_INITIALIZED;
     }
+    if (resourceWindowRegistered_ &&
+        localResourceWindow_.addr == resourceAddr &&
+        localResourceWindow_.raCtxRegistered == IsRaCtxResourceWindowRegistrationMode()) {
+        return TILEXR_SUCCESS;
+    }
     ReleaseRegisteredResourceWindow();
     if (!cachedBasicInfoValid_ || resourceAddr == 0) {
         return TILEXR_ERROR_PARA_CHECK_FAIL;
@@ -666,6 +754,127 @@ int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBuffer(uint64_t resourceAddr)
     return TILEXR_SUCCESS;
 }
 
+int TileXRCcuDirectRuntime::RegisterMemoryBuffer(
+    uint64_t addr,
+    uint64_t bytes,
+    TileXRCcuRegisteredMemoryBufferInfo* info)
+{
+    if (info == nullptr) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    *info = TileXRCcuRegisteredMemoryBufferInfo {};
+    if (!IsAvailable() || !resourceWindowRegistered_ || localResourceWindow_.raCtxHandle == nullptr) {
+        return TILEXR_ERROR_NOT_INITIALIZED;
+    }
+    if (addr == 0 || bytes == 0 || !HasRaCtxResourceWindowSymbols(loader_)) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+
+    void* tokenIdHandle = nullptr;
+    void* lmemHandle = nullptr;
+    TileXRCcuHccpTokenId allocatedToken {};
+    int ret = loader_.RaCtxTokenIdAlloc(localResourceWindow_.raCtxHandle, &allocatedToken, &tokenIdHandle);
+    if (ret != 0 || tokenIdHandle == nullptr) {
+        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(localResourceWindow_.raCtxHandle, tokenIdHandle);
+        }
+        return TILEXR_ERROR_MKIRT;
+    }
+
+    const uint64_t alignedAddr = AlignResourceWindowAddr(addr);
+    const uint64_t alignedBytes = bytes + (addr - 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(localResourceWindow_.raCtxHandle, &mr, &lmemHandle);
+    if (ret != 0 || lmemHandle == nullptr) {
+        if (lmemHandle != nullptr && loader_.RaCtxLmemUnregister != nullptr) {
+            (void)loader_.RaCtxLmemUnregister(localResourceWindow_.raCtxHandle, lmemHandle);
+        }
+        if (loader_.RaCtxTokenIdFree != nullptr) {
+            (void)loader_.RaCtxTokenIdFree(localResourceWindow_.raCtxHandle, tokenIdHandle);
+        }
+        return TILEXR_ERROR_MKIRT;
+    }
+
+    TileXRCcuRegisteredMemoryBufferInfo registeredInfo {};
+    registeredInfo.addr = addr;
+    registeredInfo.bytes = bytes;
+    registeredInfo.alignedAddr = alignedAddr;
+    registeredInfo.alignedBytes = alignedBytes;
+    registeredInfo.targetSegVa = mr.out.ub.targetSegHandle + (addr - alignedAddr);
+    registeredInfo.rawTokenId = mr.out.ub.tokenId != 0 ? mr.out.ub.tokenId : allocatedToken.tokenId;
+    registeredInfo.tokenId = registeredInfo.rawTokenId >> TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT;
+    registeredInfo.tokenValue = tokenValue;
+    registeredInfo.key = mr.out.key;
+    registeredInfo.tokenIdHandle = tokenIdHandle;
+    registeredInfo.lmemHandle = lmemHandle;
+    registeredInfo.valid = true;
+
+    registeredMemoryBuffers_.push_back(registeredInfo);
+    *info = registeredInfo;
+    return TILEXR_SUCCESS;
+}
+
+int TileXRCcuDirectRuntime::ImportRemoteMemoryBuffer(
+    const TileXRCcuRemoteMemoryBufferImportRequest& request,
+    TileXRCcuImportedRemoteMemoryBufferInfo* info)
+{
+    if (info == nullptr) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    *info = TileXRCcuImportedRemoteMemoryBufferInfo {};
+    if (!IsAvailable() || !resourceWindowRegistered_ || localResourceWindow_.raCtxHandle == nullptr) {
+        return TILEXR_ERROR_NOT_INITIALIZED;
+    }
+    if (!request.valid || request.key.size == 0 || !HasRaCtxRemoteMemoryImportSymbols(loader_)) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+
+    TileXRCcuHccpMrImportInfo mr {};
+    mr.in.key = request.key;
+    mr.in.ub.tokenValue = request.tokenValue;
+    mr.in.ub.mappingAddr = 0;
+    mr.in.ub.flags.value = 0;
+    mr.in.ub.flags.bs.access = TILEXR_CCU_HCCP_MEM_SEG_ACCESS_DEFAULT;
+
+    void* rmemHandle = nullptr;
+    const int ret = loader_.RaCtxRmemImport(localResourceWindow_.raCtxHandle, &mr, &rmemHandle);
+    if (ret != 0 || rmemHandle == nullptr || mr.out.ub.targetSegHandle == 0) {
+        if (rmemHandle != nullptr && loader_.RaCtxRmemUnimport != nullptr) {
+            (void)loader_.RaCtxRmemUnimport(localResourceWindow_.raCtxHandle, rmemHandle);
+        }
+        return TILEXR_ERROR_MKIRT;
+    }
+
+    TileXRCcuImportedRemoteMemoryBufferInfo imported {};
+    imported.addr = request.addr;
+    imported.bytes = request.bytes;
+    imported.targetSegVa = mr.out.ub.targetSegHandle + request.offset;
+    imported.rmemHandle = rmemHandle;
+    imported.valid = true;
+    importedRemoteMemoryBuffers_.push_back(imported);
+    *info = imported;
+    return TILEXR_SUCCESS;
+}
+
 int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBufferWithRaCtx(
     uint64_t resourceAddr,
     uint64_t resourceBytes)
@@ -695,13 +904,18 @@ int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBufferWithRaCtx(
     void* tokenIdHandle = nullptr;
     void* lmemHandle = nullptr;
 
+    TileXRCcuHccpDevEidInfo selectedEid {};
+    if (!SelectRaCtxResourceWindowEidInfo(options_.rank, eidInfos, &selectedEid)) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+
     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;
+    ctxAttr.ub.eidIndex = selectedEid.eidIndex;
+    ctxAttr.ub.eid = selectedEid.eid;
 
     ret = loader_.RaCtxInit(&ctxCfg, &ctxAttr, &ctxHandle);
     if (ret != 0 || ctxHandle == nullptr) {
@@ -771,9 +985,11 @@ int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBufferWithRaCtx(
     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_.eid[i] = selectedEid.eid.raw[i];
     }
-    localResourceWindow_.eidIndex = eidInfos[0].eidIndex;
+    localResourceWindow_.eidIndex = selectedEid.eidIndex;
+    localResourceWindow_.funcId = selectedEid.funcId;
+    localResourceWindow_.funcIdValid = true;
     localResourceWindow_.raCtxRegistered = true;
     resourceWindowRegistered_ = true;
     return TILEXR_SUCCESS;
@@ -781,6 +997,8 @@ int TileXRCcuDirectRuntime::RegisterCcuResourceRmaBufferWithRaCtx(
 
 void TileXRCcuDirectRuntime::ReleaseRegisteredResourceWindow()
 {
+    ReleaseImportedRemoteMemoryBuffers();
+    ReleaseRegisteredMemoryBuffers();
     ReleaseLocalEndpointRoute();
     if (localResourceWindow_.raCtxRegistered) {
         if (localResourceWindow_.lmemHandle != nullptr &&
@@ -803,6 +1021,33 @@ void TileXRCcuDirectRuntime::ReleaseRegisteredResourceWindow()
     localVerifiedEndpointRouteValid_ = false;
 }
 
+void TileXRCcuDirectRuntime::ReleaseImportedRemoteMemoryBuffers()
+{
+    if (localResourceWindow_.raCtxHandle != nullptr && loader_.RaCtxRmemUnimport != nullptr) {
+        for (auto it = importedRemoteMemoryBuffers_.rbegin(); it != importedRemoteMemoryBuffers_.rend(); ++it) {
+            if (it->rmemHandle != nullptr) {
+                (void)loader_.RaCtxRmemUnimport(localResourceWindow_.raCtxHandle, it->rmemHandle);
+            }
+        }
+    }
+    importedRemoteMemoryBuffers_.clear();
+}
+
+void TileXRCcuDirectRuntime::ReleaseRegisteredMemoryBuffers()
+{
+    if (localResourceWindow_.raCtxHandle != nullptr) {
+        for (auto it = registeredMemoryBuffers_.rbegin(); it != registeredMemoryBuffers_.rend(); ++it) {
+            if (it->lmemHandle != nullptr && loader_.RaCtxLmemUnregister != nullptr) {
+                (void)loader_.RaCtxLmemUnregister(localResourceWindow_.raCtxHandle, it->lmemHandle);
+            }
+            if (it->tokenIdHandle != nullptr && loader_.RaCtxTokenIdFree != nullptr) {
+                (void)loader_.RaCtxTokenIdFree(localResourceWindow_.raCtxHandle, it->tokenIdHandle);
+            }
+        }
+    }
+    registeredMemoryBuffers_.clear();
+}
+
 void TileXRCcuDirectRuntime::ReleasePeerEndpointImports()
 {
     if (localResourceWindow_.raCtxHandle != nullptr && loader_.RaCtxQpUnimport != nullptr) {
@@ -956,6 +1201,7 @@ int TileXRCcuDirectRuntime::CollectLocalEndpointRouteWithRaCtxOnce(
         std::cerr << "TileXRDirectCcuTrace endpointRoute qpCreate ret=" << ret
                   << " qp=" << endpointQpHandle_
                   << " keySize=" << static_cast<uint32_t>(qpInfo.key.size)
+                  << " id=" << qpInfo.ub.id
                   << " dbAddr=0x" << std::hex << qpInfo.ub.dbAddr
                   << " dbTokenId=0x" << qpInfo.ub.dbTokenId
                   << std::dec << std::endl;
@@ -1047,6 +1293,7 @@ int TileXRCcuDirectRuntime::CollectLocalEndpointRouteWithRaCtxOnce(
     collected.doorbellTokenId = qpInfo.ub.dbTokenId >> TILEXR_CCU_URMA_TOKEN_ID_RIGHT_SHIFT;
     collected.doorbellTokenValue = localResourceWindow_.tokenValue;
     collected.sqDepth = sqDepth;
+    collected.startJettyId = static_cast<uint16_t>(qpInfo.ub.id);
     collected.endpointRouteVerified = true;
     if (!HasCompleteEndpointRoute(collected)) {
         TraceEndpointRouteStep("collected route incomplete");
@@ -1059,6 +1306,7 @@ int TileXRCcuDirectRuntime::CollectLocalEndpointRouteWithRaCtxOnce(
                   << " doorbellVa=0x" << collected.doorbellVa
                   << " doorbellTokenId=0x" << collected.doorbellTokenId
                   << std::dec
+                  << " startJettyId=" << collected.startJettyId
                   << " sqDepth=" << collected.sqDepth << std::endl;
     }
     *route = collected;
@@ -1272,6 +1520,7 @@ int TileXRCcuDirectRuntime::ImportPeerEndpointRoute(
     importedRoute->doorbellTokenId = localVerifiedEndpointRoute_.doorbellTokenId;
     importedRoute->doorbellTokenValue = localVerifiedEndpointRoute_.doorbellTokenValue;
     importedRoute->sqDepth = localVerifiedEndpointRoute_.sqDepth;
+    importedRoute->startJettyId = localVerifiedEndpointRoute_.startJettyId;
     importedRoute->endpointRouteVerified = true;
     return HasCompleteEndpointRoute(*importedRoute) ? TILEXR_SUCCESS : TILEXR_ERROR_PARA_CHECK_FAIL;
 }
@@ -1302,6 +1551,8 @@ int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemot
         localResourceWindow_.rawTokenId,
         localResourceWindow_.tokenValue,
     };
+    local.funcId = localResourceWindow_.funcId;
+    local.funcIdValid = localResourceWindow_.funcIdValid;
     if (localVerifiedEndpointRouteValid_) {
         local.remoteEid = localVerifiedEndpointRoute_.remoteEid;
         local.tpn = localVerifiedEndpointRoute_.tpn;
@@ -1309,6 +1560,7 @@ int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemot
         local.doorbellTokenId = localVerifiedEndpointRoute_.doorbellTokenId;
         local.doorbellTokenValue = localVerifiedEndpointRoute_.doorbellTokenValue;
         local.sqDepth = localVerifiedEndpointRoute_.sqDepth;
+        local.startJettyId = localVerifiedEndpointRoute_.startJettyId;
         local.qpKey = endpointQpKey_;
         local.psn = endpointPsn_;
         local.endpointRouteVerified = true;
@@ -1383,6 +1635,8 @@ int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemot
         remote.memoryTokenId = peerWindow.tokenId;
         remote.rawMemoryTokenId = peerWindow.rawTokenId;
         remote.memoryTokenValue = peerWindow.tokenValue;
+        remote.localPfeId = localResourceWindow_.funcId;
+        remote.localPfeIdValid = localResourceWindow_.funcIdValid;
         remote.remoteEid = ReverseEndpointEid(peerWindow.remoteEid);
         TileXRCcuLowerLayerTransportRoute importedRoute;
         const uint64_t localTpForPeer = peerTpHandlesReady ?
@@ -1407,12 +1661,14 @@ int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemot
             remote.doorbellTokenId = importedRoute.doorbellTokenId;
             remote.doorbellTokenValue = importedRoute.doorbellTokenValue;
             remote.sqDepth = importedRoute.sqDepth;
+            remote.startJettyId = importedRoute.startJettyId;
         } else {
             remote.tpn = peerWindow.tpn;
             remote.doorbellVa = peerWindow.doorbellVa;
             remote.doorbellTokenId = peerWindow.doorbellTokenId;
             remote.doorbellTokenValue = peerWindow.doorbellTokenValue;
             remote.sqDepth = peerWindow.sqDepth;
+            remote.startJettyId = peerWindow.startJettyId;
             if (!useImportedPeerRoute && TraceEndpointRoute() && peerWindow.endpointRouteVerified) {
                 std::cerr << "TileXRDirectCcuTrace endpointRoute usePeerExportedRoute"
                           << " peerRank=" << peer
@@ -1453,10 +1709,13 @@ int TileXRCcuDirectRuntime::ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemot
                       << " remoteCcuVaOffset=0x" << remoteCcuVaOffset
                       << " doorbellVa=0x" << remote.doorbellVa
                       << " localDoorbellVa=0x" << remote.localDoorbellVa
+                      << " startJettyId=0x" << remote.startJettyId
                       << " memoryTokenId=0x" << remote.memoryTokenId
                       << " rawMemoryTokenId=0x" << remote.rawMemoryTokenId
                       << " memoryTokenValue=0x" << remote.memoryTokenValue
                       << std::dec
+                      << " localPfeId=" << remote.localPfeId
+                      << " localPfeIdValid=" << (remote.localPfeIdValid ? 1 : 0)
                       << " endpointRouteVerified=" << (remote.endpointRouteVerified ? 1 : 0)
                       << " channelResourceOwnerVerified=" << (remote.channelResourceOwnerVerified ? 1 : 0)
                       << " transportResourceExchangeVerified="
diff --git a/src/comm/ccu/tilexr_ccu_direct_runtime.h b/src/comm/ccu/tilexr_ccu_direct_runtime.h
index 3e358f06..83f73dac 100644
--- a/src/comm/ccu/tilexr_ccu_direct_runtime.h
+++ b/src/comm/ccu/tilexr_ccu_direct_runtime.h
@@ -55,6 +55,8 @@ struct TileXRCcuLocalResourceWindowInfo {
     void* lmemHandle = nullptr;
     std::array<uint8_t, TILEXR_CCU_EID_BYTES> eid {};
     uint32_t eidIndex = 0;
+    uint32_t funcId = 0;
+    bool funcIdValid = false;
     bool raCtxRegistered = false;
 };
 
@@ -70,6 +72,9 @@ struct TileXRCcuResourceWindowExchange {
     uint32_t doorbellTokenId = 0;
     uint32_t doorbellTokenValue = 0;
     uint32_t sqDepth = 0;
+    uint16_t startJettyId = 0;
+    uint32_t funcId = 0;
+    bool funcIdValid = false;
     TileXRCcuHccpQpKey qpKey {};
     uint32_t psn = 0;
     bool endpointRouteVerified = false;
@@ -77,6 +82,41 @@ struct TileXRCcuResourceWindowExchange {
     bool transportResourceExchangeVerified = false;
 };
 
+struct TileXRCcuRegisteredMemoryBufferInfo {
+    uint64_t addr = 0;
+    uint64_t bytes = 0;
+    uint64_t alignedAddr = 0;
+    uint64_t alignedBytes = 0;
+    uint64_t targetSegVa = 0;
+    uint32_t tokenId = 0;
+    uint32_t rawTokenId = 0;
+    uint32_t tokenValue = 0;
+    TileXRCcuHccpMemKey key {};
+    void* tokenIdHandle = nullptr;
+    void* lmemHandle = nullptr;
+    bool valid = false;
+};
+
+struct TileXRCcuRemoteMemoryBufferImportRequest {
+    uint64_t addr = 0;
+    uint64_t bytes = 0;
+    uint64_t alignedAddr = 0;
+    uint64_t offset = 0;
+    uint32_t tokenId = 0;
+    uint32_t rawTokenId = 0;
+    uint32_t tokenValue = 0;
+    TileXRCcuHccpMemKey key {};
+    bool valid = false;
+};
+
+struct TileXRCcuImportedRemoteMemoryBufferInfo {
+    uint64_t addr = 0;
+    uint64_t bytes = 0;
+    uint64_t targetSegVa = 0;
+    void* rmemHandle = nullptr;
+    bool valid = false;
+};
+
 struct TileXRCcuDirectRuntimeReport {
     bool initialized = false;
     bool raInitialized = false;
@@ -96,6 +136,10 @@ class TileXRCcuDirectRuntime {
     int QueryBasicInfo(uint8_t dieId, TileXRCcuBasicInfo* basicInfo, TileXRCcuDriverAdapterReport* report);
     int CreateDriverAdapter(TileXRCcuDriverAdapter* adapter, TileXRCcuDriverAdapterReport* report);
     int RegisterCcuResourceRmaBuffer(uint64_t resourceAddr);
+    int RegisterMemoryBuffer(uint64_t addr, uint64_t bytes, TileXRCcuRegisteredMemoryBufferInfo* info);
+    int ImportRemoteMemoryBuffer(
+        const TileXRCcuRemoteMemoryBufferImportRequest& request,
+        TileXRCcuImportedRemoteMemoryBufferInfo* info);
     int RefreshLocalVerifiedEndpointRoute(TileXRCcuDirectRuntimeReport* report);
     int ConfigureLocalVerifiedEndpointRoute(const TileXRCcuLowerLayerTransportRoute& route);
     int ExportLocalCcuRmaBuffer(TileXRCcuLocalResourceWindowInfo* info) const;
@@ -122,6 +166,8 @@ class TileXRCcuDirectRuntime {
         uint32_t peerPsn,
         TileXRCcuLowerLayerTransportRoute* importedRoute);
     void ReleasePeerEndpointImports();
+    void ReleaseImportedRemoteMemoryBuffers();
+    void ReleaseRegisteredMemoryBuffers();
     void ReleaseRegisteredResourceWindow();
     void ReleaseLocalEndpointRoute();
 
@@ -139,6 +185,8 @@ class TileXRCcuDirectRuntime {
     void* endpointQpHandle_ = nullptr;
     void* endpointRemoteQpHandle_ = nullptr;
     std::vector<void*> endpointPeerRemoteQpHandles_;
+    std::vector<TileXRCcuRegisteredMemoryBufferInfo> registeredMemoryBuffers_;
+    std::vector<TileXRCcuImportedRemoteMemoryBufferInfo> importedRemoteMemoryBuffers_;
     TileXRCcuHccpQpKey endpointQpKey_ = {};
     bool endpointQpKeyValid_ = false;
     bool endpointRouteBound_ = false;
diff --git a/src/comm/ccu/tilexr_ccu_driver_adapter.cpp b/src/comm/ccu/tilexr_ccu_driver_adapter.cpp
index 3c2c1bb0..f8cf6a49 100644
--- a/src/comm/ccu/tilexr_ccu_driver_adapter.cpp
+++ b/src/comm/ccu/tilexr_ccu_driver_adapter.cpp
@@ -321,6 +321,34 @@ int TileXRCcuDriverAdapter::ReadInstructions(
     return TILEXR_SUCCESS;
 }
 
+int TileXRCcuDriverAdapter::ReadMissionContext(
+    uint8_t dieId,
+    uint8_t missionId,
+    void* missionContext,
+    uint32_t missionContextBytes,
+    TileXRCcuDriverAdapterReport* report) const
+{
+    ResetReport(report);
+    if (missionContext == nullptr || missionContextBytes == 0 ||
+        missionContextBytes > TILEXR_CCU_DATA_ARRAY_SLOT_BYTES) {
+        return Fail(report, "invalid CCU mission context readback buffer");
+    }
+
+    TileXRCcuCustomChannelIn in;
+    InitRequest(dieId, TILEXR_CCU_U_OP_GET_MISSION_CTX, &in);
+    in.offsetStartIdx = missionId;
+    in.data.dataInfo.dataArraySize = 1;
+    in.data.dataInfo.dataLen = missionContextBytes;
+
+    TileXRCcuCustomChannelOut out;
+    const int ret = CallPrepared(dieId, TILEXR_CCU_U_OP_GET_MISSION_CTX, in, &out, report);
+    if (ret != TILEXR_SUCCESS) {
+        return ret;
+    }
+    std::memcpy(missionContext, &out.data.dataInfo.dataArray[0], missionContextBytes);
+    return TILEXR_SUCCESS;
+}
+
 int TileXRCcuDriverAdapter::InstallInstructions(
     uint8_t dieId,
     uint16_t instructionStartId,
@@ -393,6 +421,12 @@ int TileXRCcuDriverAdapter::InstallMsidToken(
     return CallPrepared(dieId, TILEXR_CCU_U_OP_SET_MSID_TOKEN, in, &out, report);
 }
 
+int TileXRCcuDriverAdapter::CleanTaskKillState(uint8_t dieId, TileXRCcuDriverAdapterReport* report) const
+{
+    TileXRCcuCustomChannelOut out;
+    return Call(dieId, TILEXR_CCU_U_OP_CLEAN_TASKKILL_STATE, &out, report);
+}
+
 int TileXRCcuDriverAdapter::InstallPfeCtx(
     uint8_t dieId,
     uint32_t pfeOffset,
diff --git a/src/comm/ccu/tilexr_ccu_driver_adapter.h b/src/comm/ccu/tilexr_ccu_driver_adapter.h
index 65534113..53ed1be3 100644
--- a/src/comm/ccu/tilexr_ccu_driver_adapter.h
+++ b/src/comm/ccu/tilexr_ccu_driver_adapter.h
@@ -17,7 +17,10 @@ namespace TileXR {
 constexpr uint32_t TILEXR_CCU_U_OP_GET_BASIC_INFO = 11;
 constexpr uint32_t TILEXR_CCU_U_OP_GET_DIE_WORKING = 15;
 constexpr uint32_t TILEXR_CCU_U_OP_GET_INSTRUCTION = 201;
+constexpr uint32_t TILEXR_CCU_U_OP_GET_MISSION_CTX = 208;
 constexpr uint32_t TILEXR_CCU_U_OP_SET_MSID_TOKEN = 53;
+constexpr uint32_t TILEXR_CCU_U_OP_SET_TASKKILL = 54;
+constexpr uint32_t TILEXR_CCU_U_OP_CLEAN_TASKKILL_STATE = 55;
 constexpr uint32_t TILEXR_CCU_U_OP_SET_INSTRUCTION = 251;
 constexpr uint32_t TILEXR_CCU_U_OP_SET_XN = 253;
 constexpr uint32_t TILEXR_CCU_U_OP_SET_CKE = 254;
@@ -78,6 +81,12 @@ class TileXRCcuDriverAdapter {
         uint32_t instructionCount,
         uint32_t instructionBytes,
         TileXRCcuDriverAdapterReport* report) const;
+    int ReadMissionContext(
+        uint8_t dieId,
+        uint8_t missionId,
+        void* missionContext,
+        uint32_t missionContextBytes,
+        TileXRCcuDriverAdapterReport* report) const;
     int InstallInstructions(
         uint8_t dieId,
         uint16_t instructionStartId,
@@ -99,6 +108,7 @@ class TileXRCcuDriverAdapter {
         uint32_t tokenId,
         uint32_t tokenValue,
         TileXRCcuDriverAdapterReport* report) const;
+    int CleanTaskKillState(uint8_t dieId, TileXRCcuDriverAdapterReport* report) const;
     int InstallPfeCtx(
         uint8_t dieId,
         uint32_t pfeOffset,
diff --git a/src/comm/ccu/tilexr_ccu_hccp_loader.cpp b/src/comm/ccu/tilexr_ccu_hccp_loader.cpp
index 3f28935a..fcf96e0d 100644
--- a/src/comm/ccu/tilexr_ccu_hccp_loader.cpp
+++ b/src/comm/ccu/tilexr_ccu_hccp_loader.cpp
@@ -141,6 +141,8 @@ int TileXRCcuHccpLoader::Load(TileXRCcuHccpLoaderReport* report)
     LoadOptionalSymbol(raHandle_, RaCtxTokenIdFree, "RaCtxTokenIdFree", "ra_ctx_token_id_free");
     LoadOptionalSymbol(raHandle_, RaCtxLmemRegister, "RaCtxLmemRegister", "ra_ctx_lmem_register");
     LoadOptionalSymbol(raHandle_, RaCtxLmemUnregister, "RaCtxLmemUnregister", "ra_ctx_lmem_unregister");
+    LoadOptionalSymbol(raHandle_, RaCtxRmemImport, "RaCtxRmemImport", "ra_ctx_rmem_import");
+    LoadOptionalSymbol(raHandle_, RaCtxRmemUnimport, "RaCtxRmemUnimport", "ra_ctx_rmem_unimport");
     LoadOptionalSymbol(raHandle_, RaGetSecRandom, "RaGetSecRandom", "ra_get_sec_random");
     LoadOptionalSymbol(raHandle_, RaCtxChanCreate, "RaCtxChanCreate", "ra_ctx_chan_create");
     LoadOptionalSymbol(raHandle_, RaCtxChanDestroy, "RaCtxChanDestroy", "ra_ctx_chan_destroy");
@@ -230,6 +232,8 @@ void TileXRCcuHccpLoader::Unload()
     RaCtxTokenIdFree = nullptr;
     RaCtxLmemRegister = nullptr;
     RaCtxLmemUnregister = nullptr;
+    RaCtxRmemImport = nullptr;
+    RaCtxRmemUnimport = nullptr;
     RaGetSecRandom = nullptr;
     RaCtxChanCreate = nullptr;
     RaCtxChanDestroy = nullptr;
diff --git a/src/comm/ccu/tilexr_ccu_hccp_loader.h b/src/comm/ccu/tilexr_ccu_hccp_loader.h
index 93ea1348..ef0f30ef 100644
--- a/src/comm/ccu/tilexr_ccu_hccp_loader.h
+++ b/src/comm/ccu/tilexr_ccu_hccp_loader.h
@@ -73,6 +73,8 @@ class TileXRCcuHccpLoader {
     TileXRCcuRaCtxTokenIdFreeFunc RaCtxTokenIdFree = nullptr;
     TileXRCcuRaCtxLmemRegisterFunc RaCtxLmemRegister = nullptr;
     TileXRCcuRaCtxLmemUnregisterFunc RaCtxLmemUnregister = nullptr;
+    TileXRCcuRaCtxRmemImportFunc RaCtxRmemImport = nullptr;
+    TileXRCcuRaCtxRmemUnimportFunc RaCtxRmemUnimport = nullptr;
     TileXRCcuRaGetSecRandomFunc RaGetSecRandom = nullptr;
     TileXRCcuRaCtxChanCreateFunc RaCtxChanCreate = nullptr;
     TileXRCcuRaCtxChanDestroyFunc RaCtxChanDestroy = nullptr;
diff --git a/src/comm/ccu/tilexr_ccu_hccp_types.h b/src/comm/ccu/tilexr_ccu_hccp_types.h
index abc75b1a..58460d30 100644
--- a/src/comm/ccu/tilexr_ccu_hccp_types.h
+++ b/src/comm/ccu/tilexr_ccu_hccp_types.h
@@ -225,6 +225,42 @@ struct TileXRCcuHccpMrRegInfo {
     TileXRCcuHccpMemRegInfo out;
 };
 
+union TileXRCcuHccpImportSegFlag {
+    struct {
+        uint32_t cacheable : 1;
+        uint32_t access : 6;
+        uint32_t mapping : 1;
+        uint32_t reserved : 24;
+    } bs;
+    uint32_t value;
+};
+
+struct TileXRCcuHccpMemImportAttr {
+    TileXRCcuHccpMemKey key;
+    union {
+        struct {
+            TileXRCcuHccpImportSegFlag flags;
+            uint64_t mappingAddr;
+            uint32_t tokenValue;
+        } ub;
+    };
+    uint32_t resv[4];
+};
+
+struct TileXRCcuHccpMemImportInfo {
+    union {
+        struct {
+            uint64_t targetSegHandle;
+        } ub;
+    };
+    uint32_t resv[4];
+};
+
+struct TileXRCcuHccpMrImportInfo {
+    TileXRCcuHccpMemImportAttr in;
+    TileXRCcuHccpMemImportInfo out;
+};
+
 union TileXRCcuHccpDataPlaneCstmFlag {
     struct {
         uint32_t pollCqCstm : 1;
@@ -543,6 +579,11 @@ using TileXRCcuRaCtxLmemRegisterFunc = int (*)(
     TileXRCcuHccpMrRegInfo* mr,
     void** lmemHandle);
 using TileXRCcuRaCtxLmemUnregisterFunc = int (*)(void* ctx, void* lmemHandle);
+using TileXRCcuRaCtxRmemImportFunc = int (*)(
+    void* ctx,
+    TileXRCcuHccpMrImportInfo* mr,
+    void** rmemHandle);
+using TileXRCcuRaCtxRmemUnimportFunc = int (*)(void* ctx, void* rmemHandle);
 using TileXRCcuRaGetSecRandomFunc = int (*)(TileXRCcuRaInfo* info, uint32_t* value);
 using TileXRCcuRaCtxChanCreateFunc = int (*)(void* ctx, TileXRCcuHccpChanInfo* info, void** chanHandle);
 using TileXRCcuRaCtxChanDestroyFunc = int (*)(void* ctx, void* chanHandle);
diff --git a/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.cpp b/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.cpp
index 095c38aa..054e9272 100644
--- a/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.cpp
+++ b/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.cpp
@@ -457,7 +457,9 @@ int TileXRCcuBuildLowerLayerTransportTemplate(
         result.msidToken.valid = true;
     }
     result.dieId = basicInfo.dieId;
-    result.pfeId = allocation.channels.startId;
+    result.pfeId = remoteCcuBuffers.front().localPfeIdValid ?
+        remoteCcuBuffers.front().localPfeId :
+        allocation.channels.startId;
     result.pfeOffset = SelectLowerLayerPfeOffset(basicInfo.dieId, result.pfeId);
     result.startJettyId = TILEXR_CCU_DEFAULT_START_JETTY_ID;
     result.startLocalJettyCtxId = TILEXR_CCU_DEFAULT_START_LOCAL_JETTY_CTX_ID;
@@ -469,6 +471,8 @@ int TileXRCcuBuildLowerLayerTransportTemplate(
     result.routes.reserve(remoteCcuBuffers.size());
 
     const uint16_t wqeBasicBlockStride = SelectLowerLayerWqeBasicBlockStride();
+    uint16_t verifiedStartJettyId = 0;
+    uint32_t verifiedJettyEnd = 0;
     for (uint32_t i = 0; i < remoteCcuBuffers.size(); ++i) {
         const auto& remoteCcuBuffer = remoteCcuBuffers[i];
         if (remoteCcuBuffer.remoteCcuVa == 0) {
@@ -509,10 +513,24 @@ int TileXRCcuBuildLowerLayerTransportTemplate(
             route.localDoorbellTokenId = remoteCcuBuffer.localDoorbellTokenId;
             route.localDoorbellTokenValue = remoteCcuBuffer.localDoorbellTokenValue;
             route.localSqDepth = remoteCcuBuffer.localSqDepth;
+            route.startJettyId = remoteCcuBuffer.startJettyId;
             route.endpointRouteVerified = true;
+            if (route.startJettyId != 0) {
+                verifiedStartJettyId = verifiedStartJettyId == 0 ?
+                    route.startJettyId :
+                    std::min<uint16_t>(verifiedStartJettyId, route.startJettyId);
+                verifiedJettyEnd = std::max<uint32_t>(
+                    verifiedJettyEnd,
+                    static_cast<uint32_t>(route.startJettyId) + 1U);
+            }
         }
         result.routes.push_back(route);
     }
+    if (verifiedStartJettyId != 0) {
+        result.startJettyId = verifiedStartJettyId;
+        const uint32_t requiredJettyCount = verifiedJettyEnd - verifiedStartJettyId;
+        result.pfeJettyCount = CheckedU16(std::max<uint32_t>(result.pfeJettyCount, requiredJettyCount));
+    }
 
     *snapshot = result;
     FillTemplateReport(*snapshot, report);
diff --git a/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.h b/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.h
index 26ff8c7f..c37f908a 100644
--- a/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.h
+++ b/src/comm/ccu/tilexr_ccu_lower_layer_plan_builder.h
@@ -106,6 +106,9 @@ struct TileXRCcuRemoteCcuBufferInfo {
     uint32_t doorbellTokenId = 0;
     uint32_t doorbellTokenValue = 0;
     uint32_t sqDepth = 0;
+    uint16_t startJettyId = 0;
+    uint32_t localPfeId = 0;
+    bool localPfeIdValid = false;
     uint64_t localDoorbellVa = 0;
     uint32_t localDoorbellTokenId = 0;
     uint32_t localDoorbellTokenValue = 0;
diff --git a/src/comm/ccu/tilexr_ccu_resource_allocator.cpp b/src/comm/ccu/tilexr_ccu_resource_allocator.cpp
index e9bb407b..282d9fe2 100644
--- a/src/comm/ccu/tilexr_ccu_resource_allocator.cpp
+++ b/src/comm/ccu/tilexr_ccu_resource_allocator.cpp
@@ -370,6 +370,34 @@ int TileXRCcuResourceAllocator::Release(uint64_t receiptId)
     if (it == active_.end()) {
         return TILEXR_ERROR_PARA_CHECK_FAIL;
     }
+    const auto latest = active_.rbegin();
+    if (latest == active_.rend() || latest->first != receiptId) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    const ActiveAllocation& active = it->second;
+    if (mission_.used < active.missionUsed ||
+        repository_.used < active.repositoryUsed ||
+        xn_.used < active.localXnUsed ||
+        gsa_.used < active.localGsaUsed ||
+        localWaitCke_.used < static_cast<uint16_t>(active.localWaitCkeUsed + active.sourceCkeUsed) ||
+        remoteNotifyCke_.used < active.remoteNotifyCkeUsed ||
+        channel_.used < active.channelUsed ||
+        (remoteXn_.count != 0 && remoteXn_.used < active.remoteXnUsed)) {
+        return TILEXR_ERROR_INTERNAL;
+    }
+    mission_.used = static_cast<uint16_t>(mission_.used - active.missionUsed);
+    repository_.used = static_cast<uint16_t>(repository_.used - active.repositoryUsed);
+    xn_.used = static_cast<uint16_t>(xn_.used - active.localXnUsed);
+    if (remoteXn_.count == 0) {
+        xn_.used = static_cast<uint16_t>(xn_.used - active.remoteXnUsed);
+    } else {
+        remoteXn_.used = static_cast<uint16_t>(remoteXn_.used - active.remoteXnUsed);
+    }
+    gsa_.used = static_cast<uint16_t>(gsa_.used - active.localGsaUsed);
+    localWaitCke_.used = static_cast<uint16_t>(
+        localWaitCke_.used - active.localWaitCkeUsed - active.sourceCkeUsed);
+    remoteNotifyCke_.used = static_cast<uint16_t>(remoteNotifyCke_.used - active.remoteNotifyCkeUsed);
+    channel_.used = static_cast<uint16_t>(channel_.used - active.channelUsed);
     active_.erase(it);
     return TILEXR_SUCCESS;
 }
diff --git a/src/comm/ccu/tilexr_ccu_runtime_session.cpp b/src/comm/ccu/tilexr_ccu_runtime_session.cpp
index d6c66746..a9001a20 100644
--- a/src/comm/ccu/tilexr_ccu_runtime_session.cpp
+++ b/src/comm/ccu/tilexr_ccu_runtime_session.cpp
@@ -30,19 +30,23 @@ struct TileXRThreadAllGatherState {
 static map<string, TileXRThreadAllGatherState> g_directCcuAllGatherStates;
 static std::mutex g_mtx;
 static std::mutex g_ccuDirectRuntimeMtx;
-static bool g_ccuDirectRuntimeUnavailable = false;
-static std::string g_ccuDirectRuntimeUnavailableMessage;
+static std::map<int, std::string> g_ccuDirectRuntimeUnavailableByDevice;
 
-std::string TileXRCcuRuntimeSession::ProcessDirectCcuRuntimeUnavailableMessage()
+std::string TileXRCcuRuntimeSession::DirectCcuRuntimeUnavailableMessageForDevice(int devId)
 {
     lock_guard<mutex> lock(g_ccuDirectRuntimeMtx);
-    if (!g_ccuDirectRuntimeUnavailable) {
+    const auto it = g_ccuDirectRuntimeUnavailableByDevice.find(devId);
+    if (it == g_ccuDirectRuntimeUnavailableByDevice.end()) {
         return {};
     }
-    return g_ccuDirectRuntimeUnavailableMessage.empty() ?
-        "direct CCU runtime unavailable after process-level init failure" :
-        "direct CCU runtime unavailable after process-level init failure: " +
-            g_ccuDirectRuntimeUnavailableMessage;
+    return it->second.empty() ?
+        "direct CCU runtime unavailable for device " + std::to_string(devId) :
+        "direct CCU runtime unavailable for device " + std::to_string(devId) + ": " + it->second;
+}
+
+std::string TileXRCcuRuntimeSession::DirectCcuRuntimeUnavailableMessage() const
+{
+    return DirectCcuRuntimeUnavailableMessageForDevice(devId_);
 }
 
 void TileXRCcuRuntimeSession::Shutdown()
@@ -73,6 +77,11 @@ int TileXRCcuRuntimeSession::RankSize() const
     return rankSize_;
 }
 
+int TileXRCcuRuntimeSession::DevId() const
+{
+    return devId_;
+}
+
 int TileXRCcuRuntimeSession::Init(const TileXRCcuBackendOptions &options)
 {
     Shutdown();
@@ -88,8 +97,9 @@ int TileXRCcuRuntimeSession::Init(const TileXRCcuBackendOptions &options)
     }
 
     lock_guard<mutex> lock(g_ccuDirectRuntimeMtx);
-    if (g_ccuDirectRuntimeUnavailable) {
-        TILEXR_LOG(INFO) << "direct CCU runtime skipped after previous init failure";
+    const auto unavailable = g_ccuDirectRuntimeUnavailableByDevice.find(devId_);
+    if (unavailable != g_ccuDirectRuntimeUnavailableByDevice.end()) {
+        TILEXR_LOG(INFO) << "direct CCU runtime skipped after previous init failure on device " << devId_;
         return TILEXR_SUCCESS;
     }
 
@@ -115,8 +125,7 @@ int TileXRCcuRuntimeSession::Init(const TileXRCcuBackendOptions &options)
                          << ", raInitialized " << (runtimeReport.raInitialized ? 1 : 0)
                          << ", ccuTlvInitialized " << (runtimeReport.ccuTlvInitialized ? 1 : 0)
                          << ", " << runtimeReport.message << ", direct CCU disabled";
-        g_ccuDirectRuntimeUnavailable = true;
-        g_ccuDirectRuntimeUnavailableMessage = runtimeReport.message;
+        g_ccuDirectRuntimeUnavailableByDevice[devId_] = runtimeReport.message;
         ResetDirectCcuBasicInfo();
         ccuDirectRuntime_.reset();
         return TILEXR_SUCCESS;
@@ -199,6 +208,27 @@ int TileXRCcuRuntimeSession::RegisterCcuResourceRmaBuffer(uint64_t resourceAddr)
     return ccuDirectRuntime_->RegisterCcuResourceRmaBuffer(resourceAddr);
 }
 
+int TileXRCcuRuntimeSession::RegisterMemoryBuffer(
+    uint64_t addr,
+    uint64_t bytes,
+    TileXRCcuRegisteredMemoryBufferInfo *info)
+{
+    if (ccuDirectRuntime_ == nullptr || !ccuDirectRuntime_->IsAvailable()) {
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+    return ccuDirectRuntime_->RegisterMemoryBuffer(addr, bytes, info);
+}
+
+int TileXRCcuRuntimeSession::ImportRemoteMemoryBuffer(
+    const TileXRCcuRemoteMemoryBufferImportRequest &request,
+    TileXRCcuImportedRemoteMemoryBufferInfo *info)
+{
+    if (ccuDirectRuntime_ == nullptr || !ccuDirectRuntime_->IsAvailable()) {
+        return TILEXR_ERROR_NOT_FOUND;
+    }
+    return ccuDirectRuntime_->ImportRemoteMemoryBuffer(request, info);
+}
+
 int TileXRCcuRuntimeSession::ExportLocalCcuRmaBuffer(TileXRCcuLocalResourceWindowInfo *info)
 {
     if (ccuDirectRuntime_ == nullptr || !ccuDirectRuntime_->IsAvailable()) {
@@ -315,13 +345,13 @@ int TileXRCcuRuntimeSession::DirectCcuThreadAllGather(const void *sendBuf, size_
                 return TILEXR_SUCCESS;
             }
         }
-        const std::string processUnavailableMessage = ProcessDirectCcuRuntimeUnavailableMessage();
-        if (!processUnavailableMessage.empty()) {
+        const std::string unavailableMessage = DirectCcuRuntimeUnavailableMessage();
+        if (!unavailableMessage.empty()) {
             lock_guard<mutex> lock(g_mtx);
             g_directCcuAllGatherStates.erase(key);
             TILEXR_LOG(ERROR) << "direct CCU thread allgather abort rank " << rank_ << "/" << rankSize_
                               << " uid " << uid_ << " round " << round << ", "
-                              << processUnavailableMessage;
+                              << unavailableMessage;
             return TILEXR_ERROR_NOT_FOUND;
         }
         this_thread::sleep_for(1ms);
diff --git a/src/comm/ccu/tilexr_ccu_runtime_session.h b/src/comm/ccu/tilexr_ccu_runtime_session.h
index d0010864..cde0802c 100644
--- a/src/comm/ccu/tilexr_ccu_runtime_session.h
+++ b/src/comm/ccu/tilexr_ccu_runtime_session.h
@@ -23,6 +23,7 @@ class TileXRCcuRuntimeSession {
 
     int Rank() const;
     int RankSize() const;
+    int DevId() const;
     int AllGather(const void *sendBuf, size_t sendBytes, void *recvBuf);
 
     int RefreshDirectCcuBasicInfo(uint8_t dieId = 0);
@@ -32,6 +33,10 @@ class TileXRCcuRuntimeSession {
     const TileXRCcuDriverAdapterReport &GetDirectCcuBasicInfoReport() const;
 
     int RegisterCcuResourceRmaBuffer(uint64_t resourceAddr);
+    int RegisterMemoryBuffer(uint64_t addr, uint64_t bytes, TileXRCcuRegisteredMemoryBufferInfo *info);
+    int ImportRemoteMemoryBuffer(
+        const TileXRCcuRemoteMemoryBufferImportRequest &request,
+        TileXRCcuImportedRemoteMemoryBufferInfo *info);
     int ExportLocalCcuRmaBuffer(TileXRCcuLocalResourceWindowInfo *info);
     int ExportRemoteCcuRmaBuffers(std::vector<TileXRCcuRemoteCcuBufferInfo> *buffers);
     int ExportLowerLayerTransportSnapshot(
@@ -41,10 +46,11 @@ class TileXRCcuRuntimeSession {
     int RefreshLocalVerifiedEndpointRoute(TileXRCcuDirectRuntimeReport *report);
     int CreateDriverAdapter(TileXRCcuDriverAdapter *adapter, TileXRCcuDriverAdapterReport *report);
 
-    static std::string ProcessDirectCcuRuntimeUnavailableMessage();
+    std::string DirectCcuRuntimeUnavailableMessage() const;
 
 private:
     void ResetDirectCcuBasicInfo();
+    static std::string DirectCcuRuntimeUnavailableMessageForDevice(int devId);
     static int DirectCcuAllGatherCallback(const void *sendBuf, size_t sendBytes, void *recvBuf, void *userData);
     int DirectCcuThreadAllGather(const void *sendBuf, size_t sendBytes, void *recvBuf);
 
diff --git a/src/comm/ccu/tilexr_ccu_signal_wait_program.cpp b/src/comm/ccu/tilexr_ccu_signal_wait_program.cpp
new file mode 100644
index 00000000..fb00e38a
--- /dev/null
+++ b/src/comm/ccu/tilexr_ccu_signal_wait_program.cpp
@@ -0,0 +1,221 @@
+/*
+ * Copyright (c) 2026 TileXR Project
+ * Licensed under the Apache License, Version 2.0
+ */
+
+#include "ccu/tilexr_ccu_signal_wait_program.h"
+
+namespace TileXR {
+namespace {
+
+void ResetReport(TileXRCcuBarrierProgramReport* report)
+{
+    if (report != nullptr) {
+        *report = TileXRCcuBarrierProgramReport {};
+    }
+}
+
+int Fail(
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report,
+    const std::string& message)
+{
+    if (program != nullptr) {
+        program->clear();
+    }
+    if (report != nullptr) {
+        report->message = message;
+    }
+    return TILEXR_ERROR_PARA_CHECK_FAIL;
+}
+
+bool HasSignalResources(const TileXRCcuSignalWaitProgramSpec& spec)
+{
+    return spec.channelId != 0 &&
+        spec.remoteNotifyCke != 0 &&
+        spec.remoteNotifyMask != 0 &&
+        spec.sourceCke != 0 &&
+        spec.sourceCkeMask != 0;
+}
+
+bool HasWaitResources(const TileXRCcuSignalWaitProgramSpec& spec)
+{
+    return spec.localWaitCke != 0 &&
+        spec.localWaitMask != 0;
+}
+
+TileXRCcuBarrierSyncSpec ToBarrierSpec(const TileXRCcuSignalWaitProgramSpec& spec)
+{
+    TileXRCcuBarrierSyncSpec barrier;
+    barrier.channelId = spec.channelId;
+    barrier.remoteXn = spec.remoteXn;
+    barrier.localXn = spec.localXn;
+    barrier.remoteNotifyCke = spec.remoteNotifyCke;
+    barrier.remoteNotifyMask = spec.remoteNotifyMask;
+    barrier.localWaitCke = spec.localWaitCke;
+    barrier.localWaitMask = spec.localWaitMask;
+    barrier.sourceCke = spec.sourceCke;
+    barrier.sourceCkeMask = spec.sourceCkeMask;
+    barrier.clearLocalWait = spec.clearLocalWait;
+    return barrier;
+}
+
+int AppendWaitInstruction(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report,
+    bool setWait)
+{
+    if (!HasWaitResources(spec)) {
+        return Fail(program, report, "missing wait CKE resource for direct CCU signal/wait program");
+    }
+    TileXRCcuCkeSpec wait;
+    wait.waitCkeId = spec.localWaitCke;
+    wait.waitMask = spec.localWaitMask;
+    wait.clearWait = spec.clearLocalWait;
+    TileXRCcuInstr instr;
+    const int ret = setWait ?
+        TileXRCcuEncodeSetCke(wait, &instr) :
+        TileXRCcuEncodeClearCke(wait, &instr);
+    if (ret != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode wait instruction for direct CCU signal/wait program");
+    }
+    program->push_back(instr);
+    return TILEXR_SUCCESS;
+}
+
+int AppendCommonPrelude(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report)
+{
+    if (spec.localXn == 0 || spec.sourceCke == 0 || spec.sourceCkeMask == 0) {
+        return Fail(program, report, "missing reserve resource for direct CCU signal/wait program");
+    }
+    TileXRCcuInstr instr;
+    if (TileXRCcuEncodeLoadImdToXn(spec.localXn, 0, 0, &instr) != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode signal/wait reserve XN prelude");
+    }
+    program->push_back(instr);
+
+    if (spec.localGsa != 0) {
+        if (TileXRCcuEncodeLoadImdToGsa(spec.localGsa, 0, &instr) != TILEXR_SUCCESS) {
+            return Fail(program, report, "failed to encode signal/wait reserve GSA prelude");
+        }
+    } else if (TileXRCcuEncodeLoadImdToXn(spec.localXn, 0, 0, &instr) != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode signal/wait reserve address prelude");
+    }
+    program->push_back(instr);
+
+    TileXRCcuCkeSpec init;
+    init.ckeId = spec.sourceCke;
+    init.mask = spec.sourceCkeMask;
+    init.clearWait = true;
+    if (TileXRCcuEncodeSetCke(init, &instr) != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode signal/wait reserve CKE prelude");
+    }
+    program->push_back(instr);
+    return TILEXR_SUCCESS;
+}
+
+int AppendFinish(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report)
+{
+    TileXRCcuInstr instr;
+    if (TileXRCcuEncodeLoadImdToXn(spec.localXn, 0, 0, &instr) != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode signal/wait finish instruction");
+    }
+    program->push_back(instr);
+    return TILEXR_SUCCESS;
+}
+
+int AppendSignalInstruction(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report)
+{
+    if (!HasSignalResources(spec)) {
+        return Fail(program, report, "missing signal resource for direct CCU signal/wait program");
+    }
+    TileXRCcuSyncCkeSpec post;
+    post.remoteCke = spec.remoteNotifyCke;
+    post.localCke = spec.sourceCke;
+    post.localCkeMask = spec.sourceCkeMask;
+    post.channelId = spec.channelId;
+    post.clearWait = spec.clearLocalWait;
+    TileXRCcuInstr instr;
+    if (TileXRCcuEncodeSyncCke(post, &instr) != TILEXR_SUCCESS) {
+        return Fail(program, report, "failed to encode signal instruction for direct CCU signal/wait program");
+    }
+    program->push_back(instr);
+    return TILEXR_SUCCESS;
+}
+
+int BuildWaitOnly(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report)
+{
+    program->clear();
+    const int ret = AppendWaitInstruction(spec, program, report, false);
+    if (ret != TILEXR_SUCCESS) {
+        return ret;
+    }
+    if (report != nullptr) {
+        report->postInstructionCount = 0;
+        report->waitInstructionCount = 1;
+        report->totalInstructionCount = static_cast<uint32_t>(program->size());
+        report->message = "ok";
+    }
+    return TILEXR_SUCCESS;
+}
+
+} // namespace
+
+int TileXRCcuBuildSignalWaitProgram(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report)
+{
+    ResetReport(report);
+    if (program == nullptr) {
+        return Fail(program, report, "missing output direct CCU signal/wait program");
+    }
+    program->clear();
+
+    if (spec.role == TileXRCcuSignalWaitProgramRole::Wait) {
+        return BuildWaitOnly(spec, program, report);
+    }
+
+    if (!HasSignalResources(spec)) {
+        return Fail(program, report, "missing signal resource for direct CCU signal/wait program");
+    }
+    if (spec.role == TileXRCcuSignalWaitProgramRole::SignalAndWait && !HasWaitResources(spec)) {
+        return Fail(program, report, "missing wait CKE resource for direct CCU signal/wait program");
+    }
+
+    if (AppendCommonPrelude(spec, program, report) != TILEXR_SUCCESS) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    if (AppendSignalInstruction(spec, program, report) != TILEXR_SUCCESS) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    if (spec.role == TileXRCcuSignalWaitProgramRole::SignalAndWait &&
+        AppendWaitInstruction(spec, program, report, false) != TILEXR_SUCCESS) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    if (AppendFinish(spec, program, report) != TILEXR_SUCCESS) {
+        return TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    if (report != nullptr) {
+        report->postInstructionCount = 1;
+        report->waitInstructionCount = spec.role == TileXRCcuSignalWaitProgramRole::SignalAndWait ? 1U : 0U;
+        report->totalInstructionCount = static_cast<uint32_t>(program->size());
+        report->message = "ok";
+    }
+    return TILEXR_SUCCESS;
+}
+
+} // namespace TileXR
diff --git a/src/comm/ccu/tilexr_ccu_signal_wait_program.h b/src/comm/ccu/tilexr_ccu_signal_wait_program.h
new file mode 100644
index 00000000..63cbb7e8
--- /dev/null
+++ b/src/comm/ccu/tilexr_ccu_signal_wait_program.h
@@ -0,0 +1,41 @@
+/*
+ * Copyright (c) 2026 TileXR Project
+ * Licensed under the Apache License, Version 2.0
+ */
+
+#ifndef TILEXR_CCU_SIGNAL_WAIT_PROGRAM_H
+#define TILEXR_CCU_SIGNAL_WAIT_PROGRAM_H
+
+#include "ccu/tilexr_ccu_barrier_program.h"
+
+namespace TileXR {
+
+enum class TileXRCcuSignalWaitProgramRole {
+    Signal = 0,
+    Wait = 1,
+    SignalAndWait = 2,
+};
+
+struct TileXRCcuSignalWaitProgramSpec {
+    TileXRCcuSignalWaitProgramRole role = TileXRCcuSignalWaitProgramRole::Signal;
+    uint16_t channelId = 0;
+    uint16_t remoteXn = 0;
+    uint16_t localXn = 0;
+    uint16_t localGsa = 0;
+    uint16_t remoteNotifyCke = 0;
+    uint16_t remoteNotifyMask = 0;
+    uint16_t localWaitCke = 0;
+    uint16_t localWaitMask = 0;
+    uint16_t sourceCke = 0;
+    uint16_t sourceCkeMask = 0;
+    bool clearLocalWait = true;
+};
+
+int TileXRCcuBuildSignalWaitProgram(
+    const TileXRCcuSignalWaitProgramSpec& spec,
+    std::vector<TileXRCcuInstr>* program,
+    TileXRCcuBarrierProgramReport* report);
+
+} // namespace TileXR
+
+#endif // TILEXR_CCU_SIGNAL_WAIT_PROGRAM_H
diff --git a/src/comm/tilexr_comm.cpp b/src/comm/tilexr_comm.cpp
index a2ede721..cccb7c5d 100644
--- a/src/comm/tilexr_comm.cpp
+++ b/src/comm/tilexr_comm.cpp
@@ -55,6 +55,17 @@ static std::mutex g_mtx;
 static std::mutex g_sdmaMtx;
 static bool g_sdmaUnavailable = false;
 
+constexpr const char* TILEXR_ENABLE_CCU_BACKEND_ENV = "TILEXR_ENABLE_CCU_BACKEND";
+
+bool ShouldEnableCcuBackend()
+{
+    const char *value = std::getenv(TILEXR_ENABLE_CCU_BACKEND_ENV);
+    if (value == nullptr || value[0] == '\0') {
+        return false;
+    }
+    const string text(value);
+    return text == "1" || text == "true" || text == "TRUE" || text == "yes" || text == "on";
+}
 
 // 如果是互联的链路，返回false； 对910B2C那些不互联的链路，返回true
 bool SkipUnusedChannel910B2C(int curRank, int peerRank, ChipName chipName)
@@ -210,6 +221,26 @@ int TileXRComm::InitCcuBackend()
     return ccuBackend_->Init(options);
 }
 
+int TileXRComm::InitCcuBackendIfEnabled()
+{
+    if (!ShouldEnableCcuBackend()) {
+        return TILEXR_SUCCESS;
+    }
+
+    const int ccuRet = InitCcuBackend();
+    if (ccuRet != TILEXR_SUCCESS || ccuBackend_ == nullptr || !ccuBackend_->Available()) {
+        TILEXR_LOG(WARN) << "TileXR CCU backend init failed, direct CCU disabled, ret: " << ccuRet;
+        if (ccuBackend_ != nullptr) {
+            ccuBackend_->Shutdown();
+            ccuBackend_.reset();
+        }
+        return TILEXR_SUCCESS;
+    }
+
+    TILEXR_LOG(INFO) << "TileXR CCU backend initialized";
+    return TILEXR_SUCCESS;
+}
+
 TileXRCcuBackend *TileXRComm::GetCcuBackendForCollectives()
 {
     return ccuBackend_.get();
@@ -505,6 +536,10 @@ int TileXRComm::Init()
     if (ret != TILEXR_SUCCESS) {
         return ret;
     }
+    ret = InitCcuBackendIfEnabled();
+    if (ret != TILEXR_SUCCESS) {
+        return ret;
+    }
 
     // set comm args in device.
     ret = SyncCommArgs();
@@ -577,6 +612,10 @@ int TileXRComm::InitThread(const std::string &uid)
     if (ret != TILEXR_SUCCESS) {
         return ret;
     }
+    ret = InitCcuBackendIfEnabled();
+    if (ret != TILEXR_SUCCESS) {
+        return ret;
+    }
     ret = SyncCommArgs();
     if (ret != TILEXR_SUCCESS) {
         TILEXR_LOG(ERROR) << "SyncCommArgs failed! ret: " << ret;
diff --git a/src/comm/tilexr_comm.h b/src/comm/tilexr_comm.h
index bfde341d..48dda0c3 100644
--- a/src/comm/tilexr_comm.h
+++ b/src/comm/tilexr_comm.h
@@ -83,6 +83,7 @@ class TileXRComm {
     int InitUDMA();
     int ApplyUDMACommArgsState(const TileXRUDMACommArgsState &state);
     static int ApplyUDMACommArgsStateCallback(const TileXRUDMACommArgsState &state, void *userData);
+    int InitCcuBackendIfEnabled();
     int InitSDMA();
     int UpdateCommArgsDev();
     void ResetSDMAState();
diff --git a/tests/ccu/ccu_tilexr_direct_smoke_probe.cpp b/tests/ccu/ccu_tilexr_direct_smoke_probe.cpp
index 0e1f899b..979798fc 100644
--- a/tests/ccu/ccu_tilexr_direct_smoke_probe.cpp
+++ b/tests/ccu/ccu_tilexr_direct_smoke_probe.cpp
@@ -11,7 +11,9 @@
 #include "acl/acl_rt.h"
 #include "tilexr_api.h"
 #include "tilexr_types.h"
+#include "ccu/tilexr_ccu_backend.h"
 #include "ccu/tilexr_ccu_collective_planner.h"
+#include "ccu/tilexr_ccu_driver_adapter.h"
 #include "ccu/tilexr_ccu_executor.h"
 #include "ccu/tilexr_ccu_runtime_session.h"
 #include "tools/socket/tilexr_sock_exchange.h"
@@ -39,6 +41,7 @@ struct DirectCcuSmokeContext {
     TileXR::TileXRCcuRuntimeSession session;
     TileXR::TileXRCcuCollectivePlanner planner;
     TileXR::TileXRCcuExecutor executor;
+    TileXR::TileXRCcuBackend backend;
 };
 
 using TileXRDirectCcuPrepareOptions = TileXR::TileXRCcuDirectInstallOptions;
@@ -71,6 +74,12 @@ constexpr const char* kPreSubmitDelayMsEnv = "TILEXR_CCU_DIRECT_SMOKE_PRE_SUBMIT
 constexpr const char* kP2pCcuCopyEnv = "TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY";
 constexpr const char* kExpectP2pCcuCopyEnv = "TILEXR_CCU_DIRECT_SMOKE_EXPECT_P2P_CCU_COPY";
 constexpr const char* kP2pCcuCopyBytesEnv = "TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_BYTES";
+constexpr const char* kP2pCcuCopyActiveRankEnv = "TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_ACTIVE_RANK";
+constexpr const char* kP2pCcuCopyDirectionEnv = "TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_DIRECTION";
+constexpr const char* kP2pCcuCopyResourceWindowEnv = "TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_RESOURCE_WINDOW";
+constexpr const char* kSignalWaitEnv = "TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT";
+constexpr const char* kSignalWaitSignalRankEnv = "TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK";
+constexpr const char* kSignalWaitBarrierEnv = "TILEXR_CCU_DIRECT_SMOKE_BARRIER";
 constexpr const char* kLocalWaitCkeStartEnv = "TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_START";
 constexpr const char* kLocalWaitCkeCountEnv = "TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_COUNT";
 constexpr const char* kRemoteNotifyCkeStartEnv = "TILEXR_CCU_PROBE_REMOTE_NOTIFY_CKE_START";
@@ -128,6 +137,16 @@ bool EnvFlag(const char* name)
     return text == "1" || text == "true" || text == "TRUE" || text == "yes" || text == "on";
 }
 
+bool SignalWaitSmokeEnabled()
+{
+    return EnvFlag(kSignalWaitEnv);
+}
+
+bool BarrierSmokeEnabled()
+{
+    return EnvFlag(kSignalWaitBarrierEnv);
+}
+
 bool ShouldFastExitAfterPrepareFailure(int ret)
 {
     return ret != 0 && EnvFlag(kFastExitOnPrepareFailureEnv);
@@ -182,6 +201,30 @@ int EnvInt(const char* name, int fallback)
     return ParseInt(std::getenv(name), fallback);
 }
 
+bool IsP2pCcuCopyActiveRank(int rank)
+{
+    return rank == EnvInt(kP2pCcuCopyActiveRankEnv, 0);
+}
+
+TileXR::TileXRCcuMemoryCopyDirection P2pCcuCopyDirectionFromEnv()
+{
+    const char* value = std::getenv(kP2pCcuCopyDirectionEnv);
+    if (value != nullptr && std::string(value) == "local_to_remote") {
+        return TileXR::TileXRCcuMemoryCopyDirection::LocalToRemote;
+    }
+    return TileXR::TileXRCcuMemoryCopyDirection::RemoteToLocal;
+}
+
+bool ShouldCheckInactiveP2pCcuCopyRank()
+{
+    return P2pCcuCopyDirectionFromEnv() == TileXR::TileXRCcuMemoryCopyDirection::LocalToRemote;
+}
+
+bool ShouldCheckActiveP2pCcuCopyRank()
+{
+    return P2pCcuCopyDirectionFromEnv() == TileXR::TileXRCcuMemoryCopyDirection::RemoteToLocal;
+}
+
 uint64_t ParseU64(const char* value, uint64_t fallback)
 {
     if (value == nullptr || value[0] == '\0') {
@@ -257,6 +300,32 @@ bool SyncXnLoadPostOnlyBarrierMode()
     return value != nullptr && std::string(value) == "sync_xn_load_post_only";
 }
 
+TileXR::TileXRCcuBarrierMode BarrierModeFromEnv()
+{
+    if (SyncCkeBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::SyncCke;
+    }
+    if (SyncCkeSetWaitBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::SyncCkeSetWait;
+    }
+    if (SyncCkePostOnlyBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::SyncCkePostOnly;
+    }
+    if (LocalCkeBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::LocalCke;
+    }
+    if (LocalCkePostOnlyBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::LocalCkePostOnly;
+    }
+    if (SyncXnPostOnlyBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::SyncXnPostOnly;
+    }
+    if (SyncXnLoadPostOnlyBarrierMode()) {
+        return TileXR::TileXRCcuBarrierMode::SyncXnLoadPostOnly;
+    }
+    return TileXR::TileXRCcuBarrierMode::SyncXn;
+}
+
 TileXR::TileXRCcuRepositoryInstallWindow RepositoryInstallWindowFromEnv()
 {
     const char* value = std::getenv(kRepositoryInstallWindowEnv);
@@ -476,6 +545,17 @@ void PrintP2pCcuCopy(
               << std::endl;
 }
 
+void PrintP2pCcuCopySkipped(int rank, int peer, const P2pCcuCopyState& state)
+{
+    std::cout << "tilexr_ccu_direct_smoke p2pCcuCopy skipped"
+              << " rank=" << rank
+              << " peer=" << peer
+              << " bytes=" << state.bytes
+              << " activeRank=" << EnvInt(kP2pCcuCopyActiveRankEnv, 0)
+              << " reason=\"inactive p2p CCU-copy rank\""
+              << std::endl;
+}
+
 int RunP2pCcuCopy(
     int rank,
     int peer,
@@ -487,8 +567,14 @@ int RunP2pCcuCopy(
     if (state == nullptr) {
         return TileXR::TILEXR_ERROR_PARA_CHECK_FAIL;
     }
-    if (syncRet == ACL_SUCCESS && submitRet == TileXR::TILEXR_SUCCESS) {
-        (void)CheckP2pCcuCopyState(state);
+    if (ShouldCheckActiveP2pCcuCopyRank() && syncRet == ACL_SUCCESS && submitRet == TileXR::TILEXR_SUCCESS) {
+        if (EnvFlag(kP2pCcuCopyResourceWindowEnv)) {
+            state->passed = true;
+        } else {
+            (void)CheckP2pCcuCopyState(state);
+        }
+    } else if (!ShouldCheckActiveP2pCcuCopyRank() && syncRet == ACL_SUCCESS && submitRet == TileXR::TILEXR_SUCCESS) {
+        state->passed = true;
     }
     PrintP2pCcuCopy(rank, peer, *state, prepareRet, submitRet, syncRet);
     return state->passed ? TileXR::TILEXR_SUCCESS : TileXR::TILEXR_ERROR_INTERNAL;
@@ -553,6 +639,16 @@ int InitCommForDirectCcuSmoke(int commDomain, int rankSize, int rank, int device
                   << " directOnly=" << (EnvFlag(kDirectCcuOnlyInitEnv) ? 1 : 0)
                   << std::endl;
     }
+    if (ret == TileXR::TILEXR_SUCCESS && (SignalWaitSmokeEnabled() || BarrierSmokeEnabled())) {
+        const int backendRet = context->backend.Init(options);
+        std::cout << "tilexr_ccu_signal_wait backendInit"
+                  << " rank=" << rank
+                  << " rankSize=" << rankSize
+                  << " device=" << device
+                  << " ret=" << backendRet
+                  << std::endl;
+        return backendRet;
+    }
     return ret;
 }
 
@@ -591,12 +687,100 @@ TileXRDirectCcuPrepareOptions MakePrepareOptions(int rank, int rankSize, int dev
     options.repositoryInstallOptions.dataLenMode = RepositoryInstallDataLenModeFromEnv();
     options.repositoryMemoryAllocMode = RepositoryMemoryAllocModeFromEnv();
     options.installOrder = InstallOrderFromEnv();
+    options.barrierMode = BarrierModeFromEnv();
+    options.taskTimeout = static_cast<uint16_t>(EnvInt("TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT", 0));
     options.deviceId = static_cast<uint32_t>(device);
     options.rank = static_cast<uint32_t>(rank);
     options.provider = rankSize > 0 ? "tilexr-direct-smoke-probe" : "";
     return options;
 }
 
+const char* SignalWaitRoleName(TileXR::TileXRCcuSignalWaitRole role)
+{
+    switch (role) {
+        case TileXR::TileXRCcuSignalWaitRole::Signal:
+            return "signal";
+        case TileXR::TileXRCcuSignalWaitRole::Wait:
+            return "wait";
+        case TileXR::TileXRCcuSignalWaitRole::SignalAndWait:
+            return "signal_and_wait";
+        default:
+            return "unknown";
+    }
+}
+
+TileXR::TileXRCcuSignalWaitRole SignalWaitRoleForRank(int rank)
+{
+    if (BarrierSmokeEnabled()) {
+        return TileXR::TileXRCcuSignalWaitRole::SignalAndWait;
+    }
+    const int signalRank = EnvInt(kSignalWaitSignalRankEnv, 0);
+    return rank == signalRank ? TileXR::TileXRCcuSignalWaitRole::Signal : TileXR::TileXRCcuSignalWaitRole::Wait;
+}
+
+uint32_t DefaultSignalWaitInstructionCount(TileXR::TileXRCcuSignalWaitRole role)
+{
+    switch (role) {
+        case TileXR::TileXRCcuSignalWaitRole::Signal:
+            return 5U;
+        case TileXR::TileXRCcuSignalWaitRole::Wait:
+            return 5U;
+        case TileXR::TileXRCcuSignalWaitRole::SignalAndWait:
+            return 6U;
+        default:
+            return 0U;
+    }
+}
+
+TileXR::TileXRCcuSignalWaitRequest MakeSignalWaitRequest(
+    int rank,
+    int rankSize,
+    const TileXRDirectCcuPrepareOptions& options)
+{
+    TileXR::TileXRCcuSignalWaitRequest request {};
+    request.peerRank = rankSize == 2 ? 1 - rank : (rank + 1) % rankSize;
+    request.role = SignalWaitRoleForRank(rank);
+    request.syncInstructionCount = std::getenv("TILEXR_CCU_PROBE_SYNC_INSTRUCTION_COUNT") == nullptr ?
+        0U :
+        options.syncInstructionCount;
+    request.missionStartId = options.missionStartId;
+    request.instructionStartId = options.instructionStartId;
+    request.missionInstructionStartId = options.missionInstructionStartId;
+    request.xnStartId = options.xnStartId;
+    request.remoteXnStartId = options.remoteXnStartId;
+    request.remoteXnCount = options.remoteXnCount;
+    request.ckeStartId = options.ckeStartId;
+    request.channelStartId = options.channelStartId;
+    request.localWaitCkeStartId = options.localWaitCkeStartId;
+    request.localWaitCkeCount = options.localWaitCkeCount;
+    request.remoteNotifyCkeStartId = options.remoteNotifyCkeStartId;
+    request.remoteNotifyCkeCount = options.remoteNotifyCkeCount;
+    request.timeout = static_cast<uint16_t>(EnvInt("TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT", 20));
+    request.provider = "tilexr-direct-smoke-probe-signal-wait";
+    return request;
+}
+
+TileXRDirectCcuPrepareReport SignalWaitInstallReportFromPlan(
+    const TileXR::TileXRCcuSignalWaitPlan& plan)
+{
+    TileXRDirectCcuPrepareReport report {};
+    report.pipelineBuilt = !plan.attempt.plan.taskWindows.empty();
+    report.installAttempted = plan.attempt.installReport.installAttempted;
+    report.installSucceeded = plan.attempt.installReport.installSucceeded;
+    report.submitReady = plan.attempt.providerReport.submitReady;
+    report.requiredInstallSurfaceCount = plan.attempt.installReport.requiredInstallSurfaceCount;
+    report.publicVerifiedInstallSurfaceCount = plan.attempt.installReport.publicVerifiedInstallSurfaceCount;
+    report.missingInstallSurfaceCount = plan.attempt.installReport.missingInstallSurfaceCount;
+    report.taskCount = static_cast<uint32_t>(plan.attempt.plan.taskWindows.size());
+    report.submitTaskCount = static_cast<uint32_t>(plan.submitTasks.size());
+    if (!plan.attempt.providerReport.message.empty()) {
+        report.message = plan.attempt.providerReport.message;
+    } else {
+        report.message = plan.attempt.installReport.message;
+    }
+    return report;
+}
+
 void PrintInstallReport(
     const char* prefix,
     int ret,
@@ -831,6 +1015,79 @@ void PrintPreparedTasks(TileXRDirectCcuPreparedTasksPtr prepared, uint32_t taskC
     std::cout << std::endl;
 }
 
+uint16_t LoadLe16(const uint8_t* data, size_t index)
+{
+    return static_cast<uint16_t>(data[index * 2U]) |
+        static_cast<uint16_t>(static_cast<uint16_t>(data[index * 2U + 1U]) << 8U);
+}
+
+void PrintMissionContext(
+    DirectCcuSmokeContext* context,
+    const TileXRDirectCcuTaskInfo& task,
+    const char* label)
+{
+    if (context == nullptr || label == nullptr) {
+        return;
+    }
+    TileXR::TileXRCcuDriverAdapter adapter;
+    TileXR::TileXRCcuDriverAdapterReport adapterReport;
+    int ret = context->session.CreateDriverAdapter(&adapter, &adapterReport);
+    if (ret != TileXR::TILEXR_SUCCESS) {
+        std::cerr << label << " missionCtxRead ret=" << ret
+                  << " message=\"" << adapterReport.message << "\""
+                  << std::endl;
+        return;
+    }
+
+    uint8_t raw[TileXR::TILEXR_CCU_DATA_ARRAY_SLOT_BYTES] = {};
+    ret = adapter.ReadMissionContext(
+        task.dieId,
+        task.missionId,
+        raw,
+        sizeof(raw),
+        &adapterReport);
+    if (ret != TileXR::TILEXR_SUCCESS) {
+        std::cerr << label << " missionCtxRead ret=" << ret
+                  << " opcode=" << adapterReport.opcode
+                  << " driverRet=" << adapterReport.driverRet
+                  << " opRet=" << adapterReport.opRet
+                  << " message=\"" << adapterReport.message << "\""
+                  << std::endl;
+        return;
+    }
+
+    const uint16_t part2 = LoadLe16(raw, 2);
+    const uint16_t part3 = LoadLe16(raw, 3);
+    const uint16_t part4 = LoadLe16(raw, 4);
+    const uint16_t part5 = LoadLe16(raw, 5);
+    const uint16_t part6 = LoadLe16(raw, 6);
+    const uint16_t part7 = LoadLe16(raw, 7);
+    const uint16_t status = static_cast<uint16_t>(((part3 & 0x7U) << 13U) | ((part2 >> 3U) & 0x1fffU));
+    const uint16_t currentIns = static_cast<uint16_t>(((part5 & 0x1fU) << 11U) | ((part4 >> 5U) & 0x7ffU));
+    const uint16_t endIns = static_cast<uint16_t>(((part6 & 0x1fU) << 11U) | ((part5 >> 5U) & 0x7ffU));
+    const uint16_t startIns = static_cast<uint16_t>(((part7 & 0x1fU) << 11U) | ((part6 >> 5U) & 0x7ffU));
+    const uint16_t missionVld = static_cast<uint16_t>((part7 >> 6U) & 0x1U);
+
+    std::cerr << label << " missionCtx"
+              << " dieId=" << static_cast<uint32_t>(task.dieId)
+              << " missionId=" << static_cast<uint32_t>(task.missionId)
+              << " status=0x" << std::hex << status
+              << " currentIns=" << std::dec << currentIns
+              << " startIns=" << startIns
+              << " endIns=" << endIns
+              << " missionVld=" << missionVld
+              << " rawWords=";
+    for (size_t i = 0; i < sizeof(raw) / sizeof(uint64_t); ++i) {
+        uint64_t word = 0;
+        std::memcpy(&word, raw + i * sizeof(uint64_t), sizeof(word));
+        if (i != 0) {
+            std::cerr << ",";
+        }
+        std::cerr << "0x" << std::hex << word;
+    }
+    std::cerr << std::dec << std::endl;
+}
+
 void PrintInstructionReadback(DirectCcuSmokeContext* context, TileXRDirectCcuPreparedTasksPtr prepared, uint32_t taskCount)
 {
     if (!EnvFlag(kReadbackInstructionsEnv)) {
@@ -1077,11 +1334,170 @@ bool WaitForCollectiveSubmitDone(int rank, int rankSize, int localResult)
     }
 }
 
+bool WaitForInactiveP2pCcuCopyRank(int rank, int rankSize, int localResult)
+{
+    return WaitForCollectiveSubmitDone(rank, rankSize, localResult);
+}
+
+int RunInactiveP2pCcuCopyRank(int rank, int peer, int rankSize, P2pCcuCopyState* state, int localResult)
+{
+    if (state == nullptr) {
+        return TileXR::TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    PrintP2pCcuCopySkipped(rank, peer, *state);
+    if (!WaitForInactiveP2pCcuCopyRank(rank, rankSize, localResult) && localResult == 0) {
+        return 13;
+    }
+    if (localResult != 0 || !ShouldCheckInactiveP2pCcuCopyRank()) {
+        return localResult;
+    }
+    const int checkRet = CheckP2pCcuCopyState(state);
+    PrintP2pCcuCopy(rank, peer, *state, TileXR::TILEXR_SUCCESS, TileXR::TILEXR_SUCCESS, ACL_SUCCESS);
+    return checkRet == ACL_SUCCESS ? 0 : 14;
+}
+
+int RunSignalWaitSmokeForRank(DirectCcuSmokeContext* context, int rank, int rankSize, int device)
+{
+    if (context == nullptr) {
+        return TileXR::TILEXR_ERROR_PARA_CHECK_FAIL;
+    }
+    TileXRDirectCcuPrepareOptions options = MakePrepareOptions(rank, rankSize, device);
+    options.syncResourceCount = 1;
+    options.sqeArgCount = 0;
+    const TileXR::TileXRCcuSignalWaitRole role = SignalWaitRoleForRank(rank);
+    if (std::getenv("TILEXR_CCU_PROBE_SYNC_INSTRUCTION_COUNT") == nullptr) {
+        options.syncInstructionCount = DefaultSignalWaitInstructionCount(role);
+    }
+    PrintConfig(options, rankSize);
+
+    const TileXR::TileXRCcuSignalWaitRequest request = MakeSignalWaitRequest(rank, rankSize, options);
+    std::cout << "tilexr_ccu_signal_wait config"
+              << " rank=" << rank
+              << " peer=" << request.peerRank
+              << " role=" << SignalWaitRoleName(request.role)
+              << " barrier=" << (BarrierSmokeEnabled() ? 1 : 0)
+              << std::endl;
+
+    TileXR::TileXRCcuSignalWaitPlan plan;
+    const int prepareRet = context->backend.PrepareSignalWait(request, &plan);
+    TileXRDirectCcuPrepareReport installReport = SignalWaitInstallReportFromPlan(plan);
+    PrintInstallReport("tilexr_ccu_signal_wait prepare", prepareRet, installReport);
+    PrintPreparedTasks(&plan.attempt, installReport.submitTaskCount);
+
+    int finalRet = 0;
+    const bool submitRequested = EnvFlag(kSubmitEnv);
+    const bool collectiveSubmitReady = submitRequested ?
+        WaitForCollectiveSubmitReadiness(
+            rank,
+            rankSize,
+            prepareRet == TileXR::TILEXR_SUCCESS && installReport.submitReady && plan.ready) :
+        false;
+    if (prepareRet != TileXR::TILEXR_SUCCESS) {
+        finalRet = 6;
+    } else if (submitRequested && !collectiveSubmitReady && CollectiveSubmitReadyGateConfigured()) {
+        std::cout << "tilexr_ccu_signal_wait submit skipped reason=\"collective submitReady gate did not pass\""
+                  << " localSubmitReady=" << (installReport.submitReady ? 1 : 0)
+                  << std::endl;
+    } else if (submitRequested && (!installReport.submitReady || !plan.ready)) {
+        std::cout << "tilexr_ccu_signal_wait submit skipped reason=\"prepare did not reach submitReady\""
+                  << std::endl;
+    } else if (submitRequested) {
+        aclrtStream stream = nullptr;
+        int streamRet = aclrtCreateStream(&stream);
+        if (streamRet != ACL_SUCCESS) {
+            std::cerr << "tilexr_ccu_signal_wait aclrtCreateStream ret=" << streamRet << std::endl;
+            finalRet = 7;
+        } else {
+            const int delayRank = EnvInt(kDelayRankEnv, -1);
+            const int preSubmitDelayMs = EnvInt(kPreSubmitDelayMsEnv, 0);
+            const int effectiveDelayMs = rank == delayRank && preSubmitDelayMs > 0 ? preSubmitDelayMs : 0;
+            if (effectiveDelayMs > 0) {
+                std::cout << "tilexr_ccu_signal_wait preSubmitDelay"
+                          << " rank=" << rank
+                          << " delayMs=" << effectiveDelayMs
+                          << std::endl;
+                std::this_thread::sleep_for(std::chrono::milliseconds(effectiveDelayMs));
+            }
+
+            TileXRDirectCcuSubmitReport submitReport;
+            const auto submitBegin = std::chrono::steady_clock::now();
+            const int submitRet = context->backend.SubmitSignalWait(plan, stream, &submitReport);
+            const auto submitEnd = std::chrono::steady_clock::now();
+            PrintSubmitReport("tilexr_ccu_signal_wait submit", submitRet, submitReport);
+            const auto syncBegin = std::chrono::steady_clock::now();
+            TraceLifecycle("before signal/wait aclrtSynchronizeStream");
+            const int syncRet = aclrtSynchronizeStream(stream);
+            TraceLifecycle("after signal/wait aclrtSynchronizeStream");
+            const auto syncEnd = std::chrono::steady_clock::now();
+            std::cout << "tilexr_ccu_signal_wait timing"
+                      << " rank=" << rank
+                      << " role=" << SignalWaitRoleName(request.role)
+                      << " preSubmitDelayMs=" << effectiveDelayMs
+                      << " submitRet=" << submitRet
+                      << " syncRet=" << syncRet
+                      << " submitMs="
+                      << std::chrono::duration_cast<std::chrono::milliseconds>(submitEnd - submitBegin).count()
+                      << " syncMs="
+                      << std::chrono::duration_cast<std::chrono::milliseconds>(syncEnd - syncBegin).count()
+                      << std::endl;
+            if (syncRet != ACL_SUCCESS) {
+                std::cerr << "tilexr_ccu_signal_wait aclrtSynchronizeStream ret=" << syncRet << std::endl;
+                if (!plan.submitTasks.empty()) {
+                    PrintMissionContext(context, plan.submitTasks.front(), "tilexr_ccu_signal_wait");
+                }
+                finalRet = 8;
+            } else if (submitRet != TileXR::TILEXR_SUCCESS) {
+                finalRet = 9;
+            }
+            if (!WaitForCollectiveSubmitDone(rank, rankSize, finalRet) && finalRet == 0) {
+                finalRet = 13;
+            }
+            aclrtDestroyStream(stream);
+        }
+    }
+
+    const bool passed = finalRet == 0;
+    if (passed) {
+        std::cout << "tilexr_ccu_signal_wait result passed=1"
+                  << " rank=" << rank
+                  << " role=" << SignalWaitRoleName(request.role)
+                  << " ret=" << finalRet
+                  << std::endl;
+    } else {
+        std::cout << "tilexr_ccu_signal_wait result passed=0"
+                  << " rank=" << rank
+                  << " role=" << SignalWaitRoleName(request.role)
+                  << " ret=" << finalRet
+                  << std::endl;
+    }
+    if (ShouldFastExitAfterRun()) {
+        std::cout << "tilexr_ccu_signal_wait fastExitAfterRun=1"
+                  << " ret=" << finalRet
+                  << " reason=\"skipping prepared-task cleanup to isolate cleanup hangs\""
+                  << std::endl;
+        std::fflush(stdout);
+        std::fflush(stderr);
+        std::_Exit(finalRet);
+    }
+    const int releaseRet = TileXR::TileXRCcuReleaseDirectInstallAttemptResources(plan.attempt);
+    if (releaseRet != TileXR::TILEXR_SUCCESS) {
+        std::cerr << "tilexr_ccu_signal_wait prepared destroy ret=" << releaseRet << std::endl;
+        finalRet = finalRet == 0 ? 11 : finalRet;
+    }
+    return finalRet;
+}
+
 int RunPreparedSmokeForRank(DirectCcuSmokeContext* context, int rank, int rankSize, int device)
 {
+    if (SignalWaitSmokeEnabled() || BarrierSmokeEnabled()) {
+        return RunSignalWaitSmokeForRank(context, rank, rankSize, device);
+    }
+
     TileXRDirectCcuPrepareOptions options = MakePrepareOptions(rank, rankSize, device);
     const int peer = rankSize == 2 ? 1 - rank : (rank + 1) % rankSize;
     const bool p2pCcuCopyEnabled = EnvFlag(kP2pCcuCopyEnv);
+    const bool p2pCcuCopyActiveRank = !p2pCcuCopyEnabled || IsP2pCcuCopyActiveRank(rank);
+    const TileXR::TileXRCcuMemoryCopyDirection p2pCcuCopyDirection = P2pCcuCopyDirectionFromEnv();
     P2pCcuCopyState p2pCcuCopy;
     if (p2pCcuCopyEnabled) {
         options.syncResourceCount = 1;
@@ -1104,7 +1520,7 @@ int RunPreparedSmokeForRank(DirectCcuSmokeContext* context, int rank, int rankSi
             reinterpret_cast<uint64_t>(p2pCcuCopy.destination.ptr),
             p2pCcuCopy.bytes,
             static_cast<uint32_t>(peer),
-            TileXR::TileXRCcuMemoryCopyDirection::RemoteToLocal,
+            p2pCcuCopyDirection,
             prepared,
             &installReport) :
         context->planner.PrepareDirectCcuInstallAttempt(context->session, options, prepared, &installReport);
@@ -1129,6 +1545,8 @@ int RunPreparedSmokeForRank(DirectCcuSmokeContext* context, int rank, int rankSi
     } else if (submitRequested && !installReport.submitReady) {
         std::cout << "tilexr_ccu_direct_smoke submit skipped reason=\"prepare did not reach submitReady\""
                   << std::endl;
+    } else if (submitRequested && !p2pCcuCopyActiveRank) {
+        finalRet = RunInactiveP2pCcuCopyRank(rank, peer, rankSize, &p2pCcuCopy, finalRet);
     } else if (submitRequested) {
         aclrtStream stream = nullptr;
         int streamRet = aclrtCreateStream(&stream);
@@ -1171,6 +1589,9 @@ int RunPreparedSmokeForRank(DirectCcuSmokeContext* context, int rank, int rankSi
                 std::chrono::duration_cast<std::chrono::milliseconds>(syncEnd - syncBegin).count());
             if (syncRet != ACL_SUCCESS) {
                 std::cerr << "tilexr_ccu_direct_smoke aclrtSynchronizeStream ret=" << syncRet << std::endl;
+                if (prepared != nullptr && installReport.submitTaskCount > 0) {
+                    PrintMissionContext(context, prepared->submitTasks.front(), "tilexr_ccu_direct_smoke");
+                }
                 finalRet = 8;
             } else if (submitRet != TileXR::TILEXR_SUCCESS) {
                 finalRet = 9;
@@ -1370,6 +1791,7 @@ int main()
         std::_Exit(finalRet);
     }
     TraceLifecycle("before DirectCcuSmokeContext shutdown");
+    context.backend.Shutdown();
     context.session.Shutdown();
     TraceLifecycle("after DirectCcuSmokeContext shutdown");
     TraceLifecycle("before aclrtResetDevice");
diff --git a/tests/ccu/run_tilexr_ccu_direct_smoke.sh b/tests/ccu/run_tilexr_ccu_direct_smoke.sh
index 5afcac6b..e8fac065 100644
--- a/tests/ccu/run_tilexr_ccu_direct_smoke.sh
+++ b/tests/ccu/run_tilexr_ccu_direct_smoke.sh
@@ -37,6 +37,7 @@ endpoint_fields=(
 )
 
 resource_window_token_fields=(
+    EID_INDEX
     TOKEN_ID
     RAW_TOKEN_ID
     TOKEN_VALUE
@@ -80,6 +81,12 @@ default_sync_instruction_count()
     esac
 }
 
+signal_wait_mode_enabled()
+{
+    [ "${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT:-0}" = "1" ] ||
+        [ "${TILEXR_CCU_DIRECT_SMOKE_BARRIER:-0}" = "1" ]
+}
+
 apply_p2p_ccu_copy_defaults()
 {
     if [ "${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY:-0}" != "1" ]; then
@@ -112,9 +119,43 @@ apply_p2p_ccu_copy_defaults()
     export TILEXR_CCU_PROBE_CHANNEL_START="${TILEXR_CCU_PROBE_CHANNEL_START:-2}"
     export TILEXR_CCU_DIRECT_BARRIER_MODE="${TILEXR_CCU_DIRECT_BARRIER_MODE:-sync_cke}"
     export TILEXR_CCU_DIRECT_LOWER_LAYER_WQE_MODE="${TILEXR_CCU_DIRECT_LOWER_LAYER_WQE_MODE:-hcomm_cap}"
+    export TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN="${TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN:-1}"
+}
+
+apply_signal_wait_defaults()
+{
+    if ! signal_wait_mode_enabled; then
+        return
+    fi
+
+    export TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT="${TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT:-1}"
+    export TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK="${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK:-0}"
+    export TILEXR_CCU_PROBE_MISSION_START="${TILEXR_CCU_PROBE_MISSION_START:-6}"
+    export TILEXR_CCU_PROBE_INSTRUCTION_START="${TILEXR_CCU_PROBE_INSTRUCTION_START:-475}"
+    export TILEXR_CCU_PROBE_MISSION_INSTRUCTION_START="${TILEXR_CCU_PROBE_MISSION_INSTRUCTION_START:-489}"
+    export TILEXR_CCU_PROBE_SQE_ARG_COUNT="${TILEXR_CCU_PROBE_SQE_ARG_COUNT:-0}"
+    export TILEXR_CCU_PROBE_SYNC_RESOURCE_COUNT="${TILEXR_CCU_PROBE_SYNC_RESOURCE_COUNT:-1}"
+    export TILEXR_CCU_DIRECT_REPOSITORY_INSTALL_WINDOW="${TILEXR_CCU_DIRECT_REPOSITORY_INSTALL_WINDOW:-full_repository}"
+    export TILEXR_CCU_DIRECT_REPOSITORY_DATA_LEN_MODE="${TILEXR_CCU_DIRECT_REPOSITORY_DATA_LEN_MODE:-instruction_bytes}"
+    export TILEXR_CCU_DIRECT_REPOSITORY_MEMORY_ALLOC_MODE="${TILEXR_CCU_DIRECT_REPOSITORY_MEMORY_ALLOC_MODE:-acl}"
+    export TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE="${TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE:-ra_ctx}"
+    export TILEXR_CCU_PROBE_RANK0_XN_START="${TILEXR_CCU_PROBE_RANK0_XN_START:-1961}"
+    export TILEXR_CCU_PROBE_RANK1_XN_START="${TILEXR_CCU_PROBE_RANK1_XN_START:-1961}"
+    export TILEXR_CCU_PROBE_RANK0_REMOTE_XN_START="${TILEXR_CCU_PROBE_RANK0_REMOTE_XN_START:-2361}"
+    export TILEXR_CCU_PROBE_RANK1_REMOTE_XN_START="${TILEXR_CCU_PROBE_RANK1_REMOTE_XN_START:-2361}"
+    export TILEXR_CCU_PROBE_REMOTE_XN_COUNT="${TILEXR_CCU_PROBE_REMOTE_XN_COUNT:-8}"
+    export TILEXR_CCU_PROBE_RANK0_LOCAL_WAIT_CKE_START="${TILEXR_CCU_PROBE_RANK0_LOCAL_WAIT_CKE_START:-332}"
+    export TILEXR_CCU_PROBE_RANK1_LOCAL_WAIT_CKE_START="${TILEXR_CCU_PROBE_RANK1_LOCAL_WAIT_CKE_START:-332}"
+    export TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_COUNT="${TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_COUNT:-8}"
+    export TILEXR_CCU_PROBE_RANK0_REMOTE_NOTIFY_CKE_START="${TILEXR_CCU_PROBE_RANK0_REMOTE_NOTIFY_CKE_START:-364}"
+    export TILEXR_CCU_PROBE_RANK1_REMOTE_NOTIFY_CKE_START="${TILEXR_CCU_PROBE_RANK1_REMOTE_NOTIFY_CKE_START:-364}"
+    export TILEXR_CCU_PROBE_REMOTE_NOTIFY_CKE_COUNT="${TILEXR_CCU_PROBE_REMOTE_NOTIFY_CKE_COUNT:-8}"
+    export TILEXR_CCU_PROBE_CHANNEL_START="${TILEXR_CCU_PROBE_CHANNEL_START:-2}"
+    export TILEXR_CCU_DIRECT_LOWER_LAYER_WQE_MODE="${TILEXR_CCU_DIRECT_LOWER_LAYER_WQE_MODE:-hcomm_cap}"
 }
 
 apply_p2p_ccu_copy_defaults
+apply_signal_wait_defaults
 
 if [ "${TILEXR_CCU_DIRECT_SMOKE_DRY_RUN:-0}" = "1" ]; then
     echo "tilexr_ccu_direct_smoke_runner dryRun=1 workDir=${work_dir}"
@@ -126,6 +167,9 @@ if [ "${TILEXR_CCU_DIRECT_SMOKE_DRY_RUN:-0}" = "1" ]; then
         TILEXR_CCU_DIRECT_REPOSITORY_MEMORY_ALLOC_MODE \
         TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE \
         TILEXR_CCU_DIRECT_INSTALL_ORDER \
+        TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT \
+        TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK \
+        TILEXR_CCU_DIRECT_SMOKE_BARRIER \
         TILEXR_CCU_PROBE_SQE_ARG_COUNT \
         TILEXR_CCU_PROBE_MISSION_INSTRUCTION_START; do
         diagnostic_value="${!diagnostic_var:-}"
@@ -336,6 +380,12 @@ fi
 if [ "${TILEXR_CCU_DIRECT_INSTALL_DIE_ID:-}" != "" ]; then
     common_env+=("TILEXR_CCU_DIRECT_INSTALL_DIE_ID=${TILEXR_CCU_DIRECT_INSTALL_DIE_ID}")
 fi
+if [ "${TILEXR_CCU_DIRECT_TRACE:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_TRACE=${TILEXR_CCU_DIRECT_TRACE}")
+fi
+if [ "${TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE=${TILEXR_CCU_DIRECT_TRACE_ENDPOINT_ROUTE}")
+fi
 if [ "${TILEXR_CCU_PROBE_MISSION_START:-}" != "" ]; then
     common_env+=("TILEXR_CCU_PROBE_MISSION_START=${TILEXR_CCU_PROBE_MISSION_START}")
 fi
@@ -378,6 +428,27 @@ fi
 if [ "${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_BYTES:-}" != "" ]; then
     common_env+=("TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_BYTES=${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_BYTES}")
 fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN=${TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_TRACE_LIFECYCLE:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_TRACE_LIFECYCLE=${TILEXR_CCU_DIRECT_SMOKE_TRACE_LIFECYCLE}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_RESOURCE_WINDOW:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_RESOURCE_WINDOW=${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_RESOURCE_WINDOW}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SMOKE_BARRIER:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SMOKE_BARRIER=${TILEXR_CCU_DIRECT_SMOKE_BARRIER}")
+fi
+if [ "${TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT:-}" != "" ]; then
+    common_env+=("TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT=${TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT}")
+fi
 if [ "${TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_START:-}" != "" ]; then
     common_env+=("TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_START=${TILEXR_CCU_PROBE_LOCAL_WAIT_CKE_START}")
 fi
@@ -488,7 +559,7 @@ for token_field in "${resource_window_token_fields[@]}"; do
     fi
 done
 
-echo "tilexr_ccu_direct_smoke_runner begin workDir=${work_dir} devices=${devices} commId=${comm_id} threadMode=${TILEXR_CCU_DIRECT_SMOKE_THREAD_MODE:-0} submit=${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0} barrierMode=${TILEXR_CCU_DIRECT_BARRIER_MODE:-} p2pCcuCopy=${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY:-0} timeout=${timeout_s} npuSmiTimeout=${TILEXR_CCU_SMOKE_NPU_SMI_TIMEOUT:-20}"
+echo "tilexr_ccu_direct_smoke_runner begin workDir=${work_dir} devices=${devices} commId=${comm_id} threadMode=${TILEXR_CCU_DIRECT_SMOKE_THREAD_MODE:-0} submit=${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0} barrierMode=${TILEXR_CCU_DIRECT_BARRIER_MODE:-} p2pCcuCopy=${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY:-0} signalWait=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT:-0} signalRank=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK:-0} ccuBarrier=${TILEXR_CCU_DIRECT_SMOKE_BARRIER:-0} timeout=${timeout_s} npuSmiTimeout=${TILEXR_CCU_SMOKE_NPU_SMI_TIMEOUT:-20}"
 
 if [ "${TILEXR_CCU_DIRECT_SMOKE_THREAD_MODE:-0}" = "1" ]; then
     thread_log="${work_dir}/ccu_thread.log"
@@ -503,27 +574,56 @@ if [ "${TILEXR_CCU_DIRECT_SMOKE_THREAD_MODE:-0}" = "1" ]; then
         echo "thread log: ${thread_log}" >&2
         exit 4
     fi
-    if [ "$(grep -c "tilexr_ccu_direct_smoke prepare ret=0" "${thread_log}")" -lt 2 ]; then
-        echo "ERROR: direct CCU thread-mode prepare did not return success for both ranks" >&2
-        exit 5
-    fi
-    if [ "$(grep -c "installSucceeded=1" "${thread_log}")" -lt 2 ]; then
-        echo "ERROR: direct CCU thread-mode prepare did not complete install attempt for both ranks" >&2
-        exit 6
-    fi
-    if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]; then
-        if [ "$(grep -c "submitReady=1" "${thread_log}")" -lt 2 ]; then
-            echo "ERROR: direct CCU thread-mode submit requested but prepare did not reach submitReady=1" >&2
+    if signal_wait_mode_enabled; then
+        if [ "$(grep -c "tilexr_ccu_signal_wait prepare ret=0" "${thread_log}")" -lt 2 ]; then
+            echo "ERROR: direct CCU signal/wait thread-mode prepare did not return success for both ranks" >&2
+            exit 5
+        fi
+        if [ "$(grep -c "installSucceeded=1" "${thread_log}")" -lt 2 ]; then
+            echo "ERROR: direct CCU signal/wait thread-mode prepare did not complete install attempt for both ranks" >&2
             exit 6
         fi
-        if [ "$(grep -c "tilexr_ccu_direct_smoke submit ret=0" "${thread_log}")" -lt 2 ]; then
-            echo "ERROR: direct CCU thread-mode submit did not return success for both ranks" >&2
-            exit 7
+        if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]; then
+            if [ "$(grep -c "submitReady=1" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU signal/wait thread-mode submit requested but prepare did not reach submitReady=1" >&2
+                exit 6
+            fi
+            if [ "$(grep -c "tilexr_ccu_signal_wait submit ret=0" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU signal/wait thread-mode submit did not return success for both ranks" >&2
+                exit 7
+            fi
+            if [ "$(grep -c "tilexr_ccu_signal_wait timing" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU signal/wait thread-mode timing was not reported for both ranks" >&2
+                exit 8
+            fi
         fi
-        if [ "$(grep -c "tilexr_ccu_direct_smoke submitTiming" "${thread_log}")" -lt 2 ]; then
-            echo "ERROR: direct CCU thread-mode submit timing was not reported for both ranks" >&2
+        if [ "$(grep -c "tilexr_ccu_signal_wait result passed=1" "${thread_log}")" -lt 2 ]; then
+            echo "ERROR: direct CCU signal/wait thread-mode result did not pass for both ranks" >&2
             exit 8
         fi
+    else
+        if [ "$(grep -c "tilexr_ccu_direct_smoke prepare ret=0" "${thread_log}")" -lt 2 ]; then
+            echo "ERROR: direct CCU thread-mode prepare did not return success for both ranks" >&2
+            exit 5
+        fi
+        if [ "$(grep -c "installSucceeded=1" "${thread_log}")" -lt 2 ]; then
+            echo "ERROR: direct CCU thread-mode prepare did not complete install attempt for both ranks" >&2
+            exit 6
+        fi
+        if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]; then
+            if [ "$(grep -c "submitReady=1" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU thread-mode submit requested but prepare did not reach submitReady=1" >&2
+                exit 6
+            fi
+            if [ "$(grep -c "tilexr_ccu_direct_smoke submit ret=0" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU thread-mode submit did not return success for both ranks" >&2
+                exit 7
+            fi
+            if [ "$(grep -c "tilexr_ccu_direct_smoke submitTiming" "${thread_log}")" -lt 2 ]; then
+                echo "ERROR: direct CCU thread-mode submit timing was not reported for both ranks" >&2
+                exit 8
+            fi
+        fi
     fi
     if [ "${TILEXR_CCU_DIRECT_SMOKE_EXPECT_P2P_CCU_COPY:-0}" = "1" ]; then
         if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" != "1" ]; then
@@ -566,16 +666,36 @@ if [ "${rank0_status}" -ne 0 ] || [ "${rank1_status}" -ne 0 ]; then
     exit 4
 fi
 
-for log in "${rank0_log}" "${rank1_log}"; do
-    if ! grep -q "tilexr_ccu_direct_smoke prepare ret=0" "${log}"; then
-        echo "ERROR: direct CCU prepare did not return success in ${log}" >&2
-        exit 5
-    fi
-    if ! grep -q "installSucceeded=1" "${log}"; then
-        echo "ERROR: direct CCU prepare did not complete install attempt in ${log}" >&2
-        exit 6
-    fi
-done
+if signal_wait_mode_enabled; then
+    for log in "${rank0_log}" "${rank1_log}"; do
+        if ! grep -q "tilexr_ccu_signal_wait prepare ret=0" "${log}"; then
+            echo "ERROR: direct CCU signal/wait prepare did not return success in ${log}" >&2
+            exit 5
+        fi
+        if ! grep -q "installSucceeded=1" "${log}"; then
+            echo "ERROR: direct CCU signal/wait prepare did not complete install attempt in ${log}" >&2
+            exit 6
+        fi
+    done
+else
+    for log in "${rank0_log}" "${rank1_log}"; do
+        if ! grep -q "tilexr_ccu_direct_smoke prepare ret=0" "${log}"; then
+            echo "ERROR: direct CCU prepare did not return success in ${log}" >&2
+            exit 5
+        fi
+        if ! grep -q "installSucceeded=1" "${log}"; then
+            echo "ERROR: direct CCU prepare did not complete install attempt in ${log}" >&2
+            exit 6
+        fi
+    done
+fi
+
+rank_skipped_p2p_ccu_copy_submit()
+{
+    local log="$1"
+    [ "${TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY:-0}" = "1" ] &&
+        grep -q "tilexr_ccu_direct_smoke p2pCcuCopy skipped" "${log}"
+}
 
 if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]; then
     for log in "${rank0_log}" "${rank1_log}"; do
@@ -585,17 +705,43 @@ if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]; then
         fi
     done
     for log in "${rank0_log}" "${rank1_log}"; do
+        if signal_wait_mode_enabled; then
+            if ! grep -q "tilexr_ccu_signal_wait submit ret=0" "${log}"; then
+                echo "ERROR: direct CCU signal/wait submit did not return success in ${log}" >&2
+                exit 7
+            fi
+            if ! grep -q "tilexr_ccu_signal_wait timing" "${log}"; then
+                echo "ERROR: direct CCU signal/wait timing was not reported in ${log}" >&2
+                exit 8
+            fi
+            continue
+        fi
         if ! grep -q "tilexr_ccu_direct_smoke submit ret=0" "${log}"; then
+            if rank_skipped_p2p_ccu_copy_submit "${log}"; then
+                continue
+            fi
             echo "ERROR: direct CCU submit did not return success in ${log}" >&2
             exit 7
         fi
         if ! grep -q "tilexr_ccu_direct_smoke submitTiming" "${log}"; then
+            if rank_skipped_p2p_ccu_copy_submit "${log}"; then
+                continue
+            fi
             echo "ERROR: direct CCU submit timing was not reported in ${log}" >&2
             exit 8
         fi
     done
 fi
 
+if signal_wait_mode_enabled; then
+    for log in "${rank0_log}" "${rank1_log}"; do
+        if ! grep -q "tilexr_ccu_signal_wait result passed=1" "${log}"; then
+            echo "ERROR: direct CCU signal/wait result did not pass in ${log}" >&2
+            exit 9
+        fi
+    done
+fi
+
 if [ "${TILEXR_CCU_DIRECT_SMOKE_EXPECT_BARRIER_WAIT:-0}" = "1" ]; then
     if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" != "1" ]; then
         echo "ERROR: direct CCU barrier wait check requires TILEXR_CCU_DIRECT_SMOKE_SUBMIT=1" >&2
@@ -610,7 +756,7 @@ if [ "${TILEXR_CCU_DIRECT_SMOKE_EXPECT_BARRIER_WAIT:-0}" = "1" ]; then
     fi
     wait_sync_ms="$(
         awk '
-            /tilexr_ccu_direct_smoke submitTiming/ {
+            /tilexr_ccu_direct_smoke submitTiming|tilexr_ccu_signal_wait timing/ {
                 for (i = 1; i <= NF; ++i) {
                     if ($i ~ /^syncMs=/) {
                         split($i, parts, "=");
diff --git a/tests/ccu/test_tilexr_ccu_backend_boundary.py b/tests/ccu/test_tilexr_ccu_backend_boundary.py
index 1985c58b..55e645b8 100644
--- a/tests/ccu/test_tilexr_ccu_backend_boundary.py
+++ b/tests/ccu/test_tilexr_ccu_backend_boundary.py
@@ -9,6 +9,7 @@
 
 REPO_ROOT = Path(__file__).resolve().parents[2]
 COMM_HEADER = REPO_ROOT / "src" / "comm" / "tilexr_comm.h"
+PUBLIC_API_HEADER = REPO_ROOT / "src" / "include" / "tilexr_api.h"
 BACKEND_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_backend.h"
 BACKEND_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_backend.cpp"
 RUNTIME_SESSION_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_runtime_session.h"
@@ -54,11 +55,21 @@ def test_tilexr_comm_header_owns_only_opaque_backend(self):
 
     def test_backend_header_exposes_facade_not_public_c_api(self):
         header = BACKEND_HEADER.read_text(encoding="utf-8")
+        public_api = PUBLIC_API_HEADER.read_text(encoding="utf-8")
         self.assertIn("class TileXRCcuBackend", header)
         self.assertIn("struct TileXRCcuBackendOptions", header)
         self.assertIn("TileXRSockExchange *exchange", header)
         self.assertIn("PrepareCollective", header)
         self.assertIn("SubmitCollective", header)
+        for needle in [
+            "enum class TileXRCcuSignalWaitRole",
+            "struct TileXRCcuSignalWaitRequest",
+            "struct TileXRCcuSignalWaitPlan",
+            "PrepareSignalWait",
+            "SubmitSignalWait",
+        ]:
+            with self.subTest(internal=needle):
+                self.assertIn(needle, header)
         for needle in [
             "TileXRDirectCcuPreparedTasksPtr",
             "TileXRCommPrepareDirectCcu",
@@ -66,6 +77,14 @@ def test_backend_header_exposes_facade_not_public_c_api(self):
         ]:
             with self.subTest(needle=needle):
                 self.assertNotIn(needle, header)
+                self.assertNotIn(needle, public_api)
+        for needle in [
+            "TileXRCcuSignalWait",
+            "PrepareSignalWait",
+            "SubmitSignalWait",
+        ]:
+            with self.subTest(public_needle=needle):
+                self.assertNotIn(needle, public_api)
 
     def test_split_sources_own_restored_direct_ccu_runtime_glue(self):
         source = BACKEND_SOURCE.read_text(encoding="utf-8")
@@ -92,13 +111,18 @@ def test_split_sources_own_restored_direct_ccu_runtime_glue(self):
         for needle in [
             "#include \"ccu/tilexr_ccu_repository.h\"",
             "PrepareDirectCcuInstallAttempt",
+            "PrepareSignalWait",
             "PrepareDirectCcuLowerLayerPlanCallback",
             "TileXRCcuRunDirectInstallAttempt",
+            "TileXRCcuRunDirectSignalWaitInstallAttempt",
             "TileXRCcuMakeRepositoryDeviceMemoryOps",
         ]:
             with self.subTest(needle=needle):
                 self.assertIn(needle, planner)
 
+        self.assertIn("planner_->PrepareSignalWait", source)
+        self.assertNotIn("return TILEXR_ERROR_NOT_SUPPORT;", source[source.index("PrepareSignalWait"):])
+
         planner_header = PLANNER_HEADER.read_text(encoding="utf-8")
         executor_header = EXECUTOR_HEADER.read_text(encoding="utf-8")
         self.assertRegex(
@@ -125,6 +149,80 @@ def test_split_sources_own_restored_direct_ccu_runtime_glue(self):
             with self.subTest(fake_ready=fake_ready):
                 self.assertNotIn(fake_ready, source + "\n" + runtime + "\n" + planner + "\n" + executor)
 
+    def test_tilexr_comm_can_auto_initialize_ccu_backend_without_blocking_comm_init(self):
+        source = (REPO_ROOT / "src" / "comm" / "tilexr_comm.cpp").read_text(encoding="utf-8")
+
+        self.assertIn('constexpr const char* TILEXR_ENABLE_CCU_BACKEND_ENV = "TILEXR_ENABLE_CCU_BACKEND"', source)
+        self.assertIn("bool ShouldEnableCcuBackend()", source)
+        self.assertIn("int TileXRComm::InitCcuBackendIfEnabled()", source)
+        self.assertIn("const int ccuRet = InitCcuBackend()", source)
+        self.assertIn("TileXR CCU backend init failed, direct CCU disabled", source)
+        self.assertIn("TileXR CCU backend initialized", source)
+
+        process_init = source[source.index("int TileXRComm::Init()"): source.index("int TileXRComm::InitThread")]
+        thread_init = source[source.index("int TileXRComm::InitThread"): source.index("int TileXRComm::EnablePeerAccess")]
+
+        for body_name, body in [("process", process_init), ("thread", thread_init)]:
+            with self.subTest(body=body_name):
+                self.assertIn("ret = InitCcuBackendIfEnabled();", body)
+                self.assertLess(body.index("ret = InitCcuBackendIfEnabled();"), body.index("ret = SyncCommArgs();"))
+                self.assertIn("if (ret != TILEXR_SUCCESS) {", body)
+                self.assertIn("return ret;", body)
+
+        helper = source[
+            source.index("int TileXRComm::InitCcuBackendIfEnabled()"):
+            source.index("TileXRCcuBackend *TileXRComm::GetCcuBackendForCollectives")
+        ]
+        self.assertIn("if (!ShouldEnableCcuBackend())", helper)
+        self.assertIn("return TILEXR_SUCCESS;", helper)
+        self.assertNotIn("return ccuRet;", helper)
+
+    def test_p2p_ccu_copy_process_token_is_not_urma_shifted(self):
+        planner = PLANNER_SOURCE.read_text(encoding="utf-8")
+        token_query = planner[
+            planner.index("int QueryDirectCcuProcessMemoryToken"):
+            planner.index("int BuildDirectCcuLocalMemoryCopyEndpoint")
+        ]
+
+        self.assertIn("rtUbDevQueryInfo(QUERY_PROCESS_TOKEN, &info)", token_query)
+        self.assertIn("const uint32_t tokenId = info.tokenId;", token_query)
+        self.assertIn("TileXRCcuPackMemoryToken(tokenId, info.tokenValue, true)", token_query)
+        self.assertNotIn("info.tokenId >>", token_query)
+
+    def test_p2p_ccu_copy_uses_original_va_for_microcode_and_imported_segva_for_route(self):
+        planner = PLANNER_SOURCE.read_text(encoding="utf-8")
+        endpoint_builder = planner[
+            planner.index("int BuildDirectCcuLocalMemoryCopyEndpoint"):
+            planner.index("void TileXRCcuCollectivePlanner::Reset")
+        ]
+        prepare_copy = planner[
+            planner.index("int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt"):
+            planner.index("int TileXRCcuCollectivePlanner::RefreshDirectCcuLowerLayerPlan")
+        ]
+
+        self.assertIn("session.RegisterCcuResourceRmaBuffer(basicInfo->resourceAddr)", prepare_copy)
+        self.assertIn("session.RegisterMemoryBuffer(sourceAddr, bytes, &sourceInfo)", endpoint_builder)
+        self.assertIn("session.RegisterMemoryBuffer(destinationAddr, bytes, &destinationInfo)", endpoint_builder)
+        self.assertIn("endpoint->sourceAddr = sourceInfo.addr", endpoint_builder)
+        self.assertIn("endpoint->destinationAddr = destinationInfo.addr", endpoint_builder)
+        self.assertNotIn("endpoint->sourceAddr = sourceInfo.targetSegVa", endpoint_builder)
+        self.assertNotIn("endpoint->destinationAddr = destinationInfo.targetSegVa", endpoint_builder)
+        self.assertIn("endpoint->sourceRemoteImport.key = sourceInfo.key", endpoint_builder)
+        self.assertIn("endpoint->destinationRemoteImport.key = destinationInfo.key", endpoint_builder)
+        self.assertIn("TileXRCcuPackMemoryToken(sourceInfo.tokenId, sourceInfo.tokenValue, true)", endpoint_builder)
+        self.assertIn(
+            "TileXRCcuPackMemoryToken(destinationInfo.tokenId, destinationInfo.tokenValue, true)",
+            endpoint_builder,
+        )
+        self.assertIn("remoteImportRequest = peerEndpoint.sourceRemoteImport", prepare_copy)
+        self.assertIn("remoteImportRequest = peerEndpoint.destinationRemoteImport", prepare_copy)
+        self.assertIn("session.ImportRemoteMemoryBuffer(remoteImportRequest, &importedRemoteBuffer)", prepare_copy)
+        self.assertIn("memoryCopy.remoteAddr = remoteImportRequest.addr", prepare_copy)
+        self.assertIn("SetDirectCcuRemoteRouteMemoryOverride(", prepare_copy)
+        self.assertIn("importedRemoteBuffer.targetSegVa", prepare_copy)
+        self.assertNotIn("QueryDirectCcuProcessMemoryToken(sourceAddr", endpoint_builder)
+        self.assertNotIn("QueryDirectCcuProcessMemoryToken(destinationAddr", endpoint_builder)
+
 
 if __name__ == "__main__":
     unittest.main()
diff --git a/tests/ccu/test_tilexr_ccu_barrier_program.py b/tests/ccu/test_tilexr_ccu_barrier_program.py
index 254a1b82..cbecf46a 100644
--- a/tests/ccu/test_tilexr_ccu_barrier_program.py
+++ b/tests/ccu/test_tilexr_ccu_barrier_program.py
@@ -242,7 +242,7 @@ def test_barrier_program_can_emit_hcomm_like_synccke_post_and_clear_wait(self):
                     std::cerr << "source CKE init mismatch\n";
                     return 3;
                 }
-                if (program[1].words[0] != 0x000102210330100bULL ||
+                if (program[1].words[0] != 0xffff02210330100bULL ||
                     program[1].words[1] != 0x0000000000000002ULL ||
                     program[1].words[2] != 0x0001000000000000ULL ||
                     program[1].words[3] != 0) {
@@ -309,7 +309,7 @@ def test_barrier_program_can_emit_synccke_post_and_set_wait(self):
                     std::cerr << "source CKE init mismatch\n";
                     return 3;
                 }
-                if (program[1].words[0] != 0x000102210330100bULL ||
+                if (program[1].words[0] != 0xffff02210330100bULL ||
                     program[1].words[1] != 0x0000000000000002ULL ||
                     program[1].words[2] != 0x0001000000000000ULL ||
                     program[1].words[3] != 0) {
@@ -374,7 +374,7 @@ def test_barrier_program_can_emit_synccke_post_only_diagnostic(self):
                     std::cerr << "source CKE init mismatch\n";
                     return 3;
                 }
-                if (program[1].words[0] != 0x000102210330100bULL ||
+                if (program[1].words[0] != 0xffff02210330100bULL ||
                     program[1].words[1] != 0x0000000000000002ULL ||
                     program[1].words[2] != 0x0001000000000000ULL ||
                     program[1].words[3] != 0) {
diff --git a/tests/ccu/test_tilexr_ccu_direct_orchestrator.py b/tests/ccu/test_tilexr_ccu_direct_orchestrator.py
index d45dc384..b2bb6587 100644
--- a/tests/ccu/test_tilexr_ccu_direct_orchestrator.py
+++ b/tests/ccu/test_tilexr_ccu_direct_orchestrator.py
@@ -27,6 +27,7 @@
 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"
+SIGNAL_WAIT_PROGRAM_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_signal_wait_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"
@@ -137,6 +138,7 @@ def compile_and_run(self, code: str):
                 str(BARRIER_SOURCE),
                 str(MICROCODE_SOURCE),
                 str(MEMORY_PROGRAM_SOURCE),
+                str(SIGNAL_WAIT_PROGRAM_SOURCE),
                 str(RUNTIME_SOURCE),
                 str(LOWER_LAYER_PLAN_SOURCE),
                 str(LOWER_LAYER_PAYLOAD_SOURCE),
@@ -778,6 +780,29 @@ def test_direct_submit_default_runtime_trace_indexes_final_tasks(self):
         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_task_timeout_override_is_applied_before_install_trace_and_manifest(self):
+        source = DIRECT_SOURCE.read_text(encoding="utf-8")
+        body = source[
+            source.index("int RunDirectInstallAttemptImpl("):
+            source.index("int TileXRCcuRunDirectMemoryCopyInstallAttempt(")
+        ]
+
+        first_override = body.index("ApplyTaskTimeoutOverride(options.taskTimeout, attempt)")
+        self.assertLess(first_override, body.index("TraceDirectInstallAttempt(*attempt)"))
+        self.assertLess(first_override, body.index("TileXRCcuBuildInstallManifest"))
+        self.assertLess(first_override, body.index("TileXRCcuInstallHardware"))
+
+    def test_signal_wait_uses_notify_mask_for_source_cke_reserve(self):
+        source = DIRECT_SOURCE.read_text(encoding="utf-8")
+        body = source[
+            source.index("int BuildDirectSignalWaitLaunchPackage("):
+            source.index("void FillReportFromAttempt(")
+        ]
+
+        self.assertIn("spec.remoteNotifyMask = resource.remoteNotifyMask == 0 ? 1U : resource.remoteNotifyMask", body)
+        self.assertIn("spec.sourceCkeMask = spec.remoteNotifyMask", body)
+        self.assertNotIn("spec.sourceCkeMask = resource.sourceCkeMask", body)
+
     def test_direct_install_attempt_can_prepare_lower_layer_plan_after_allocation(self):
         code = textwrap.dedent(
             r'''
@@ -1664,6 +1689,7 @@ def test_direct_orchestrator_is_wired_and_has_no_private_runtime_surface(self):
         self.assertIn("ccu/tilexr_ccu_direct_orchestrator.cpp", cmake)
         self.assertIn("struct TileXRCcuDirectInstallOptions", header)
         self.assertIn("struct TileXRCcuDirectMemoryCopySpec", header)
+        self.assertIn("struct TileXRCcuDirectSignalWaitSpec", header)
         self.assertIn("struct TileXRCcuDirectInstallAttempt", header)
         self.assertIn("struct TileXRCcuDirectInstallReport", header)
         self.assertIn("struct TileXRCcuDirectSubmitReport", header)
@@ -1677,18 +1703,24 @@ def test_direct_orchestrator_is_wired_and_has_no_private_runtime_surface(self):
         self.assertIn("TileXRCcuSubmitPreparedTasks", header)
         self.assertIn("TileXRCcuRunDirectInstallAttempt", header)
         self.assertIn("TileXRCcuRunDirectMemoryCopyInstallAttempt", header)
+        self.assertIn("TileXRCcuRunDirectSignalWaitInstallAttempt", 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("TileXRCcuBuildSignalWaitProgram", source)
+        self.assertIn("BuildDirectSignalWaitLaunchPackage", source)
+        self.assertIn("spec.localWaitCke = resource.localWaitCke == 0 ? resource.notifyCke : resource.localWaitCke", 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("uint16_t taskTimeout = 0", header)
+        self.assertIn("ApplyTaskTimeoutOverride(options.taskTimeout, attempt)", source)
         self.assertIn("remote XN install provider is missing", source)
         self.assertIn("TraceDecodedInstr", source)
         self.assertIn("TraceDecodedPfeCtx", source)
@@ -1719,6 +1751,18 @@ def test_direct_install_options_default_to_lower_layer_first(self):
         )
         self.assertIn("installRequest.installOrder = options.installOrder", source)
 
+    def test_direct_memory_copy_submit_tasks_keep_planned_runtime_payload(self):
+        source = DIRECT_SOURCE.read_text(encoding="utf-8")
+        memory_copy_body = source[
+            source.index("int BuildDirectMemoryCopyLaunchPackage"):
+            source.index("void FillReportFromAttempt")
+        ]
+
+        self.assertIn("TileXRCcuBuildTasks(attempt->plan, &tasks, &planReport)", memory_copy_body)
+        self.assertIn("attempt->package.tasks = tasks", memory_copy_body)
+        self.assertNotIn("NormalizeDirectMemoryCopySubmitTasks", source)
+        self.assertNotIn("task.argSize = 1", source)
+
 
 if __name__ == "__main__":
     unittest.main()
diff --git a/tests/ccu/test_tilexr_ccu_direct_smoke_probe.py b/tests/ccu/test_tilexr_ccu_direct_smoke_probe.py
index 7c8279bc..960cdadc 100644
--- a/tests/ccu/test_tilexr_ccu_direct_smoke_probe.py
+++ b/tests/ccu/test_tilexr_ccu_direct_smoke_probe.py
@@ -242,10 +242,71 @@ def test_p2p_ccu_copy_mode_prepares_memory_copy_task(self):
         self.assertIn("RunP2pCcuCopy", source)
         self.assertIn("PrepareDirectCcuMemoryCopyInstallAttempt", source)
         self.assertIn("TileXRCcuMemoryCopyDirection::RemoteToLocal", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_DIRECTION", source)
+        self.assertIn("P2pCcuCopyDirectionFromEnv", source)
+        self.assertIn("TileXRCcuMemoryCopyDirection::LocalToRemote", source)
         self.assertIn("ACL_MEMCPY_HOST_TO_DEVICE", source)
         self.assertIn("ACL_MEMCPY_DEVICE_TO_HOST", source)
         self.assertIn("p2pCcuCopy", source)
 
+    def test_prepare_options_read_submit_timeout(self):
+        source = PROBE_SOURCE.read_text(encoding="utf-8")
+        prepare_options = source[
+            source.index("TileXRDirectCcuPrepareOptions MakePrepareOptions"):
+            source.index("const char* SignalWaitRoleName")
+        ]
+
+        self.assertIn(
+            'options.taskTimeout = static_cast<uint16_t>(EnvInt("TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT", 0));',
+            prepare_options,
+        )
+
+    def test_local_to_remote_p2p_ccu_copy_inactive_rank_checks_destination(self):
+        source = PROBE_SOURCE.read_text(encoding="utf-8")
+        prepared_body = source[
+            source.index("int RunPreparedSmokeForRank"):
+            source.index("int RunThreadModeSmoke")
+        ]
+
+        self.assertIn("ShouldCheckInactiveP2pCcuCopyRank", source)
+        self.assertIn("RunInactiveP2pCcuCopyRank", source)
+        self.assertIn('"local_to_remote"', source)
+        self.assertIn("CheckP2pCcuCopyState", source)
+        inactive_branch = prepared_body[
+            prepared_body.index("!p2pCcuCopyActiveRank"):
+            prepared_body.index("} else if (submitRequested) {", prepared_body.index("!p2pCcuCopyActiveRank"))
+        ]
+        self.assertIn("RunInactiveP2pCcuCopyRank", inactive_branch)
+        self.assertIn("CheckP2pCcuCopyState", source[source.index("RunInactiveP2pCcuCopyRank"):])
+        self.assertNotIn("TileXRDirectCcuSubmitPrepared", inactive_branch)
+
+    def test_p2p_ccu_copy_submit_is_limited_to_active_rank(self):
+        source = PROBE_SOURCE.read_text(encoding="utf-8")
+        prepared_body = source[
+            source.index("int RunPreparedSmokeForRank"):
+            source.index("int RunThreadModeSmoke")
+        ]
+
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_P2P_CCU_COPY_ACTIVE_RANK", source)
+        self.assertIn("IsP2pCcuCopyActiveRank", source)
+        self.assertIn("p2pCcuCopyActiveRank", prepared_body)
+        self.assertIn("!p2pCcuCopyActiveRank", prepared_body)
+        self.assertIn("PrintP2pCcuCopySkipped", source)
+        self.assertIn("inactive p2p CCU-copy rank", source)
+        self.assertIn("WaitForInactiveP2pCcuCopyRank", source)
+        inactive_branch = prepared_body[
+            prepared_body.index("!p2pCcuCopyActiveRank"):
+            prepared_body.index("} else if (submitRequested) {", prepared_body.index("!p2pCcuCopyActiveRank"))
+        ]
+        inactive_helper = source[
+            source.index("int RunInactiveP2pCcuCopyRank"):
+            source.index("int RunPreparedSmokeForRank")
+        ]
+        self.assertIn("RunInactiveP2pCcuCopyRank", inactive_branch)
+        self.assertIn("PrintP2pCcuCopySkipped", inactive_helper)
+        self.assertIn("WaitForInactiveP2pCcuCopyRank", inactive_helper)
+        self.assertNotIn("TileXRDirectCcuSubmitPrepared", inactive_branch)
+
     def test_thread_mode_path_uses_single_process_init_and_never_rank_ipc_init(self):
         source = PROBE_SOURCE.read_text(encoding="utf-8")
 
@@ -259,6 +320,29 @@ def test_thread_mode_path_uses_single_process_init_and_never_rank_ipc_init(self)
         self.assertNotIn("TileXRCommInitRankWithDomain", thread_mode_body)
         self.assertNotIn("TILEXR_COMM_ID", thread_mode_body)
 
+    def test_probe_can_run_signal_wait_through_internal_backend(self):
+        source = PROBE_SOURCE.read_text(encoding="utf-8")
+        prepared_body = source[
+            source.index("int RunPreparedSmokeForRank"):
+            source.index("int RunThreadModeSmoke")
+        ]
+
+        self.assertIn("ccu/tilexr_ccu_backend.h", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_BARRIER", source)
+        self.assertIn("TileXRCcuBackend backend", source)
+        self.assertIn("context->backend.Init", source)
+        self.assertIn("RunSignalWaitSmokeForRank", source)
+        self.assertIn("TileXRCcuSignalWaitRequest", source)
+        self.assertIn("PrepareSignalWait", source)
+        self.assertIn("SubmitSignalWait", source)
+        self.assertIn("tilexr_ccu_signal_wait prepare", source)
+        self.assertIn("tilexr_ccu_signal_wait submit", source)
+        self.assertIn("tilexr_ccu_signal_wait timing", source)
+        self.assertIn("tilexr_ccu_signal_wait result passed=1", source)
+        self.assertIn("if (SignalWaitSmokeEnabled() || BarrierSmokeEnabled())", prepared_body)
+
     def test_thread_mode_worker_sets_device_before_direct_ccu_submit(self):
         source = PROBE_SOURCE.read_text(encoding="utf-8")
         thread_mode_body = source[source.index("int RunThreadModeSmoke"): source.index("int main()")]
@@ -460,6 +544,19 @@ def test_probe_defaults_to_lower_layer_first_install_order(self):
             install_order_body.index("TileXRCcuInstallOrder::RepositoryFirst"),
         )
 
+    def test_probe_wires_barrier_mode_env_into_prepare_options(self):
+        source = PROBE_SOURCE.read_text(encoding="utf-8")
+        prepare_options_body = source[
+            source.index("TileXRDirectCcuPrepareOptions MakePrepareOptions"):
+            source.index("const char* SignalWaitRoleName")
+        ]
+
+        self.assertIn("TileXR::TileXRCcuBarrierMode BarrierModeFromEnv()", source)
+        self.assertIn("options.barrierMode = BarrierModeFromEnv();", prepare_options_body)
+        self.assertIn('return TileXR::TileXRCcuBarrierMode::LocalCkePostOnly;', source)
+        self.assertIn('return TileXR::TileXRCcuBarrierMode::SyncCkePostOnly;', source)
+        self.assertIn('return TileXR::TileXRCcuBarrierMode::SyncXnLoadPostOnly;', source)
+
     def test_probe_default_sync_instruction_count_includes_hcomm_style_task1_prelude(self):
         source = PROBE_SOURCE.read_text(encoding="utf-8")
         default_body = source[
diff --git a/tests/ccu/test_tilexr_ccu_direct_smoke_runner.py b/tests/ccu/test_tilexr_ccu_direct_smoke_runner.py
index 85c5be50..c1248231 100644
--- a/tests/ccu/test_tilexr_ccu_direct_smoke_runner.py
+++ b/tests/ccu/test_tilexr_ccu_direct_smoke_runner.py
@@ -141,11 +141,45 @@ def test_runner_p2p_mode_applies_direct_ccu_resource_defaults(self):
         self.assertIn('TILEXR_CCU_PROBE_REMOTE_NOTIFY_CKE_COUNT="${TILEXR_CCU_PROBE_REMOTE_NOTIFY_CKE_COUNT:-8}"', source)
         self.assertIn('TILEXR_CCU_PROBE_CHANNEL_START="${TILEXR_CCU_PROBE_CHANNEL_START:-2}"', source)
         self.assertIn('TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE="${TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE:-ra_ctx}"', source)
+        self.assertIn('TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN="${TILEXR_CCU_DIRECT_SMOKE_FAST_EXIT_AFTER_RUN:-1}"', source)
         self.assertIn('common_env+=("TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE=${TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE}")', source)
         self.assertIn("p2p_passed_count=0", source)
         self.assertIn('grep -q "tilexr_ccu_direct_smoke p2pCcuCopy skipped"', source)
         self.assertIn("direct CCU P2P CCU-copy produced no passing receiver result", source)
 
+    def test_runner_signal_wait_mode_uses_internal_backend_defaults_and_asserts_result(self):
+        source = RUNNER.read_text(encoding="utf-8")
+
+        self.assertIn("apply_signal_wait_defaults", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK", source)
+        self.assertIn("TILEXR_CCU_DIRECT_SMOKE_BARRIER", source)
+        self.assertIn('TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT="${TILEXR_CCU_DIRECT_SMOKE_DIRECT_CCU_ONLY_INIT:-1}"', source)
+        self.assertIn('TILEXR_CCU_PROBE_MISSION_START="${TILEXR_CCU_PROBE_MISSION_START:-6}"', source)
+        self.assertIn('TILEXR_CCU_PROBE_INSTRUCTION_START="${TILEXR_CCU_PROBE_INSTRUCTION_START:-475}"', source)
+        self.assertIn('TILEXR_CCU_PROBE_MISSION_INSTRUCTION_START="${TILEXR_CCU_PROBE_MISSION_INSTRUCTION_START:-489}"', source)
+        self.assertIn('TILEXR_CCU_PROBE_CHANNEL_START="${TILEXR_CCU_PROBE_CHANNEL_START:-2}"', source)
+        self.assertIn('common_env+=("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_WAIT}")', source)
+        self.assertIn('common_env+=("TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK=${TILEXR_CCU_DIRECT_SMOKE_SIGNAL_RANK}")', source)
+        self.assertIn('common_env+=("TILEXR_CCU_DIRECT_SMOKE_BARRIER=${TILEXR_CCU_DIRECT_SMOKE_BARRIER}")', source)
+        self.assertIn('common_env+=("TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT=${TILEXR_CCU_DIRECT_SUBMIT_TIMEOUT}")', source)
+        self.assertIn('grep -q "tilexr_ccu_signal_wait prepare ret=0"', source)
+        self.assertIn('grep -q "tilexr_ccu_signal_wait submit ret=0"', source)
+        self.assertIn('grep -q "tilexr_ccu_signal_wait result passed=1"', source)
+
+    def test_runner_allows_inactive_p2p_rank_to_skip_submit(self):
+        source = RUNNER.read_text(encoding="utf-8")
+        submit_check = source[
+            source.index('if [ "${TILEXR_CCU_DIRECT_SMOKE_SUBMIT:-0}" = "1" ]', source.index('if [ "${rank0_status}"')):
+            source.index('if [ "${TILEXR_CCU_DIRECT_SMOKE_EXPECT_BARRIER_WAIT:-0}" = "1" ]')
+        ]
+
+        self.assertIn("rank_skipped_p2p_ccu_copy_submit", source)
+        self.assertIn('grep -q "tilexr_ccu_direct_smoke p2pCcuCopy skipped"', source)
+        self.assertIn("rank_skipped_p2p_ccu_copy_submit", submit_check)
+        self.assertIn("continue", submit_check)
+        self.assertIn("direct CCU submit did not return success", submit_check)
+
     def test_runner_default_run_skips_without_hardware(self):
         with tempfile.TemporaryDirectory() as temp_dir:
             env = os.environ.copy()
@@ -230,6 +264,15 @@ def test_runner_dry_run_shows_rank_specific_resource_window_token_overrides_comm
         self.assertNotIn("npu-smi rc=", result.stdout + result.stderr)
         self.assertNotIn("ccu_rank0.log", result.stdout)
 
+    def test_runner_passes_rank_specific_resource_window_eid_index(self):
+        source = RUNNER.read_text(encoding="utf-8")
+
+        self.assertIn("EID_INDEX", source[source.index("resource_window_token_fields=("):])
+        self.assertIn('rank0_token_var="TILEXR_CCU_DIRECT_RESOURCE_WINDOW_${token_field}_RANK0"', source)
+        self.assertIn('rank1_token_var="TILEXR_CCU_DIRECT_RESOURCE_WINDOW_${token_field}_RANK1"', source)
+        self.assertIn('echo "dryRun rank0 TILEXR_CCU_DIRECT_RESOURCE_WINDOW_${token_field}=${rank0_token_value}"', source)
+        self.assertIn('echo "dryRun rank1 TILEXR_CCU_DIRECT_RESOURCE_WINDOW_${token_field}=${rank1_token_value}"', source)
+
     def test_runner_dry_run_shows_repository_install_diagnostic_variants(self):
         with tempfile.TemporaryDirectory() as temp_dir:
             env = os.environ.copy()
diff --git a/tests/ccu/test_tilexr_ccu_driver_adapter.py b/tests/ccu/test_tilexr_ccu_driver_adapter.py
index c2b6863d..5244d89e 100644
--- a/tests/ccu/test_tilexr_ccu_driver_adapter.py
+++ b/tests/ccu/test_tilexr_ccu_driver_adapter.py
@@ -950,6 +950,8 @@ def test_driver_adapter_is_wired_and_does_not_reference_hcomm_runtime_surface(se
         self.assertIn("TILEXR_CCU_U_OP_GET_BASIC_INFO", header)
         self.assertIn("TILEXR_CCU_U_OP_GET_DIE_WORKING", header)
         self.assertIn("TILEXR_CCU_U_OP_SET_MSID_TOKEN", header)
+        self.assertIn("TILEXR_CCU_U_OP_SET_TASKKILL", header)
+        self.assertIn("TILEXR_CCU_U_OP_CLEAN_TASKKILL_STATE", header)
         self.assertIn("TILEXR_CCU_U_OP_SET_INSTRUCTION", header)
         self.assertIn("TILEXR_CCU_U_OP_SET_XN", header)
         self.assertIn("TILEXR_CCU_U_OP_SET_CKE", header)
@@ -965,6 +967,8 @@ def test_driver_adapter_is_wired_and_does_not_reference_hcomm_runtime_surface(se
         self.assertIn("GetDieEnabled", header)
         self.assertIn("InstallInstructions", header)
         self.assertIn("InstallMsidToken", header)
+        self.assertIn("CleanTaskKillState", header)
+        self.assertIn("TILEXR_CCU_U_OP_CLEAN_TASKKILL_STATE, &out, report", source)
         self.assertIn("InstallPfeCtx", header)
         self.assertIn("InstallJettyCtx", header)
         self.assertIn("InstallChannelCtxV1", header)
diff --git a/tests/ccu/test_tilexr_ccu_lower_layer_plan_builder.py b/tests/ccu/test_tilexr_ccu_lower_layer_plan_builder.py
index 8cd8c8b8..8db3c238 100644
--- a/tests/ccu/test_tilexr_ccu_lower_layer_plan_builder.py
+++ b/tests/ccu/test_tilexr_ccu_lower_layer_plan_builder.py
@@ -553,6 +553,75 @@ def test_transport_template_can_use_hcomm_fe_id_pfe_partition_for_direct_ccu_exp
         self.assertEqual("", result.stderr)
         self.assertEqual(0, result.returncode, result.stdout + result.stderr)
 
+    def test_transport_template_uses_local_pfe_id_from_ra_eid_func_id(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_lower_layer_plan_builder.h"
+
+            #include <iostream>
+            #include <vector>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                TileXRCcuBasicInfo basic;
+                basic.dieId = 1;
+                basic.msId = 0x45;
+                basic.msidToken.tokenId = 0x1234;
+                basic.msidToken.valid = true;
+
+                TileXRCcuResourceAllocation allocation;
+                allocation.channels = {1, 2, 1};
+                allocation.localXn = {1, 0x120, 1};
+                allocation.remoteXn = {1, 0x240, 1};
+                allocation.notifyCke = {1, 0x330, 1};
+
+                std::vector<TileXRCcuRemoteCcuBufferInfo> remoteCcuBuffers;
+                TileXRCcuRemoteCcuBufferInfo remote;
+                remote.remoteCcuVa = 0x0000009234000000ULL;
+                remote.memoryTokenId = 0x23456;
+                remote.localPfeId = 7;
+                remote.localPfeIdValid = true;
+                remoteCcuBuffers.push_back(remote);
+
+                TileXRCcuLowerLayerTransportSnapshot snapshot;
+                TileXRCcuLowerLayerPlanBuilderReport report;
+                if (TileXRCcuBuildLowerLayerTransportTemplate(
+                        basic, allocation, remoteCcuBuffers, &snapshot, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "template build failed: " << report.message << "\n";
+                    return 1;
+                }
+                if (snapshot.pfeId != 7 || snapshot.pfeOffset != 23) {
+                    std::cerr << "local PFE id not applied: pfeId=" << snapshot.pfeId
+                              << " pfeOffset=" << snapshot.pfeOffset << "\n";
+                    return 2;
+                }
+                return 0;
+            }
+            '''
+        )
+        env = os.environ.copy()
+        env["TILEXR_CCU_DIRECT_LOWER_LAYER_PFE_OFFSET_SOURCE"] = "hcomm_die"
+
+        result = self.compile_and_run(code, env=env)
+
+        self.assertEqual("", result.stderr)
+        self.assertEqual(0, result.returncode, result.stdout + result.stderr)
+
+    def test_transport_template_prefers_ra_eid_func_id_over_channel_id_in_source(self):
+        source = BUILDER_SOURCE.read_text(encoding="utf-8")
+        build_body = source[
+            source.index("int TileXRCcuBuildLowerLayerTransportTemplate"):
+            source.index("int TileXRCcuOverlayVerifiedEndpointRoutes")
+        ]
+
+        self.assertIn("remoteCcuBuffers.front().localPfeId", build_body)
+        self.assertIn("allocation.channels.startId", build_body)
+        self.assertLess(
+            build_body.index("remoteCcuBuffers.front().localPfeId"),
+            build_body.index("SelectLowerLayerPfeOffset"))
+
     def test_builds_lower_layer_install_plan_from_transport_snapshot(self):
         code = textwrap.dedent(
             r'''
@@ -2412,6 +2481,147 @@ def test_direct_ccu_runtime_can_register_resource_window_with_public_ra_ctx(self
         self.assertEqual("", result.stderr)
         self.assertEqual(0, result.returncode, result.stdout + result.stderr)
 
+    def test_direct_runtime_source_supports_selecting_ra_ctx_resource_window_eid(self):
+        source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8")
+
+        self.assertIn("TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX", source)
+        self.assertIn("SelectRaCtxResourceWindowEidInfo", source)
+        self.assertIn("SelectRankedEnv(TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_ENV", source)
+        self.assertIn("TraceRaCtxEidInfos", source)
+        self.assertIn("ctxAttr.ub.eidIndex = selectedEid.eidIndex", source)
+
+    def test_direct_ccu_runtime_can_select_ra_ctx_resource_window_eid_by_env(self):
+        code = textwrap.dedent(
+            r'''
+            #include <cstdint>
+            #include <cstring>
+            #include <iostream>
+
+            #define private public
+            #include "ccu/tilexr_ccu_direct_runtime.h"
+            #undef private
+
+            using namespace TileXR;
+            constexpr uint64_t kResourceAddr = 0x10001234ULL;
+            constexpr uint64_t kAlignedResourceAddr = 0x10001000ULL;
+            constexpr uint64_t kAlignedResourceBytes =
+                TILEXR_CCU_RESOURCE_WINDOW_BYTES + (kResourceAddr - kAlignedResourceAddr);
+
+            int FakeRaCustomChannel(TileXRCcuRaInfo, TileXRCcuCustomChannelIn*, TileXRCcuCustomChannelOut*)
+            {
+                return 0;
+            }
+
+            int FakeRaGetDevEidInfoNum(TileXRCcuRaInfo, uint32_t* num)
+            {
+                *num = 2;
+                return 0;
+            }
+
+            int FakeRaGetDevEidInfoList(TileXRCcuRaInfo, TileXRCcuHccpDevEidInfo* list, uint32_t* num)
+            {
+                if (list == nullptr || num == nullptr || *num != 2) {
+                    return -1;
+                }
+                list[0].eidIndex = 3;
+                list[1].eidIndex = 9;
+                for (uint32_t i = 0; i < TILEXR_CCU_EID_BYTES; ++i) {
+                    list[0].eid.raw[i] = static_cast<uint8_t>(0xa0 + i);
+                    list[1].eid.raw[i] = static_cast<uint8_t>(0xc0 + i);
+                }
+                return 0;
+            }
+
+            int FakeRaCtxInit(TileXRCcuHccpCtxInitCfg*, TileXRCcuHccpCtxInitAttr* attr, void** ctx)
+            {
+                if (attr == nullptr || attr->ub.eidIndex != 9 || attr->ub.eid.raw[0] != 0xc0) {
+                    return -1;
+                }
+                *ctx = reinterpret_cast<void*>(0x1000);
+                return 0;
+            }
+
+            int FakeRaCtxDeinit(void*) { return 0; }
+            int FakeRaCtxTokenIdAlloc(void*, TileXRCcuHccpTokenId* token, void** tokenHandle)
+            {
+                token->tokenId = 0x12345600U;
+                *tokenHandle = reinterpret_cast<void*>(0x2000);
+                return 0;
+            }
+            int FakeRaCtxTokenIdFree(void*, void*) { return 0; }
+            int FakeRaGetSecRandom(TileXRCcuRaInfo*, uint32_t* value)
+            {
+                *value = 0xabcdef01U;
+                return 0;
+            }
+            int FakeRaCtxLmemRegister(void*, TileXRCcuHccpMrRegInfo* mr, void** handle)
+            {
+                if (mr == nullptr || mr->in.mem.addr != kAlignedResourceAddr ||
+                    mr->in.mem.size != kAlignedResourceBytes) {
+                    return -1;
+                }
+                mr->out.ub.tokenId = 0x12345600U;
+                mr->out.ub.targetSegHandle = 0x4455667788ULL;
+                *handle = reinterpret_cast<void*>(0x3000);
+                return 0;
+            }
+            int FakeRaCtxLmemUnregister(void*, void*) { return 0; }
+
+            int main()
+            {
+                TileXRCcuDirectRuntime runtime;
+                runtime.initialized_ = true;
+                runtime.loader_.RaCustomChannel = FakeRaCustomChannel;
+                runtime.loader_.loaded_ = true;
+                runtime.loader_.RaGetDevEidInfoNum = FakeRaGetDevEidInfoNum;
+                runtime.loader_.RaGetDevEidInfoList = FakeRaGetDevEidInfoList;
+                runtime.loader_.RaCtxInit = FakeRaCtxInit;
+                runtime.loader_.RaCtxDeinit = FakeRaCtxDeinit;
+                runtime.loader_.RaCtxTokenIdAlloc = FakeRaCtxTokenIdAlloc;
+                runtime.loader_.RaCtxTokenIdFree = FakeRaCtxTokenIdFree;
+                runtime.loader_.RaGetSecRandom = FakeRaGetSecRandom;
+                runtime.loader_.RaCtxLmemRegister = FakeRaCtxLmemRegister;
+                runtime.loader_.RaCtxLmemUnregister = FakeRaCtxLmemUnregister;
+                runtime.devicePhyId_ = 0x55;
+                runtime.options_.rank = 1;
+                runtime.cachedBasicInfoValid_ = true;
+                runtime.cachedBasicInfo_.resourceAddr = kResourceAddr;
+                runtime.cachedBasicInfo_.caps.cap1 = 7U << 16U;
+
+                if (runtime.RegisterCcuResourceRmaBuffer(kResourceAddr) != TILEXR_SUCCESS) {
+                    std::cerr << "ra ctx resource window register failed\n";
+                    return 1;
+                }
+                TileXRCcuLocalResourceWindowInfo local;
+                if (runtime.ExportLocalCcuRmaBuffer(&local) != TILEXR_SUCCESS ||
+                    local.eidIndex != 9 || local.eid[0] != 0xc0) {
+                    std::cerr << "wrong eid selection index=" << local.eidIndex
+                              << " first=" << static_cast<uint32_t>(local.eid[0]) << "\n";
+                    return 2;
+                }
+                runtime.Shutdown();
+                return 0;
+            }
+            '''
+        )
+        env = os.environ.copy()
+        env["TILEXR_CCU_DIRECT_RESOURCE_WINDOW_REGISTRATION_MODE"] = "ra_ctx"
+        env["TILEXR_CCU_DIRECT_RESOURCE_WINDOW_EID_INDEX_RANK1"] = "9"
+
+        result = self.compile_and_run(
+            code,
+            env=env,
+            extra_sources=[
+                DIRECT_RUNTIME_SOURCE,
+                REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_hccp_loader.cpp",
+                REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_ra_custom_channel_provider.cpp",
+                REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_driver_adapter.cpp",
+            ],
+            extra_link_flags=["-ldl"])
+
+        self.assertEqual("", result.stderr)
+        self.assertEqual(0, result.returncode, result.stdout + result.stderr)
+
     def test_direct_ccu_runtime_collects_ra_ctx_endpoint_route_when_resource_window_uses_ra_ctx(self):
         code = textwrap.dedent(
             r'''
@@ -3703,11 +3913,11 @@ def test_tilexr_comm_prepares_direct_ccu_install_attempt_without_submitting(self
         ]:
             self.assertNotIn(forbidden, combined)
 
-    def test_tilexr_comm_direct_ccu_prepare_fails_fast_after_process_init_failure(self):
+    def test_tilexr_comm_direct_ccu_prepare_fails_fast_after_device_init_failure(self):
         runtime_session_source = CCU_RUNTIME_SESSION_SOURCE.read_text(encoding="utf-8")
         planner_source = CCU_PLANNER_SOURCE.read_text(encoding="utf-8")
 
-        self.assertIn("g_ccuDirectRuntimeUnavailableMessage", runtime_session_source)
+        self.assertIn("g_ccuDirectRuntimeUnavailableByDevice", runtime_session_source)
         init_body = runtime_session_source[
             runtime_session_source.index("int TileXRCcuRuntimeSession::Init("):
             runtime_session_source.index("void TileXRCcuRuntimeSession::ResetDirectCcuBasicInfo")
@@ -3717,11 +3927,11 @@ def test_tilexr_comm_direct_ccu_prepare_fails_fast_after_process_init_failure(se
             planner_source.index("int TileXRCcuCollectivePlanner::PrepareDirectCcuMemoryCopyInstallAttempt")
         ]
 
-        self.assertIn("g_ccuDirectRuntimeUnavailableMessage = runtimeReport.message", init_body)
-        self.assertIn("direct CCU runtime unavailable after process-level init failure", runtime_session_source)
-        self.assertIn("ProcessDirectCcuRuntimeUnavailableMessage()", prepare_body)
+        self.assertIn("g_ccuDirectRuntimeUnavailableByDevice[devId_] = runtimeReport.message", init_body)
+        self.assertIn("direct CCU runtime unavailable for device", runtime_session_source)
+        self.assertIn("session.DirectCcuRuntimeUnavailableMessage()", prepare_body)
         self.assertLess(
-            prepare_body.index("ProcessDirectCcuRuntimeUnavailableMessage()"),
+            prepare_body.index("session.DirectCcuRuntimeUnavailableMessage()"),
             prepare_body.index("RefreshDirectCcuBasicInfo(installDieId)"),
         )
 
@@ -3735,7 +3945,7 @@ def test_tilexr_comm_direct_ccu_runtime_init_serializes_ra_initialization(self):
         lock_pos = init_body.index("lock_guard<mutex> lock(g_ccuDirectRuntimeMtx);")
         allocation_pos = init_body.index("ccuDirectRuntime_.reset(new (nothrow) TileXRCcuDirectRuntime())")
         runtime_init_pos = init_body.index("ccuDirectRuntime_->Init(runtimeOptions, &runtimeReport)")
-        unavailable_set_pos = init_body.index("g_ccuDirectRuntimeUnavailable = true")
+        unavailable_set_pos = init_body.index("g_ccuDirectRuntimeUnavailableByDevice[devId_] = runtimeReport.message")
 
         self.assertLess(lock_pos, allocation_pos)
         self.assertLess(allocation_pos, runtime_init_pos)
@@ -3755,17 +3965,17 @@ def test_tilexr_comm_direct_ccu_prepare_can_select_install_die_for_diagnostics(s
         self.assertIn("basicInfo->dieId != installDieId", prepare_body)
         self.assertNotIn("RefreshDirectCcuBasicInfo(0)", prepare_body)
 
-    def test_tilexr_comm_direct_ccu_thread_allgather_aborts_after_process_init_failure(self):
+    def test_tilexr_comm_direct_ccu_thread_allgather_aborts_after_device_init_failure(self):
         runtime_session_source = CCU_RUNTIME_SESSION_SOURCE.read_text(encoding="utf-8")
         thread_allgather_body = runtime_session_source[
             runtime_session_source.index("int TileXRCcuRuntimeSession::DirectCcuThreadAllGather"):
             runtime_session_source.index("} // namespace TileXR")
         ]
 
-        self.assertIn("ProcessDirectCcuRuntimeUnavailableMessage()", thread_allgather_body)
+        self.assertIn("DirectCcuRuntimeUnavailableMessage()", thread_allgather_body)
         self.assertIn("direct CCU thread allgather abort", thread_allgather_body)
         self.assertLess(
-            thread_allgather_body.index("ProcessDirectCcuRuntimeUnavailableMessage()"),
+            thread_allgather_body.index("DirectCcuRuntimeUnavailableMessage()"),
             thread_allgather_body.index("TILEXR_INIT_TIMEOUT"),
         )
 
diff --git a/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py b/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py
index 36c2068c..f03bfc52 100644
--- a/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py
+++ b/tests/ccu/test_tilexr_ccu_ra_custom_channel_loader.py
@@ -1262,6 +1262,17 @@ def test_direct_ccu_runtime_queries_basic_info_through_ccu_hccp_loader(self):
             with self.subTest(needle=needle):
                 self.assertNotIn(needle, combined)
 
+    def test_direct_ccu_runtime_init_cleans_sticky_taskkill_state(self):
+        source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8")
+
+        init_body = source[
+            source.index("int TileXRCcuDirectRuntime::Init("):
+            source.index("void TileXRCcuDirectRuntime::Shutdown()")
+        ]
+        self.assertIn("adapter.CleanTaskKillState", init_body)
+        self.assertIn("TILEXR_CCU_DIRECT_DEFAULT_DIE_ID", init_body)
+        self.assertIn("TraceTaskKillCleanup", source)
+
     def test_direct_ccu_runtime_keeps_ra_custom_channel_provider_alive_for_created_adapters(self):
         header = DIRECT_RUNTIME_HEADER.read_text(encoding="utf-8")
         source = DIRECT_RUNTIME_SOURCE.read_text(encoding="utf-8")
diff --git a/tests/ccu/test_tilexr_ccu_resource_allocator.py b/tests/ccu/test_tilexr_ccu_resource_allocator.py
index 17c105f4..b9f2ca28 100644
--- a/tests/ccu/test_tilexr_ccu_resource_allocator.py
+++ b/tests/ccu/test_tilexr_ccu_resource_allocator.py
@@ -751,6 +751,103 @@ def test_allocator_rejects_resource_exhaustion_and_double_release(self):
         self.assertEqual("", result.stderr)
         self.assertEqual(0, result.returncode, result.stdout + result.stderr)
 
+    def test_allocator_release_reclaims_latest_receipt_and_rejects_non_lifo_release(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_resource_allocator.h"
+
+            #include <iostream>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                TileXRCcuResourceSpec spec;
+                spec.dieId = 0;
+                spec.missionKey = 0x12345678U;
+                spec.missionStartId = 10;
+                spec.missionCount = 4;
+                spec.instructionStartId = 100;
+                spec.instructionCount = 64;
+                spec.xnStartId = 200;
+                spec.xnCount = 16;
+                spec.ckeStartId = 300;
+                spec.ckeCount = 8;
+                spec.channelStartId = 20;
+                spec.channelCount = 4;
+
+                TileXRCcuResourceRequest request;
+                request.sqeArgCount = 0;
+                request.syncResourceCount = 1;
+                request.syncInstructionCount = 2;
+                request.bindingsPerSyncResource = 1;
+
+                TileXRCcuResourceAllocator allocator;
+                if (allocator.Init(spec) != TILEXR_SUCCESS) {
+                    std::cerr << "init failed\n";
+                    return 1;
+                }
+
+                TileXRCcuProducerPlan firstPlan;
+                TileXRCcuProducerPlan secondPlan;
+                TileXRCcuProducerPlan thirdPlan;
+                TileXRCcuResourceAllocation first;
+                TileXRCcuResourceAllocation second;
+                TileXRCcuResourceAllocation third;
+                TileXRCcuResourceAllocatorReport report;
+                if (allocator.Allocate(request, &firstPlan, &first, &report) != TILEXR_SUCCESS ||
+                    allocator.Allocate(request, &secondPlan, &second, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "initial allocate failed: " << report.message << "\n";
+                    return 2;
+                }
+                if (first.mission.startId != 10 || second.mission.startId != 11 ||
+                    first.repository.startId != 100 || second.repository.startId != 102 ||
+                    first.localXn.startId == second.localXn.startId ||
+                    first.channels.startId == second.channels.startId) {
+                    std::cerr << "initial allocations did not advance resource cursors\n";
+                    return 3;
+                }
+                if (allocator.Release(first.receiptId) != TILEXR_ERROR_PARA_CHECK_FAIL) {
+                    std::cerr << "non-LIFO release was accepted\n";
+                    return 4;
+                }
+                if (allocator.Release(second.receiptId) != TILEXR_SUCCESS) {
+                    std::cerr << "latest release failed\n";
+                    return 5;
+                }
+                if (allocator.Allocate(request, &thirdPlan, &third, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "third allocate failed after release: " << report.message << "\n";
+                    return 6;
+                }
+                if (third.mission.startId != second.mission.startId ||
+                    third.repository.startId != second.repository.startId ||
+                    third.localXn.startId != second.localXn.startId ||
+                    third.remoteXn.startId != second.remoteXn.startId ||
+                    third.localWaitCke.startId != second.localWaitCke.startId ||
+                    third.remoteNotifyCke.startId != second.remoteNotifyCke.startId ||
+                    third.channels.startId != second.channels.startId) {
+                    std::cerr << "released resource ranges were not reused\n";
+                    return 7;
+                }
+                if (allocator.Release(first.receiptId) != TILEXR_ERROR_PARA_CHECK_FAIL) {
+                    std::cerr << "old non-tail receipt release should still be rejected\n";
+                    return 8;
+                }
+                if (allocator.Release(third.receiptId) != TILEXR_SUCCESS ||
+                    allocator.Release(first.receiptId) != TILEXR_SUCCESS) {
+                    std::cerr << "tail releases failed\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_allocator_rejects_sync_instruction_window_too_small_for_barrier_program(self):
         code = textwrap.dedent(
             r'''
diff --git a/tests/ccu/test_tilexr_ccu_signal_wait_program.py b/tests/ccu/test_tilexr_ccu_signal_wait_program.py
new file mode 100644
index 00000000..f008749b
--- /dev/null
+++ b/tests/ccu/test_tilexr_ccu_signal_wait_program.py
@@ -0,0 +1,264 @@
+#!/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]
+SIGNAL_WAIT_HEADER = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_signal_wait_program.h"
+SIGNAL_WAIT_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_signal_wait_program.cpp"
+BARRIER_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_barrier_program.cpp"
+MICROCODE_SOURCE = REPO_ROOT / "src" / "comm" / "ccu" / "tilexr_ccu_microcode.cpp"
+COMM_CMAKE = REPO_ROOT / "src" / "comm" / "CMakeLists.txt"
+INCLUDE_DIR = REPO_ROOT / "src" / "include"
+COMM_DIR = REPO_ROOT / "src" / "comm"
+
+
+class TileXRCcuSignalWaitProgramTest(unittest.TestCase):
+    def compile_and_run(self, code: str):
+        compiler = shutil.which("g++") or shutil.which("clang++") or shutil.which("c++")
+        if compiler is None:
+            self.skipTest("no local C++ compiler found; remote CANN compile covers signal/wait C++ syntax")
+        with tempfile.TemporaryDirectory() as temp_dir:
+            temp_path = Path(temp_dir)
+            test_cpp = temp_path / "signal_wait_program_test.cpp"
+            test_bin = temp_path / "signal_wait_program_test"
+            test_cpp.write_text(code, encoding="utf-8")
+            subprocess.run(
+                [
+                    compiler,
+                    "-std=c++14",
+                    "-I",
+                    str(INCLUDE_DIR),
+                    "-I",
+                    str(COMM_DIR),
+                    str(test_cpp),
+                    str(SIGNAL_WAIT_SOURCE),
+                    str(BARRIER_SOURCE),
+                    str(MICROCODE_SOURCE),
+                    "-o",
+                    str(test_bin),
+                ],
+                cwd=REPO_ROOT,
+                check=True,
+                text=True,
+                capture_output=True,
+            )
+            return subprocess.run([str(test_bin)], cwd=REPO_ROOT, check=False, text=True, capture_output=True)
+
+    def test_signal_wait_program_initializes_source_cke_before_signal_only_post(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_signal_wait_program.h"
+
+            #include <iostream>
+            #include <vector>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                TileXRCcuSignalWaitProgramSpec spec;
+                spec.role = TileXRCcuSignalWaitProgramRole::Signal;
+                spec.channelId = 2;
+                spec.remoteXn = 2361;
+                spec.localXn = 1961;
+                spec.remoteNotifyCke = 364;
+                spec.remoteNotifyMask = 1;
+                spec.sourceCke = 0x101;
+                spec.sourceCkeMask = 0xffff;
+
+                std::vector<TileXRCcuInstr> program;
+                TileXRCcuBarrierProgramReport report;
+                if (TileXRCcuBuildSignalWaitProgram(spec, &program, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "signal program build failed: " << report.message << "\n";
+                    return 1;
+                }
+                if (program.size() != 5 || report.postInstructionCount != 1 ||
+                    report.waitInstructionCount != 0 || report.totalInstructionCount != 5) {
+                    std::cerr << "unexpected signal report\n";
+                    return 2;
+                }
+                if (program[2].words[0] != 0xffff010100010802ULL ||
+                    program[3].words[0] != 0xffff0101016c100bULL ||
+                    program[3].words[1] != 0x0000000000000002ULL ||
+                    program[3].words[2] != 0x0001000000000000ULL) {
+                    std::cerr << "unexpected signal instructions\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_signal_wait_program_builds_wait_only_clear_cke_wait(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_signal_wait_program.h"
+
+            #include <iostream>
+            #include <vector>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                TileXRCcuSignalWaitProgramSpec spec;
+                spec.role = TileXRCcuSignalWaitProgramRole::Wait;
+                spec.localXn = 0x7a9;
+                spec.localWaitCke = 0x220;
+                spec.localWaitMask = 1;
+
+                std::vector<TileXRCcuInstr> program;
+                TileXRCcuBarrierProgramReport report;
+                if (TileXRCcuBuildSignalWaitProgram(spec, &program, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "wait program build failed: " << report.message << "\n";
+                    return 1;
+                }
+                if (program.size() != 1 || report.postInstructionCount != 0 ||
+                    report.waitInstructionCount != 1 || report.totalInstructionCount != 1) {
+                    std::cerr << "unexpected wait report\n";
+                    return 2;
+                }
+                if (program[0].words[0] != 0x0000000000010804ULL ||
+                    program[0].words[1] != 0x0000000000010220ULL) {
+                    std::cerr << "unexpected wait instruction\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_signal_wait_program_builds_signal_and_wait_for_barrier(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_signal_wait_program.h"
+
+            #include <iostream>
+            #include <vector>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                TileXRCcuSignalWaitProgramSpec spec;
+                spec.role = TileXRCcuSignalWaitProgramRole::SignalAndWait;
+                spec.channelId = 2;
+                spec.remoteXn = 2361;
+                spec.localXn = 1961;
+                spec.remoteNotifyCke = 364;
+                spec.remoteNotifyMask = 1;
+                spec.localWaitCke = 0x220;
+                spec.localWaitMask = 1;
+                spec.sourceCke = 0x101;
+                spec.sourceCkeMask = 0xffff;
+
+                std::vector<TileXRCcuInstr> program;
+                TileXRCcuBarrierProgramReport report;
+                if (TileXRCcuBuildSignalWaitProgram(spec, &program, &report) != TILEXR_SUCCESS) {
+                    std::cerr << "barrier program build failed: " << report.message << "\n";
+                    return 1;
+                }
+                if (program.size() != 6 || report.postInstructionCount != 1 ||
+                    report.waitInstructionCount != 1 || report.totalInstructionCount != 6) {
+                    std::cerr << "unexpected signal_and_wait report\n";
+                    return 2;
+                }
+                if (program[2].words[0] != 0xffff010100010802ULL ||
+                    program[3].words[0] != 0xffff0101016c100bULL ||
+                    program[4].words[0] != 0x0000000000010804ULL ||
+                    program[4].words[1] != 0x0000000000010220ULL) {
+                    std::cerr << "unexpected signal_and_wait instructions\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_signal_wait_program_rejects_missing_role_resources(self):
+        code = textwrap.dedent(
+            r'''
+            #include "ccu/tilexr_ccu_signal_wait_program.h"
+
+            #include <iostream>
+            #include <vector>
+
+            using namespace TileXR;
+
+            int main()
+            {
+                std::vector<TileXRCcuInstr> program(1);
+                TileXRCcuBarrierProgramReport report;
+                TileXRCcuSignalWaitProgramSpec spec;
+                spec.role = TileXRCcuSignalWaitProgramRole::Signal;
+                if (TileXRCcuBuildSignalWaitProgram(spec, &program, &report) !=
+                    TILEXR_ERROR_PARA_CHECK_FAIL) {
+                    std::cerr << "missing signal resources accepted\n";
+                    return 1;
+                }
+                if (!program.empty() || report.message.find("signal") == std::string::npos) {
+                    std::cerr << "weak signal diagnostic: " << report.message << "\n";
+                    return 2;
+                }
+                spec.role = TileXRCcuSignalWaitProgramRole::Wait;
+                if (TileXRCcuBuildSignalWaitProgram(spec, &program, &report) !=
+                    TILEXR_ERROR_PARA_CHECK_FAIL) {
+                    std::cerr << "missing wait resources accepted\n";
+                    return 3;
+                }
+                if (!program.empty() || report.message.find("wait") == std::string::npos) {
+                    std::cerr << "weak wait diagnostic: " << report.message << "\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_signal_wait_program_is_wired(self):
+        cmake = COMM_CMAKE.read_text(encoding="utf-8")
+        header = SIGNAL_WAIT_HEADER.read_text(encoding="utf-8")
+        source = SIGNAL_WAIT_SOURCE.read_text(encoding="utf-8")
+
+        self.assertIn("ccu/tilexr_ccu_signal_wait_program.h", cmake)
+        self.assertIn("ccu/tilexr_ccu_signal_wait_program.cpp", cmake)
+        self.assertIn("enum class TileXRCcuSignalWaitProgramRole", header)
+        self.assertIn("struct TileXRCcuSignalWaitProgramSpec", header)
+        self.assertIn("TileXRCcuBuildSignalWaitProgram", header)
+        self.assertIn("TileXRCcuEncodeLoadImdToXn", source)
+        self.assertIn("TileXRCcuEncodeSetCke", source)
+        self.assertIn("TileXRCcuEncodeSyncCke", source)
+        self.assertIn("TileXRCcuEncodeClearCke", source)
+
+
+if __name__ == "__main__":
+    unittest.main()