[Metal:Feature] Add context release API and fix IDST buffer ownership

GitOrigin-RevId: 3c93aa7573c0d10a45028ab1a5b1095858360bf3

Author: Tangesion

include/MNN/MNNSharedContext.h

@@ -52,6 +52,7 @@ struct MNNMetalTensorContent {
 };
 
 MNN_PUBLIC int MNNMetalGetTensorContent(MNNMetalTensorContent* content, void* tensor);
+MNN_PUBLIC int MNNMetalReleaseContext(void);
 #endif
 
 #ifdef MNN_USER_SET_DEVICE

skills/general-debug/SKILL.md

@@ -17,7 +17,7 @@ description: MNN 各类正确性/回归 bug 的排查方法论集合。按 bug 
 
 | # | 类别 | 典型症状 | 章节 |
 |---|------|---------|------|
-| 1 | **内存别名 / 生命周期** | 数值错乱、乱码 token、NaN；代码逻辑看着正确、指针地址合法；换后端结果不同；关掉某优化就好；单/多线程行为差异 | [§1](#1-内存别名--生命周期错误) |
+| 1 | **内存别名 / 生命周期** | 数值错乱、乱码 token、NaN；反复创建/销毁后物理内存线性增长；代码逻辑看着正确、指针地址合法；换后端结果不同；关掉某优化就好；单/多线程行为差异 | [§1](#1-内存别名--生命周期错误) |
 | 2 | **量化误差 / 导出侧权重损坏** | 低 bit（Q4）下输出乱码或退化、Q8/更高 bit 正常；**所有推理后端都错**（CPU/Metal 一致地错）；torch 侧 `--test` 正常但 MNN 推理错；只有大模型/大 vocab 触发 | [§2](#2-量化误差--导出侧权重损坏) |
 | 3 | **并发 / 线程竞争**（GPU 侧已覆盖） | 结果不稳定、每次运行不同；单线程稳定但多线程随机错。**GPU 上逐次不同**已有专门案例（融合引入的别名竞争 + 逐 op commit 二分法） | [§1.6](#16-参考案例融合引入的别名竞争layernorm-折进-conv1x12026-08-03)；CPU 多线程分片类仍待补 |
 | 4 | *（待补：图优化回归）* | 某个 converter pass 之后模型跑错，disable 该 pass 后正常 | *待补* |
@@ -52,6 +52,7 @@ description: MNN 各类正确性/回归 bug 的排查方法论集合。按 bug 
 - 关掉某个新加的 op / 优化 pass 就好，看着又没写错逻辑；
 - `onResize` / buffer 分配相关代码改过后开始回归；
 - 加 `printf` 或改 buffer size 就"好了"。
+- 同一 workload 反复创建/销毁后，物理内存按固定步长增长，且常规 backend GC 无效。
 
 ### 1.1 核心心法
 
@@ -265,7 +266,20 @@ re-home 失败则**保守跳过该次融合**（fail-safe 方向）。
   「fp32 bit-identical + fp16 确定性 + 质量/回归」，并在文档里明确它与「byte-identical」的差别，
   不要含糊带过。
 
-### 1.8 相关文件索引
+### 1.8 物理内存增长：先按 VM 区域和分配栈归因
+
+GPU workload 的 `phys_footprint` 增长不等于 GPU 资源泄漏。若 backend 自报的活跃分配已回落，
+但进程物理内存仍线性增长，应先用系统 VM/heap 工具回答“活着的是哪类区域、谁分配的”，再改
+context、allocator 或缓存策略：
+
+1. 在每轮完整 teardown 后同时记录进程物理内存和 backend 活跃分配量；两者分离时不要继续只查 GPU。
+2. 用 VM 区域汇总区分 IOAccelerator/Metal 映射与 `MALLOC_LARGE` 等 CPU heap；用 heap 大小分布找出
+   与每轮增量匹配的重复 allocation。
+3. 对一个代表性活块抓分配回溯，沿栈检查返回缓冲的所有权。调用方提供了外部 buffer，**不代表**
+   callee 返回的指针一定与它别名；所有权判断必须基于实际返回指针或显式契约，而不能只看“曾传入非空指针”。
+4. 修复后既要验证多轮 teardown 平台化，也要补覆盖“返回外部 buffer”和“返回新临时 buffer”两条路径的单测。
+
+### 1.9 相关文件索引
 
 | 文件 | 作用 |
 |------|------|

skills/test-ci/ios-llm-bench.md

@@ -72,12 +72,15 @@ xcrun devicectl device info processes --device <id>
 ## 已知陷阱
 
 - **shell 环境的 `SDKROOT` / `CPATH` 指向 MacOSX.sdk 会打爆整个 iOS 编译**（2026-07-30 实锤）：症状是几百个 `<cstddef> tried including <stddef.h> but didn't find libc++'s <stddef.h>`，libc++ 头来自 iPhoneOS.sdk 而 C 头来自 MacOSX.sdk。且 `buildiOS.sh` 失败后仍 `exit 0`，只留下残缺 framework（仅 Headers + Info.plist，无二进制、`Headers/llm` 为空），下游 App 编译报 `MNN/llm/llm.hpp not found` 误导排查方向。**处置**：跑本脚本一律 `env -u SDKROOT -u CPATH DEVELOPER_DIR=/Applications/Xcode.app/Contents/Developer sh ios_llm_bench.sh ...`；怀疑 framework 残缺时先 `ls MNN.framework/MNN` 确认二进制存在。
-- **Team ID 必须查本机，不能抄文档**：`--team` 用错会报 `No Account for Team`。查法：`security find-identity -v -p codesigning` 看证书，或 `security cms -D -i ~/Library/Developer/Xcode/UserData/Provisioning\ Profiles/*.mobileprovision | grep -A2 TeamIdentifier`。本机（jbyang，2026-07-30）为 `3TLG5LZ643`。
+- **Team ID 必须查本机，不能抄文档**：`--team` 用错会报 `No Account for Team`。证书名称末尾括号中的值不一定是 provisioning 的 Team ID；应以实际 profile 的 `TeamIdentifier` / `Entitlements.com.apple.developer.team-identifier` 为准：`security cms -D -i <profile.mobileprovision> | plutil -p -`。Automatic signing 在 Xcode 未登录账号时仍可能复用缓存的 Xcode-managed profile，但 bundle id 和 Team 必须与该 profile 完全匹配。
 - **Personal Team 可能没有可列出的 codesign identity**：`security find-identity` 返回 0 时，可从 `defaults read com.apple.dt.Xcode` 的 `IDEProvisioningTeamByIdentifier` 找当前 Personal Team，再用 `codesign -dvvv <app>` 核对实际 `TeamIdentifier`。复用已安装的 bundle id 时，新旧 App 的 Team 也必须一致，否则会报 `MismatchedApplicationIdentifierEntitlement`。
 - **免费开发者证书每台设备最多 3 个 App**：安装报 `CoreDeviceError 3002` + `maximum number of installed apps using a free developer profile`，错误信息会列出占位的 3 个 bundle id。**处置**：优先用 `--bundle-id` 复用其中同 Team 的旧 bench App 原地覆盖（如 `com.jiuqi.mnn-llm-bench`），不必删设备上的 App。
 - **schema 变更后的旧模型不能用于性能对比**：旧 FlatBuffer 可能不会给出清晰的“版本不兼容”错误，而是在加载阶段表现为 `std::bad_alloc` / `SIGABRT` / `SIGSEGV`。若多个分支都未进入 `[MNN_BENCH] run=...` 就崩溃，先用当前 schema 和 `MNNConvert` 重导模型，并检查 `export_args.json` 及 `llm.mnn.json` 的融合 Op，不要将加载崩溃误判为性能回归。
 - **签名未信任的启动失败是"秒失败"，但脚本会傻等满 `--timeout`（默认 1800s）**：换新 bundle id / 新 profile 首次安装后必须在 iPad 上手动信任（设置 > 通用 > VPN与设备管理 > 开发者App）。日志特征：`FBSOpenApplicationErrorDomain error 3` + `its profile has not been explicitly trusted by the user`。看到 TIMEOUT 先翻 `bench_logs/*.log` 头部有没有这个错误，别真等 30 分钟。
 - **设备锁屏**：锁屏时 `devicectl` 无法启动 App（FBSOpenApplicationErrorDomain error 7 "Locked"），脚本会立即报 `app failed to launch (device locked?)`。测试前保持屏幕解锁（建议 设置 > 显示与亮度 > 自动锁定 设为"永不"）。
+- **`devicectl --log-output` 不等于 App stdout 归档**：`device process launch --console --log-output` 的文件可能只包含 CoreDevice 自身日志，终端里看到的 App 输出未必写入该文件。长测和内存复现必须把关键 sample/done/error 标记同步写进 App 的 Documents，再用 `device copy from --domain-type appDataContainer --domain-identifier <bundle-id>` 拉回；console 只用于实时观察。
+- **先验证单轮峰值，再开始多轮泄漏测试**：普通开发签名下，大模型可能在第一轮 session 创建时就被 iOS `SIGKILL(9)`，此时没有资格用“六轮未增长”判断修复有效。先做一轮 smoke 并记录最后成功阶段；若设备额度无法承载完整模型，只能另加一个明确经过目标分配/所有权路径的最小 A/B 压测，并在报告中把它与“完整模型未跑通”分开陈述。
+- **Diffusion framework 还需要 tokenizer 支持**：即使只跑 SD1.5，iOS framework 也要同时启用 `MNN_BUILD_DIFFUSION=ON`、`MNN_BUILD_OPENCV=ON`、`MNN_IMGCODECS=ON` 和 `MNN_BUILD_LLM=ON`；缺少最后一项时 `tokenizer.mtok` 会在 `Diffusion::load()` 直接失败，尚未进入被测推理路径。
 - **iOS 26.5 Metal4 Tensor API 探测（本 skill 相关 bugfix）**：MPP `matmul2d` 要求 M/N 至少一个是 16 的倍数、静态 K 是 16 的倍数。探测 kernel 描述符需用 `(16, 8, dynamic_extent)`；同时 `MetalAttentionShader.hpp` 中 legacy 16x16x8 tensor 路径（静态 K=8）必须保持禁用（宏 `MNN_METAL_TENSOR_OPS_LEGACY_8X8`），否则探测通过但运行时反复编译失败，prefill 反而大幅回退（953 → 717 tok/s）。完整修复后 tensor API 生效，prefill 953 → 1884 tok/s（Qwen3.5-2B，prompt=512）。
 - **GPU 开关**：通过 `devicectl` 启动时 App 处于 Inactive 状态，Metal backend 若在此时创建，必须监听 `UIApplicationDidBecomeActiveNotification`（而非 WillEnterForeground）才能恢复 GPU，否则 bench 卡死。
 - **xcode-select 指向 CommandLineTools**：cmake iOS toolchain 会报 `get_filename_component` 错误；脚本已自动设置 `DEVELOPER_DIR`，手动编译时需 `export DEVELOPER_DIR=/Applications/Xcode.app/Contents/Developer`。

source/backend/metal/MetalBackend.mm

@@ -1684,6 +1684,7 @@ static void _execute(id<MTLComputeCommandEncoder> encoder, const MetalBackend::C
     std::mutex pLock;
     MNNMetalContext* pContext;
     id<MTLDevice> pDevice;
+    size_t runtimeCount = 0;
 };
 static MetalContext* gContext = nullptr;
 MetalRuntime* MetalRuntime::create(const Backend::Info& info) {
@@ -1711,6 +1712,7 @@ static void _execute(id<MTLComputeCommandEncoder> encoder, const MetalBackend::C
     }
     auto mContext = (__bridge_retained void *)(gContext->pContext);
     auto rt = new MetalRuntime(mContext);
+    ++gContext->runtimeCount;
     rt->setGpuMode(info.gpuMode);
     if (nil != sharedContext.queue) {
         rt->setCommandQueue(sharedContext.queue, true);
@@ -1795,6 +1797,11 @@ static void _execute(id<MTLComputeCommandEncoder> encoder, const MetalBackend::C
 }
 
 MetalRuntime::~ MetalRuntime() {
+    {
+        std::unique_lock<std::mutex> _l(gContext->pLock);
+        MNN_ASSERT(gContext->runtimeCount > 0);
+        --gContext->runtimeCount;
+    }
     if(mContext) {
         CFRelease(mContext);
     }
@@ -2098,6 +2105,19 @@ virtual void onRelease(MemChunk chunk) override {
     id<MTLDevice> mDevice;
 };
 
+static int releaseMetalContext() {
+    if (gContext == nullptr) {
+        return 0;
+    }
+    std::unique_lock<std::mutex> lock(gContext->pLock);
+    if (gContext->runtimeCount != 0) {
+        return 1;
+    }
+    gContext->pContext = nil;
+    gContext->pDevice = nil;
+    return 0;
+}
+
 void registerMetalRuntimeCreator() {
     // according to
     // https://developer.apple.com/library/archive/documentation/DeviceInformation/Reference/iOSDeviceCompatibility/HardwareGPUInformation/HardwareGPUInformation.html
@@ -2117,6 +2137,10 @@ void registerMetalRuntimeCreator() {
     }
 }
 } // namespace MNN
+
+extern "C" int MNNMetalReleaseContext(void) {
+    return MNN::releaseMetalContext();
+}
 #else
 namespace MNN {
 void registerMetalRuntimeCreator() {
@@ -2125,5 +2149,8 @@ void registerMetalRuntimeCreator() {
 int MNNMetalGetTensorContent(MNNMetalTensorContent* content, void* tensor) {
     return -1;
 }
+extern "C" int MNNMetalReleaseContext(void) {
+    return 0;
+}
 
 #endif

source/core/ConvolutionCommon.cpp

@@ -747,7 +747,9 @@ std::shared_ptr<ConvolutionCommon::Int8Common> ConvolutionCommon::load(const Op*
                 MNN_PRINT("Alloc memory error for extract idst int8\n");
                 return nullptr;
             }
-            if(weightPtr != nullptr) {
+            // Low-bit decoding may return a temporary buffer even when the caller supplied
+            // weightPtr. Only the caller-owned buffer must be marked non-owning.
+            if (buffer == weightPtr) {
                 result->weight.set(buffer, false);
             } else {
                 result->weight.set(buffer, (int)weightLength);