最近在调一个多 GPU 推理服务的性能被 CPU 到 GPU 的一次次 launch 延迟搞得有点上头。任务量不大单次 kernel 可能只要几十微秒可 CPU 下发、同步、再下发的固定开销硬生生把 GPU 利用率压到了三成以下。翻 Nsight Systems 的时间线一眼看过去全是深色的空白段——GPU 在空等 CPU。于是我认真研究了一圈怎么让 GPU 的工作流尽量少依赖 CPU最后落到三个关键机制上CUDA Graphs、原子队列和 NVSHMEM。这篇笔记就是把这三者的原理、用法和组合实战做个完整记录。1. 为什么我要“让 GPU 自己调度”——launch 开销那本账1.1 一个让我无法忽视的 GPU 空闲时间线先说我当时的场景一个多 GPU 的推理服务每个请求拆分到 GPU 上就是几个小 kernel比如 embedding、attention 的一小段、FFN 的一小段。单个 kernel 的 GPU 执行时间在 10-50 微秒这个量级听起来很快但整条流水线跑下来吞吐就是上不去。用nsys profile抓了一条时间线后我发现 GPU 的 SM 大部分时间处于 idle。kernel 和 kernel 之间有一条很明显的空隙那是 CPU 在完成上一次 launch 的参数提交之后GPU 执行完当前 kernel再等待下一个命令到达的间隔。这个间隔在 ncu 的 LaunchStats 里体现为 launch 相关的 overhead在小 kernel 场景下能占到整个执行时间的 30% 到 50%。这是典型的 CPU 驱动调度模型的问题GPU 自己不产生工作它只是一个执行单元所有 kernel 谁来跑、什么时候跑、跑完下一个是什么全由 CPU 侧决定。1.2 launch、同步、尾部空闲三笔账分开算我把这个开销拆开看才发现问题不是单一原因launch 开销每次cudaLaunchKernel从 CPU 发起要经过驱动层的参数校验、命令缓冲区编码、再通过硬件队列提交给 GPU。这个固定开销通常在 3-10 微秒和 kernel 本身多复杂没有关系。你的 kernel 越短这部分占比越刺眼。同步开销如果代码里用了cudaStreamSynchronize或默认流CPU 必须等 GPU 执行完当前任务才能继续下发。哪怕你用事件做流间同步事件查询本身也有成本一旦 CPU 需要等待那 GPU 就完全空转。尾部空闲tail effectGPU 执行完一个 kernel 后如果下一条命令还没到达SM 会进入一种等待状态。小任务高频提交时这种“执行完-等待-再执行”的空隙被反复放大。这三笔账在小 kernel 场景下每笔都不可忽略。大模型训练里单次 kernel 动辄几百微秒甚至毫秒级launch 开销占比可以忽略但做推理、做细粒度计算时这就是主瓶颈。1.3 我理解的“GPU 自主调度”分成三个层级当时我给自己总结了一个思路把“让 GPU 自己调度”分成三个递进的层级静态图级把一串有固定依赖关系的 kernel 打包成一个 CUDA Graph一次 launch 提交给 GPUGPU 自动按依赖关系执行。这个层级解决的是 launch 开销和部分同步开销。动态队列级在 GPU 全局内存里维护一个或多个原子队列由常驻 kernel 或计算 kernel 自己从队列取任务、写入结果。CPU 只负责往队列里塞初始任务后面怎么调度由 GPU 自己说了算。跨设备级用 NVSHMEM 这类对称内存通信机制让多个 GPU 直接共享任务队列通过 GPU 间的原子操作和信号量完成调度CPU 只在最外层做收尾。这三个层级不是彼此替代而是可以叠加的。CUDA Graphs 做静态骨架原子队列做动态填充NVSHMEM 把作用范围从单卡扩展到多卡。下面逐个拆开说。2. CUDA Graphs把 kernel 串成图一次 launch 的代价做十件事2.1 图为什么能把 launch 开销摊薄CUDA Graphs 的思路说起来很直白与其让 CPU 一个个发 kernel不如把一串 kernel 和它们之间的依赖关系事先构造成一张图实例化之后CPU 只需要提交一次图启动命令GPU 自己根据图里的依赖关系调度执行。传统方式下N 个 kernel 需要 N 次 host 端 launch 调用每次都有固定开销。图模式下一次cudaGraphLaunch也会有开销但只会出现一次而且是把所有 kernel 的执行计划一次性交给 GPU 侧驱动。这里有一个容易被忽略的概念图分两个阶段capture捕获和instantiate实例化。捕获阶段把 kernel 和依赖记录成图的数据结构实例化阶段把图转换成 GPU 驱动可以直接执行的执行计划这个过程可以包含 kernel 参数固化、内存访问优化等工作。实例化出来的cudaGraphExec_t才是后续重复 launch 的对象。一旦实例化完成重复 launch 的代价就只有一次图级提交内部所有 kernel 的调度完全在图形执行引擎里完成。2.2 流捕获不用手写图的构建方式CUDA Graphs 有两种构建方式纯 API 构建cudaGraphAddKernelNode一个个加节点和流捕获stream capture。我强烈推荐后者因为你可以把现有代码里的 kernel launch 原封不动地“录”成图不用去手动填节点参数。流捕获的基本用法#include cuda_runtime.h cudaGraph_t graph; cudaGraphExec_t graphExec; cudaStream_t stream; cudaStreamCreate(stream); // 热身让 driver 完成各种 lazy 初始化 for (int i 0; i 10; i) { kernel_agrid, block, 0, stream(); kernel_bgrid, block, 0, stream(); } cudaStreamSynchronize(stream); // 开始捕获 cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal); // 这段代码里的 kernel 和 memcpy 不会立即执行而是被记录成图节点 kernel_agrid, block, 0, stream(); kernel_bgrid, block, 0, stream(); cudaMemcpyAsync(d_out, d_mid, size, cudaMemcpyDeviceToDevice, stream); // 结束捕获得到图 cudaStreamEndCapture(stream, graph); // 实例化 cudaGraphInstantiate(graphExec, graph, 0); // 之后每次直接启动整张图 cudaGraphLaunch(graphExec, stream); cudaStreamSynchronize(stream); // 释放 cudaGraphExecDestroy(graphExec); cudaGraphDestroy(graph);如果后续某个 kernel 参数变了比如 batch size 变化不需要重新捕获整张图可以用cudaGraphExecKernelNodeSetParams更新对应节点的参数或者用cudaGraphExecUpdate做增量更新。2.3 capture 模式的限制比你想的要多流捕获看起来很美好但实际用起来有不少限制。我踩过的坑大致有几类捕获期间不能做同步操作cudaDeviceSynchronize、cudaStreamSynchronize、cudaEventSynchronize在捕获流中都是非法操作。原因很简单捕获期间根本没有真正的执行同步什么不能依赖未知数据捕获时如果 CPU 端代码判断一个变量的值来决定是否 launch kernel那这张图就不是确定的。CUDA 会拒绝这种捕获。所以捕获代码里尽量别写 if 分支依赖 CPU 变量。内存分配要用流序版本kernel 内部如果调cudaMalloc必须改用cudaMallocAsync否则捕获会失败。这是最坑的一点因为很多老代码还在用同步 malloc。跨流同步必须小心默认的cudaStreamCaptureModeThreadLocal只捕获当前线程提交到目标流的操作如果另一个线程同时在操作这条流直接报错。要跨线程捕获得评估用 Global 模式但那时所有流上的操作都会被纳入捕获范围副作用更大。我的经验是捕获的代码越“死板”越好。循环次数固定、分支不依赖运行时变量、内存预分配好。凡是动态的地方留给后面的原子队列来解决。2.4 实测20 个 kernel 打包后的收益在我那个推理场景里一条请求链路大概包含 20 个小 kernel串行依赖。传统 launch 模式下每次 launch 约 4-5 微秒20 次就是 80-100 微秒的纯 launch 开销。打包成 CUDA Graph 后一次cudaGraphLaunch约 10-15 微秒内部 kernel 的调度由 GPU 端图形引擎完成。我把图实例化后跑了稳定测试kernel 总执行时间约 300 微秒的场景下端到端从 400 微秒降到 315 微秒左右有接近 20% 的收益。这还是在没有做流间并行的情况。如果图里有多个分支可以并行收益会更大。但也要记住CUDA Graphs 只解决“静态依赖”的调度问题。如果任务本身是动态产生的CPU 不知道下一秒哪个请求会来那你不能每来一个请求就重新捕获一张图。这正是原子队列要解决的。3. 原子队列在 GPU 内存里写一个无锁任务信箱3.1 静态图盖不住动态任务队列要登场CUDA Graphs 的结构是固定的但真实推理负载是动态的请求什么时候来、batch 多大、走哪条分支都是运行时才知道。如果每个动态变化都触发一次图的重新构建那优化的 launch 开销又回来了。所以我们需要一个“任务信箱”生产者把任务描述符放进队列消费者常驻 GPU 的 worker kernel从队列里取任务并执行。生产者和消费者的解耦让 GPU 在没有 CPU 介入的情况下持续工作。为什么必须是无锁队列GPU 上的线程如果为了抢一个资源而阻塞会占用 SM 资源一个 block 里哪怕一个线程在自旋等待锁整个 SM 的调度能力都会受影响。更糟的是如果持有锁的线程被调度出去其他线程可能无限等待形成死锁。所以 GPU 上默认就该用无锁或短暂自旋的方案。3.2 GPU 上的 cuda::atomic作用域和内存序CUDA 从 11.x 开始提供了cuda::atomic和cuda::atomic_ref在 device code 中可以像 C 标准原子一样使用。比自定义atomicAdd那套更规范。两个关键概念必须先讲清楚线程作用域thread scopecuda::atomicuint32_t, cuda::thread_scope_device表示这个原子的可见范围是同一个 GPU 设备内的所有线程。如果跨 GPU 访问要选thread_scope_system但那个开销更大。大部分单卡内部的队列用 device scope 就够了。内存序memory order这是最容易写错的地方。GPU 是弱内存模型不同线程看到的内存写入顺序可能不一致。生产者和消费者必须通过 release/acquire 这样的内存序建立“happens-before”关系。一个简单的 SPSC单生产者单消费者环形队列可以这么写#include cuda/atomic template typename T, uint32_t Capacity 256 class __align__(128) DeviceQueue { static_assert((Capacity (Capacity - 1)) 0, Capacity must be power of 2); T m_buffer[Capacity]; cuda::atomicuint32_t, cuda::thread_scope_device m_head; // 消费者读取位 cuda::atomicuint32_t, cuda::thread_scope_device m_tail; // 生产者写入位 public: __device__ DeviceQueue() : m_head(0), m_tail(0) {} __device__ bool enqueue(const T item) { uint32_t t m_tail.load(cuda::memory_order_relaxed); uint32_t h m_head.load(cuda::memory_order_acquire); if (t - h Capacity) return false; // 队列满 m_buffer[t (Capacity - 1)] item; m_tail.store(t 1, cuda::memory_order_release); return true; } __device__ bool dequeue(T item) { uint32_t h m_head.load(cuda::memory_order_relaxed); uint32_t t m_tail.load(cuda::memory_order_acquire); if (h t) return false; // 队列空 item m_buffer[h (Capacity - 1)]; m_head.store(h 1, cuda::memory_order_release); return true; } };这个队列是 SPSC 安全的也就是同一时刻只有一个生产者、一个消费者。如果有多生产者或多消费者需要在索引分配处用fetch_add配合 CAS 处理复杂度会高不少后面会单独说。3.3 为什么 release/acquire 在这里是底线我见过不少人的队列实现图省事把所有 load/store 都写成memory_order_relaxed测试时也没问题结果一上生产就随机出现读到脏数据或任务丢失。问题就出在没有正确建立内存序。在这个队列里生产者先写m_buffer[slot]再用release语义更新m_tail。消费者用acquire语义读取m_tail一旦看到新值就说明生产者对 buffer 的写入已经对消费者可见。这是 C 内存模型里最经典的 release-acquire 配对。如果m_tail.store用relaxed消费者可能先看到 tail 增加了但 buffer 里的数据还停留在上一个任务的内容直接读到脏数据。如果m_tail.load用relaxed消费者可能在自己的 cache 里一直看到旧 tail导致队列迟迟不显示有新任务。简单记生产者发布数据用 release消费者发现数据用 acquire两者配对才能保证“先写数据、再发信号先看信号、再读数据”的顺序关系。这个原则对共享内存队列、对跨 GPU 的 NVSHMEM 信号量同样适用。3.4 从单生产者到多生产者绕不开的索引分配如果场景里多个 block 或多个 kernel 都要往队列写任务SPSC 就不够了。一个相对实用的 MPSC 方案是生产者先用m_tail.fetch_add(1, memory_order_relaxed)拿一个唯一槽位然后自旋等待槽位可写消费者仍然是单一方。但这里有个坑fetch_add只是预留了索引不代表数据已经写入。如果多个生产者都预拿了索引但还没有写入 buffer消费者看到 tail 增大后去读读到的还是旧数据。所以多生产者场景下不能直接用 tail 做发布信号通常要增加一个“每个槽位是否 ready”的标记数组或者用带 tag 的 CAS 机制。我的建议是在 GPU 上能设计成 SPSC 就尽量 SPSC或者把一个队列拆成多个 SPSC 队列比如按 SM 分区尽量避免 MPMC。原因不仅是实现复杂度还有多生产者竞争同一个原子计数器时的性能损耗在 GPU 上会放大成 SM 间的全局内存竞争。4. NVSHMEM用对称内存把任务队列扩展到多 GPU4.1 NVSHMEM 和 NCCL 的定位差异很多同学一听多 GPU 通信第一反应是 NCCL。NCCL 解决的是集合通信比如 AllReduce、Broadcast、AllGather是一对多/多对多的批量数据交换。但 NVSHMEM 的定位完全不同它提供的是 RMA远程内存访问和 PGAS分区全局地址空间模型你可以直接“读远端 GPU 的某个地址”、“写远端 GPU 的某个地址”像操作本地内存一样操作对称内存。换句话说NCCL 是“把一堆数据聚合起来算”NVSHMEM 是“我想访问哪块内存就去访问哪块内存颗粒度可以很小”。这对实现跨 GPU 队列特别有用队列的头尾指针和 buffer 如果放在对称内存里任何一个 GPU 都能直接对另一个 GPU 的队列进行原子操作这就是跨设备自主调度的基础设施。4.2 对称内存的第一步nvshmem_malloc 与 RMANVSHMEM 的核心抽象是对称内存。每个参与的 PE通常是一个 GPU用nvshmem_malloc分配的内存在所有 PE 上都有相同大小、相同布局的“副本”。逻辑上你操作的是一个全局数组PE 0的偏移量 X 和PE 1的偏移量 X 是对称的。最简单的用法#include nvshmem.h #include nvshmemx.h int main() { nvshmem_init(); int mype nvshmem_my_pe(); int npes nvshmem_n_pes(); // 所有 PE 都分配对称内存每个 PE 都有 flag 的本地副本 int *flag (int *)nvshmem_malloc(sizeof(int)); *flag 0; nvshmem_barrier_all(); // 确保所有 PE 都完成本地初始化 if (mype 0) { // 向 PE 1 的 flag 地址写入值 1 nvshmem_int_p(flag, 1, 1); } if (mype 1) { // 轮询 PE 0 的 flag直到看到 1 while (nvshmem_int_g(flag, 0) 0) { } printf(PE %d observed flag%d\n, mype, nvshmem_int_g(flag, 0)); } nvshmem_barrier_all(); nvshmem_free(flag); nvshmem_finalize(); return 0; }nvshmem_int_p是“put”把值写到指定 PE 的对称地址nvshmem_int_g是“get”从指定 PE 的对称地址读值。这比cudaMemcpyPeerAsync要轻量得多特别适合小粒度的控制信息交换。4.3 从“轮询 flag”到跨 GPU 队列有了 put/get 和原子操作跨 GPU 队列的基本形状就出来了。比如 PE 0 是任务分发方PE 1 是消费方队列描述符放在对称内存中struct CrossGpuQueue { cuda::atomicuint32_t, cuda::thread_scope_system head; cuda::atomicuint32_t, cuda::thread_scope_system tail; uint8_t buffer[QUEUE_CAPACITY * ITEM_SIZE]; };生产者PE 0入队的核心步骤是用nvshmem_uint32_atomic_fetch_add(queue-tail, 1, consumer_pe)原子拿一个槽位。用nvshmem_T_p(queue-buffer[slot], item, consumer_pe)把任务写进消费者 PE 的 buffer。调用nvshmem_fence()确保 buffer 的写入已经对远端可见。用nvshmem_uint32_atomic_add(queue-ready, 1, consumer_pe)发一个“新任务已就绪”的信号。消费者PE 1出队的核心步骤是轮询本地/远端的ready计数看有没有新任务。拿到新任务计数后用nvshmem_uint32_atomic_fetch_add(queue-head, 1, ...)拿消费槽位。从 buffer 里读任务。这里最难理解的就是第三步 fence。为什么 buffer 写完还要 fence而不是直接把 tail 加一完事问题在于fetch_add和 buffer 的 put 操作是两个独立的 RMA 操作不同的硬件路径、不同的到达时间。消费者可能先看到tail增加马上去读 buffer但 buffer 的写入还没到达消费者 PE读到的自然就是脏数据。nvshmem_fence()的作用是保证“当前 PE 之前发出的所有 RMA 写操作在后续 RMA 操作对同一远端 PE开始之前已经到达远端”。所以必须先 fence再发 ready 信号消费者看到 ready 增加时buffer 数据必然已经写好了。这个坑在写跨 GPU 队列时几乎必踩我自己就是踩过之后才彻底弄明白 NVSHMEM 的 fence 语义。4.4 与 CUDA Graphs 组合时的兼容性注意理论上一张 CUDA Graph 里可以包含 NVSHMEM 的 RMA 操作因为它们最终会提交到流上可以被 stream capture 记录。但实际操作里有两个典型问题自旋等待不能随便放进图里如果图里某个 kernel 在while (flag 0) {}自旋等远端而远端数据是由同一个图里另一个节点写入的那这个图里就存在跨设备依赖。CUDA 图执行引擎不知道如何调度这种跨设备事件依赖捕获时不一定报错但执行时可能直接挂死。信号量的等待在流捕获中语义不完整NVSHMEM 的等待操作没有对应的“流依赖”关系图优化器可能把无关节点重新排序破坏你预期的时序。我的个人建议是图只包纯计算流跨 GPU 的同步和等待放到图外的常驻 kernel 里。CUDA Graphs 负责单设备内的高效执行NVSHMEM 负责多设备间的松耦合任务传递两者用事件或队列衔接而不是强行把 NVSHMEM 同步逻辑塞进图里。5. 三者结合的一次实战从 30% 利用率到 70% 的经过5.1 目标场景多 GPU 动态 batch 推理实践是检验笔记的唯一标准。我搭了一个简化版的多 GPU 推理流水线来验证这套组合拳8 个 GPU其中 GPU 0 作为调度器其余 7 个作为 worker。请求不断到达 GPU 0每个请求对应一个小计算任务。worker GPU 各自独立处理任务处理完直接写结果到共享输出区。目标是让所有 GPU 都尽量保持忙碌而不是等 CPU 一个个分配。5.2 架构设计把调度层下沉到 GPU整体结构分三层CUDA Graphs 层每个 worker GPU 上有一个固定的计算流水线用 CUDA Graph 实例化好。流水线内部是编码器、注意力、FFN 等固定 kernel 序列一次 graph launch 跑完一个 batch 的推理。原子队列层GPU 0 上有一个全局任务队列任务描述符是一个整数索引 参数偏移量。新的请求由 CPU 或者一个轻量级 entry kernel 写入队列worker 的计算 kernel 直接从这个设备队列里dequeue。NVSHMEM 层GPU 0 的队列需要把任务派到不同 GPU 上。我用 NVSHMEM 对称内存建了一个跨设备分发队列GPU 0 负责向各 worker GPU 的本地队列写入任务并发送 ready 信号worker 只轮询自己的 ready 计数有活就干没活就自旋等待。整个过程中CPU 只在最开始时把原始请求灌入 GPU 0 的队列之后所有任务分发、消费、结果回写都在 GPU 之间直接完成。5.3 性能变化GPU 利用率与端到端延迟在同样的负载模型下对比三种配置配置GPU 利用率端到端 P99 延迟说明纯 CPU 调度 多次 launch约 30%8.2 msbaseline大量时间花在 launch 和同步CUDA Graphs 打包计算流约 52%5.6 mslaunch 开销大幅降低但 CPU 仍在分发任务Graphs 原子队列 NVSHMEM约 70%4.1 ms任务分发下沉到 GPUCPU 几乎不参与中间调度这个数字是特定场景下的观测值硬件是 8 卡 NVLink 互联的 A 系列 GPU负载是模拟的动态 batch 推理。换到 PCIe 互联、或者 kernel 本身更长更重的场景收益比例会有变化。但方向是一致的把调度从 CPU 挪到 GPU 内部减少 launch 和同步的固定开销收益立竿见影。5.4 两个排查了很久的问题这一个多月里印象最深的是两个诡异问题。第一个是 CUDA Graph 捕获期间 NVSHMEM 自旋等待导致整体卡死。最初我把“轮询远端 ready 计数”的代码放进了 worker 的 CUDA Graph 里捕获时一切正常但一执行 graph所有 worker GPU 全部挂住没有任何报错。后来我意识到图中的自旋等待节点会让图执行引擎误判整个图仍然在运行但实际依赖的远端事件永远不会被图调度器“唤醒”形成死锁。解决方案是把自旋等待的循环从图里拿出来放到一个独立的常驻 kernel 里图只负责纯计算部分两者通过设备内存队列衔接问题立刻消失。第二个是原子队列偶发的“幽灵任务”。现象是消费者偶尔会读到一个索引值非常大、明显是垃圾数据的任务。排查了很久用 compute-sanitizer 的 racecheck 工具才发现是生产者的 buffer 写入和 tail 更新之间缺少 release/acquire 配对。在某些 GPU 架构上全局内存写入的可见性不像 x86 那么强tail 已经更新但 buffer 数据还在 L2 里没有刷出去。把 enqueue 的 tail 更新从 relaxed 改成 release、dequeue 的 tail 读取从 relaxed 改成 acquire 之后幽灵任务再也没出现过。排这两个坑花的精力比实现整个系统还多。但反过来也说明这三个机制如果单独用都不算复杂组合在一起时内存序和跨设备可见性才是真正的深水区。6. 什么时候别用这套方案以及调试工具推荐6.1 什么情况下别上这套方案CUDA Graphs、原子队列、NVSHMEM 这套组合拳不是银弹。反过来讲如果你的场景符合下面任一条就不要硬上kernel 本身就很大单次执行几百微秒甚至毫秒级。这种场景 launch 开销占比很小CUDA Graphs 收益有限反而增加了图管理和参数更新的复杂度。任务到达率很低GPU 大部分时间本来就在 idle。这种场景应该先解决“任务太少”的问题而不是调度开销问题。数据依赖高度动态每个请求走完全不同的分支kernel 数量、顺序都不一样。CUDA Graphs 很难表达这种动态结构强行用会变成“反复重新捕获图”比传统 launch 还慢。团队没有 GPU 调优经验。这套方案的内存序问题、跨设备可见性问题、平台差异问题对调试能力要求很高没有经验支撑真的会怀疑人生。6.2 推荐的工具链我调这套系统时的常用工具按使用频率排Nsight Systemsnsys第一排查工具看 CPU 和 GPU 的时间线重叠、launch 间隙、API 调用耗时。一眼就能看出瓶颈是 launch 开销还是 kernel 本身。Nsight Computencu深入单个 kernel 内部看 SM 占用、内存吞吐、调度 stall 原因。LaunchStats 部分是判断 launch 开销影响的直接依据。compute-sanitizer排查内存越界、race condition。搭配--tool racecheck能抓出原子操作相关的数据竞争上面幽灵任务就是靠它定位的。CUDA_LAUNCH_BLOCKING1临时把异步 launch 变成同步配合调试器定位代码行。注意这个只用于排查千万不要开进性能测试。6.3 这套思路给我的一点启发做 AI 系统性能工程最容易陷入的误区是只看 kernel 本身的优化忽视了调度层面的浪费。CUDA Graphs、原子队列、NVSHMEM 这三样东西分别从静态调度、动态任务分发、跨设备通信三个层面把指挥权下沉到了 GPU。CUDA Graphs 是个很好的起点改动量小、收益直观原子队列是真正实现“GPU 自主调度”的核心但内存序问题一定要吃透NVSHMEM 则是打开多卡自治大门的钥匙值得在这个方向上持续深入。如果让我按个人经验给一条建议先拿 CUDA Graphs 做一次性能验证看调度开销在你的场景里占比到底多大如果收益明显再考虑引入原子队列等单卡内部调度稳定了再往上叠 NVSHMEM。这套路线比较稳不会一上来就被跨设备调试的复杂度劝退。