Skip to main content

Sync And Async

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