Skip to main content
Sync And Async
面临的问题
- 多种类型的硬件单元需要进行同步
- 不确定的循环次数
- N to N的同步需求
- 无缝的同步,无缝的并行
- 频繁的同步需求,频繁的状态pulling,低latency
- 灵活的抽象适应所有的同步需求
- transformer的Flash-attention就需要在L1内做多次fusion,不是简单的DMA和算力的同步
- 方便的软件使用
采用Global调度的逻辑进行确保并行和同步
- 软核实现
- 类似一个全局锁,成为性能瓶颈点
- 调度颗粒度大,难以实现精细化的控制
- DMA、Kernel之间的同步
- 难以做到无缝的调度
采用专有的中心化硬件
- 采用专有的,多层级的Sync硬件,自带计数和复位功能,配合mailbox实现N to N异步依赖
- 可以实现复杂的producer consumer模型 Pipeline
- 软件编写困难
- 软件需要根据需求管理全局硬件资源得使用
- 不同得同步需求,使用不同得同步硬件,增加软件复杂性
- 中心化的控制模式,灵活度不够,不能实现一些比较精细得控制
- 性能不是最理想
- 虽然可以做到比较快的同步,但是每次发同步命令(write mbx,wait mbx)都要经过多次写总线并ack,latency和带宽占用都不理想
去中心化的全异步模型 mbarrier
- 去中心化,每个硬件单元都可以同步别人和被别人同步,每个单元都能实现对应得标准得M barrier的功能
- 提高灵活性
- 统一的同步方式,简化设计和使用
- 传输数据的同时,传递同步信号
- 数据传递带上依赖信息,目的存储器可以识别依赖信息,并正确得处理,而不用返回数据发送方
- 避免没有必要的fence
- 改进前 Core0 Store() Core0 Fence() memory Ack() Core0 Signal(Core1 mbx)
- 改进后 Core0 Store_mbarrier() Core1 Wait(mbarrier)
- 灵活性
- 每一轮的同步都可以根据实际运算数据的大小,动态设置需要的同步单元和数据搬运数量
- 性能好
- 每个单元只用wait自己local的mbarrier单元的counter信号就可以保证同步,latency小
- 需要的信号transaction最小化
No Comments