AI加速芯片 C++ SIMD The support for these instructions is wide but not universal. Both Intel and AMD support the compatible version of FMA, called FMA 3, in their CPUs released since 2012-2013. See hardware support section for more info. Another caveat, the latency of FMA is not great, 4-5 CPU cycles on modern CPUs. If you are computing dot product or similar, have an inner loop which updates the accumulator, the loop will throttle to 4-5 cycles per iteration due to data dependency chain. To resolve, unroll the loop by a small factor like 4, use 4 independent accumulators, and sum them after the loop. This way each iteration of the loop handles 4 vectors independently, and the code should saturate the throughput instead of stalling on latency. See this stackoverflow answer for the sample code which computes dot product of two FP32 vectors. http://const.me/articles/simd/simd.pdf Open GPGPU Ventus GitHub - THU-DSP-LAB/ventus-gpgpu: GPGPU processor supporting RISCV-V extension, developed with Chisel HDL rvgpu https://gitee.com/rvgpu Vortex https://github.com/vortexgpgpu/vortex https://github.com/vortexgpgpu/vortex_tutorials https://github.com/cupbop/CuPBoP [![image.png](Open GPGPU/acvimage-png.png)](Open GPGPU/acvimage-png.png) [![image.png](Open GPGPU/GkWimage-png.png)](Open GPGPU/GkWimage-png.png) CUDA指令和架构 A100显卡上的tensorcore有自己的私有寄存器吗，微架构探索 https://zhuanlan.zhihu.com/p/620257581 乘影GPGPU架构文档手册v2.01.pdf https://zhuanlan.zhihu.com/p/166180054 https://www.tinyedi.com/cuda_learning/#pipeline https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html https://zhuanlan.zhihu.com/p/486224812 虽然异步屏障仍然是 Hopper 编程模型的一部分，但 Hopper 添加了一种新形式的屏障，称为Asynchronous Transaction Barrier。Asynchronous Transaction Barrier与异步屏障非常相似。它也是一个拆分屏障，但它不仅计算线程到达，还计算事务。 Hopper 包含一个用于写入共享内存的新命令，该命令传递要写入的数据和事务计数。事务计数本质上是一个字节计数。异步事务屏障将阻塞等待命令的线程，直到所有生产者线程都执行了到达，并且所有事务计数的总和达到预期值。异步事务屏障是用于异步内存副本或数据交换的强大新原语。 Predicate 有时候代指指令前加的形如 @P0, @!P3 等的predication，PTX文档里也把这个叫guard predicate。有时候也会特指那个predicate register，如 P0, P6, PT ，反正都差不多这个意思， 不混淆就行。在之前的文章也提到过，在现有的所有架构中，每个指令都有4bit的编码来指定每 个predicate，3bit用来指定索引（所以每线程有2^3=8个predicate register P0~P7 ，其中 P7=PT 为恒 True ），1bit表示是否取反。如果是 @PT ，那就会在反汇编中省去不显示。那如果是 @! PT 呢？嗯，大家自己想…… Predicate是控制某个线程是否执行某个指令的两种方式之一，另一种就是conditional branch。两 者的区别在于用predicate时，可以让warp内的所有线程名义上走同一路径而省去跳转的开销， 从而也避免了可能的divergence。因为branch的latency比较长，还涉及到instruction cache的问 题，一般很短的分支是不太愿意跳转的。当然branch在出现divergence的时候，内部也有一个 mask，表明当前这个thread是否active，但是用户不能直接修改这个mask。PTX中可以通过warp vote或是load特殊寄存器 %lanemask_* 之类的方法获得当前warp内的mask情况。 Cache control 指令：这个主要意义是我知道我将要访问的某个元素位于某个cache line，但 是又不确定具体要哪一个值，所以没法先load。所以可以先把整个cacheline放到cache里， 等到要用的时候从cache里取。这样load的latency可以更好的被隐藏。PTX有相关的控制指 令，但是我在Turing上测试过几次，只测到过TLB的latency被消除，实际访问latency好像 没变，存疑。 Control code 主要有reuse、read barrier、write barrier、wait barrier、yield hint、stall count等几个 域。reuse是唯一能在反汇编文本里看到的。它可以有限的解决一些GPR的bank conflict问题，也 许同时还能减少GPR的读写，节省一点功耗。关于barrier的几个主要控制thread内的依赖性问 题，这会影响指令发射的仲裁过程。yield和stall count主要影响的是warp调度的逻辑。 memory类的指令是性能优化的重点中的重点。绝大部分未经优化的程序都会是memory bound， 多数实际应用优化完了还是memory bound。所以，通过合理的选择相应指令，搭配合适的内存 排布，从而更好的隐藏内存访问的latency，或者是减少相应访问的开销，是性能优化中的主要 课题之一。 SASS里并没有直接的条件跳转指令，所有的条件跳转是用predicate实现的。 每个指令能否发射还要满足相应的依赖关系和资源需求。比如指令 LDG.E R6, [R2] ; 首先需要等待之前写入 R[2:3] 的指令完成，其次需要当前memory IO的queue还有空位，否则指令也无法下发。还有一些指令可能有conflict的情况，比如shared memory的bank conflict，register的bank conflict，atomic的地址conflict，constant memory在同一warp内地址不统一的conflict等等，这些都有可能导致指令re-issue（甚至cache miss也可能导致指令replay）。这些情况会不会重复占用dispatch port发射带宽 当前CUDA的所有架构都没有乱序执行（Out of order） Register Reuse Cache : 编译器hint，明确一个寄存器的值会被下面的指令使用到，可以放到cache里面去，避免再次read Wait Dependency Barrier : 有6bit，每个bit表示是否需要等待对应的dependency barrier。每个线程有6个dependency barrier，每个barrier都可以被后面的Read或Write操作设置上。设置wait dependency barrier是等待依赖的其中一种方式。SASS里面还有一个对应的指令，如 DEPBAR.LE SB0, 0x0, {2,1} ; Read dependency barrier : 有3bit，表示需要设置的6个barrier中对应的索引（0~5，对应barrier 1-6，如果不需要设置barrier，就设置为0b111）。Read dependency barrier主要是一些指令不会在一开始就把所有操作数读进去，所以需要hold住GPR的值，防止后面的指令在它读取其内容之前把GPR改掉。 使用Read dependency barrier的主要是memory类的指令 ，但是一些转换指令如 F2I/I2F 之类好像偶尔也能见到。 Write dependency barrier : 与read dependency很类似，也是3bit，后面跟barrier索引。注意Read和Write两者用的dependency barrier资源是一样的，也都是上面wait的那6个。Write dependency barrier比较好理解，就是某个指令要把操作结果保存到某个GPR或是predicate中， 使用barrier进行保序可以防止出现data race 。不过这主要针对的是不定长latency指令。如果一个指令的latency是确定的（或者有不太长的上限），那用后面提到的stall cycle停足够长时间就可以保证没有race。 Yield hint flag : 是1bit。如果Yield，就表示下一个cycle会优先发射其他warp的指令。也许yield这个bit就是stall count的高位。只是假如这个bit不为0，那stall的cycle会>16，相当于warp被切换的概率也会大大增加。 Stall count : 有4bit，表示当前指令后需要stall指令发射的cycle数，然后再决定是不是要继续发射。 CUDA性能优化原则 寻找并行化顺序代码的方法。 最大限度地减少主机和设备之间的数据传输。 调整内核启动配置以最大限度地提高设备利用率。Adjust kernel launch configuration to maximize device utilization. 确保全局内存访问被合并。 尽可能减少对全局内存的冗余访问。 Minimize redundant accesses to global memory whenever possible. 避免同一扭曲内的线程进行长序列的发散执行。Avoid long sequences of diverged execution by threads within the same warp. TMA TMA 允许应用程序在全局内存和共享内存之间以及同一集群中不同 SM 的共享内存区域之间双向传输 1D 和最多 5D 张量（请参阅 线程块集群 ）。此外，对于从共享内存到全局内存的写入，它允许指定按元素归约操作，例如 add/min/max 以及按位和/或最常见的数据类型。 memcpy_async 的API有三个变形, 分别是group,barrier,pipeline. group: 搭配 cooperative_groups::wait(group) 使用,注意不是 group.sync() , 从逻辑上说，相当于动态为group插入了一个异步线程. cooperative_groups::wait 只是让当前线程等待所有异步线程完成,并不会有 group.sync() 的效果. barrier: 搭配 barrier.arrive_and_wait() 使用,从实现角度讲,相当于动态增大了barrier的count,这个count只有在async_memcpy完成时才会下降 pipeline: 约等于线程安全的queue SIMT With Vector DMA 2D算力的表达和设计 微架构和ISA的配合，软件控制流水线，硬件hzd检查简单高效 GS和Cache系统的设计 左右支的复用 RO WO 存储类型的利用 [![image.png](SIMT With Vector/sfeimage-png.png)](SIMT With Vector/sfeimage-png.png) Cuda Pipeline 同步机制 pipeline 它实现上是一个proxy pattern, cuda::pipeline 是每个thread访问 pipeline_shared_state 的proxy pipeline_shared_state的模板参数也仅仅是描述pipeline会被 共享的范围 ,和barrier类似. pipeline_shared_state 需要在共享的内存区域创建 thread_scope的pipeline是性能最优秀的, 它不使用任何共享资源, 用 cuda::pipeline pipeline = cuda::make_pipeline() 直接创建 make_pipeline是一个同步操作,它用于初始化 pipeline_shared_state ,为当前线程确定role, 并通过线程间通信确定group内producer/consumer的数量 pipeline逻辑上是一个 fifo, head in, tail out, 这个pipeline的元素称为stage pipeline这个proxy有三种可能的角色, consumer, producer, both fifo的最大容量是编译期创建pipeline_shared_state时指定的,当fifo中的stage满时,后续的producer将在acquire时被阻塞. 对于role为producer的pipeline: pipeline.producer_acquire(); 表明当前thread开始尝试向fifo中push数据,并需要lock相关的资源 当producer_acquire成功后,当前thread的后续async指令都可以发射到acquire获得的stage中 pipeline.producer_commit();提交当前thread的async任务 当group内所有的producer都commit后, 对应的stage将正式被push到fifo中,这个stage将会在内部所有async任务完成后被标记为ready 对于role为consumer的pipeline pipeline.consumer_wait(); 表明当前thread需要从fifo中取数据,并需要锁定相关资源 当consumer_wait()成功后,说明group中的producer提交的一个stage已经ready,可以开始处理数据 pipeline.consumer_release(); 表明当前thread的任务已经完成. 当group内所有的consumer都release后,这个stage将被pop出fifo 注意: stage只有执行顺序的逻辑意义, fifo tail对应的是哪些数据搬运或者计算需要由用户自行track // 申请一个pipeline，同步API，会自动记录producer和consumer的数量 cuda::pipeline pipeline = cuda::make_pipeline(block, &shared_state, thread_role); if (thread_role == cuda::pipeline_role::producer) { // Only the producer threads schedule asynchronous memcpys: pipeline.producer_acquire(); size_t shared_idx = fetch_batch % stages_count; size_t batch_idx = fetch_batch; size_t global_batch_idx = block_batch(batch_idx) + thread_idx; size_t shared_batch_idx = shared_offset[shared_idx] + thread_idx; cuda::memcpy_async(shared + shared_batch_idx, global_in + global_batch_idx, sizeof(int), pipeline); // 同步接口？ pipeline.producer_commit(); } if (thread_role == cuda::pipeline_role::consumer) { // Only the consumer threads compute: // 同步接口？ pipeline.consumer_wait(); size_t shared_idx = compute_batch % stages_count; size_t global_batch_idx = block_batch(compute_batch) + thread_idx; size_t shared_batch_idx = shared_offset[shared_idx] + thread_idx; compute(global_out + global_batch_idx, *(shared + shared_batch_idx)); pipeline.consumer_release(); } Cuda Tensor Core 要保持张量核心持续运行并不容易。 研究人员发现GPU硬件具有一些特性，对于保持矩阵乘法的运行非常重要： WGMMA指令虽然是必要的，但使用起来颇为麻烦。 共享内存的速度并不如预期的快，使用时还需格外注意。 生成地址的成本较高。 保持高占用率对于提升性能是有益的，寄存器至关重要 早期GPU中的张量核心指令如wmma.mma.sync和mma.sync，要求SM一个子单元内的32个线程的一个warp同步传输数据块至张量核心并等待结果。 wgmma.mma_async指令则不同。它允许128个连续线程跨SM所有子单元协作同步，并从共享内存及寄存器（可选）异步启动矩阵乘法。这使得这些warp在等待矩阵乘法结果时可以处理其他任务。 然而，这些指令的内存布局极其复杂。未重排的共享内存布局合并性差，需要额外的L2带宽。重排的内存布局记录不准确，研究人员花费了大量时间才弄明白。最终发现，这些布局只适用于特定矩阵形状，并与wgmma.mma_async指令的其他部分不兼容，例如硬件仅在未重排的布局下转置子矩阵。 此外，未重排的wgmma布局内存合并性差且有bank conflicts。尽管TMA和L2缓存在如flash attention这类内核上能较好地掩盖这些问题，但要充分利用硬件，必须精心控制内存请求的合并和避免bank conflicts。 TMA允许用户在全局和共享内存中指定多维张量布局，命令其异步提取张量的一部分，并在完成后触发一个屏障。这大大节省了地址生成的开销，并简化了pipelines的构建。 Volta Turing Ampere Hopper Base Size* 64个F16(4x4x4) 64个F16(4x4x4) 256个F16(8x4x8) 512个F16(8x4x16) tensorCore个数 672(8/SM) 576(8/SM) 512(4/SM) 576(4/SM) SM内算力* 512个F16 512个F16 1024个F16 2048个F16 总算力 FOPS(F16) 125T 130.5T 312T 1000T boost clock(MHz) 1530 1455 1410 ≈A100x1.3 FP64 - - √ √ TF32 - - √ √ BF16 - - √ √ FP16 √ √ √ √ FP8 - - - √ INT8/UINT8 - √ √ √ INT4/UINT4 - √ √ - INT1 - √ √ - Volta Turing Ampere Hopper SM/TPC 2 2 2 2 process blocks/SM 4 4 4 4 FP64/PB 8 - 8 16 FP32/PB 16 16 16 32 INT32/PB 16 16 16 16 tensorCore/PB 2 2 1 1 LSU/PB 8 4 8 8 register file 64KB(16384*32bit) 64KB(16384*32bit) 64KB(16384*32bit) 64KB(16384*32bit) L0 ICache/PB 1 1 1 1 L1/SHM/SM 128KB 96KB 192KB 256KB warp scheduler/PB 1(32thread/clk) 1(32thread/clk) 1(32thread/clk) 1(32thread/clk) dispatch uint/PB 1(32thread/clk) 1(32thread/clk) 1(32thread/clk) 1(32thread/clk) [![image.png](Cuda Tensor Core/Ubsimage-png.png)](Cuda Tensor Core/Ubsimage-png.png) tensorCore的启动是依靠SM内的thread（warp）共同完成的，每个thread能提供的数据个数是相同的，因此统计这个算力能有助于分析SM内部的thread和tensorCore之间的调度方式 边缘视觉的软件/硬件方案 软件 iree的路线 https://github.com/openxla/iree 兼容CUDA/OpenCL的路线， CuPBoP   ，  Vortex Open GPGPU Triton  https://openai.com/research/triton Halide Openxla  https://github.com/openxla/xla UXL Group 的开发工作主要集中在英特尔OneAPI的软件工具包上。OneAPI 基于名为 SYCL 的早期框架（SYCL是一种Khronos开放标准 https://flow-computing.com/technology/ mojo开发方案？ 模型 Segment Anything 主流AI加速大芯片 NVIDIA Sohu 芯片初创公司Etched近日宣布推出了一款 针对 Transformer架构 专用的AISC芯片 “Sohu”，并声称其在AI大语言模型（LLM）推理性能方面击败了NVIDIA最新的B200 GPU，AI性能达到了H100的20倍。 Google TPU Sambanova SN40L是第三代芯片（加了HBM）。**前两代芯片依靠了Dataflow的Spatial编程特点，降低了对DRAM高带宽的需求，走了大容量的DDR路线。**而第三代芯片，是在此基础上，又加入了64G的HBM， 既要带宽又要容量 。从PR的内容来看，应该是加在了DDR4和PMU SRAM之间。 最新一代芯片 RDU SN40L 的设计还加入了520MB SRAM (300 TB/s)、64GB HBM3 (3 TB/s) 和 1.5 TB DDR5 (0.3 TB/s)三级内存方案。 SN40L 则进一步针对大模型进行优化，有望成为 Serving LLM （Continual Pre-training、Finetune 和 Inference）性价比最高的设备。Sambanova 也推出了自己的软件服务，为非科技企业客户进行模型咨询和 fine-tune，进而搭建出一套软硬件产品。 私有大模型市场： 软硬件结合独立开发的难度高 共用模型难确保数据隐私 企业的个性化需求满足 自主掌控信息的时效性 模型可审计性和所有权归属界定 Pattern Compute Unit （PCU）：计算单元； Pattern Memory Unit （PMU） ：由 SRAM 制成的内存单元； Switch：负责高效连接 PCU 和 PMU； Address Generator Units （AGU） 和 Coalescing Units （CU） ：共同负责连接电脑的其他部分，比如 off-chip 的 DRAM 内存、硬盘或网卡等资源。 Cerebras 46,225平方毫米，1.2万亿个晶体管 400,000核 18 GB的片上内存 9 PetaByte 每秒的内存带宽 100 Petabit 每秒的架构带宽 台积电16nm工艺制造 Graphcore Graphcore联合创始人兼CEO Nigel Toon去年在接受电子创新网等媒体采访时曾表示对于CPU、GPU、FPGA和ASIC而言，Graphcore的IPU处理器是与它们完全不同的，Graphcore 的IPU特点可概括为： 1、同时支持训练和推理 2、采用同构多核（many-core）架构，有超过1000个独立的处理器； 3、支持 all-to-all的核间通信，采用Bulk Synchronous Parallel的同步计算模型； 4、采用大量片上SRAM，不需要外部DRAM。 Programming model IPU Programmer Groq 软件定义的硬件 Groq不寻常的软件优先方法始于构建一个原型编译器，而不是硬件原型。围绕编译器，再构建硬件结构，由此产生的TSP有一个简化的硬件设计，但所有执行程序都在软件中进行。软件实质上协调了所有数据流和时序，从而确保计算不停顿，而且延迟和性能都是可预测的。 Groq首席架构师Dennis Abts解释道：“我们将大量的控制权交给了编译器，这样就能够在软-硬件接口上进行一些折中…… 从而获得了确定性执行的结果，所有的latency都是明确的 。” Abts是一位在Google数据中心工作了12年的资深人士，他还曾在科雷(Cray)公司担任硬件架构师十多年。Abts解释说，编译器即可以控制程序执行，也可以控制功耗状况，因此，在 编译时可以准确预测到精确、可重复的执行时间，以及每个模型运行的功耗 。 Abts说：“我们认为这使我们在易用性方面更具优势。” 编译时即了解执行时间和功耗情况意味着“可以从模型开发的角度进行快速试验，部署系统时对所能达到的性能胸有成竹。” 编译器可以完全控制芯片，无论是动态的，还是静态的。 他说：“没有什么动态分析代码之说，因为静态与动态是一样的，这样可以实现一些非常好的特性。” Habana CAISA CAISA定制数据流架构为鲲云自主研发的源头性技术，是为深度学习神经网络定制的高性能AI计算架构。 CAISA架构通过数据流流动次序来控制计算顺序，消除指令操作导致的额外时间开销，让CNN网络的算子级数据流图可以实现高效流水线运算。 同时CAISA可并行执行数据访问和数据计算，进一步减少计算单元的空闲时间，最大化地利用芯片的计算资源，从而为客户提供更高的实测算力， 算力性价比领先市场。 ASIC设计成固定的功能块，通过数据输入来触发计算，而不需要指令。所有的功能块都是固定功能。（功能块内部也是有指令存储？可以通过更改内部存储的指令来更改功能块的功能）同一个功能块在一个神经网络计算内，可以被重复使用，而不需要更改指令。 Xilinx Versal acap Adaptive Compute Acceleration Platform ，FPGA+ARM+AI Engine。 该AI Engine是一组SIMD核阵列，每个核都包含了完整的RISC处理器、定点SIMD处理单元、浮点SIMD处理单元以及本地内存。每个核之间还可以通过片上网络（NoC）连接到一起，从而可以实现高度灵活的数据流。 Versal AI Core系列芯片中，将会集成128-400个AI Engine，从而实现43-133TOPS的INT8定点计算能力。 Triton 高层次Kernel开发语言Triton The aim of Triton is to provide an open-source environment to write fast code at higher productivity than CUDA, but also with higher flexibility than other existing DSLs. https://github.com/openai/triton https://triton-lang.org/main/index.html Triton更像是一个面向AI加速器算子开发的领域开发语言，为了能够将用户使用Triton语言开发的kernel映射到具体硬件上的执行码，需要设计开发相应的Triton compiler来完成这层映射。所以当我们说Triton的时候，其实隐指了 Triton语言+Triton编译器 这两个事物的综合体。 Triton的核心设计思想---- Block-wise编程，Block上面的归用户，Block内部的归Triton compiler自动化处理 。相应地，Block内部的优化细节，也交由Triton compiler处理了。 优化Passes在Triton当前的实现里主要包括了NV GPU计算kernel优化的一些常见技巧，包括用于辅助向量化访存的 coalescing 、用于缓解计算访存差异的 pipeline / prefetch ，用于避免shared memory访问bank-conflict的 swizzling 。 硬件开发及仿真工具 LogicSIM生成RTL 时间定义只对reg有作用 同一时间reg只能被一个信号写 logic数据要存到reg才能继续被下一次使用 reg只有在clk（统一的时钟）的上升沿才会被触发写 logic的一次执行不会跨越clk（统一的时钟）的上升沿 所有的latch简化为get上一个cycle的数据，set下一个数据，不用复杂的map和锁 module 定义一堆latch和logic的集合 throughput都是1 latency可以>1 CIRCT项目的软件栈 PyChip vs Chisel AI加速芯片架构的动态性支持讨论 动态性的级别划分 数据内容不同动态 算子数据的尺寸动态 大部分非GPGPU的DSA架构都支持得不好，有各种问题 同步开销大，同步资源管理难，存储资源浪费，程序代码大 等问题 不同的数据地址的程序代码动态 通过条件执行、跳转执行支持的程序动态 GPGPU类支持 自动生成、调用不用的单元的代码动态 很少在加速卡上实现，大部分是CPU可以实现 不同架构对动态性支持的区别 SIMT类型的GPGPU 华为和寒武纪类的DSA NPU推理DSA 类似燧原的SPMT 主流推理小芯片 NVIDIA Jetson Orin HAILO Hailo-8 15 等等 domain-specific-dataflow-processing 据悉，它能够在功耗低于5W的情况下，以每秒生成10 个token（TPS）的速度流畅运行Llama2-7B模型。同时，在处理Stable Diffusion 2.1模型（该模型可根据文本提示生成图像）时，Hailo-10同样在超低功耗范围内实现了每张图像评级低于5秒的优秀表现。 Hailo-10的性能高达每秒40万亿次（TOPS）。根据最新公布的基准测试数据，Hailo-10在性能上超越了集成神经处理单元（NPU）解决方案，同时在功耗方面表现更为出色。与Intel的Core Ultra NPU相比，Hailo-10在保持功耗减半的同时，实现了至少2倍的性能提升。 算能 1684X 支持高达 17.6T 的INT8峰值算力 LPDDR4x 68.3GB/s，最大支持16GB 17W 爱芯元智 地瓜机器人 https://d-robotics.cc 聆思科技 https://www.listenai.com/products/chips/csk6 AMD旗下的Versal自适应片上系统(SoC)产品升级全新第二代，包括 面向AI驱动型嵌入式系统第二代的Versal AI Edge系列、面向经典嵌入式系统的第二代Versal Prime系列。 Sophgo SG2380 超星未来 https://www.novauto.com.cn/ sifive-intelligence-x280 the Vector Coprocessor Interface Extension (VCIX) 现代工作负载和应用程序通常需要最高的性能，但需要在有限的功率环境中实现。由于标准 RISC-V Vector ISA 和 SiFive Intelligence Extensions 提供的出色矢量计算能力，设计人员能够将各种专用 DSP 加速器功能整合到单个 X280 处理器设计中，从而保持更简单的系统设计和易编程性，同时仍能实现所需的性能和效率目标。 在某些情况下，例如需要对矢量数据进行高度密集计算的工作负载，设计人员仍然需要使用经过高度优化的定制加速器来从主处理器卸载这些任务。这种方法的挑战在于，定制加速器需要设计为与主处理器共存，这会带来一些微妙的设计后果，即加速器： RVV The X280 processor implements a 512-bit vector length architecture (VLEN), fully supporting the vector extension standard, with dynamic variable vector length operations. The vector ALU and load/store architecture data width (DLEN) is 256-bits. 云天励飞的新一代AI芯片DeepEdge10搭载公司自研的神经网络处理器NNP400T Meta的MTIA，该芯片将拥有 256MB 片上内存，频率为 1.3GHz，而 v1 的片上内存为 128MB 和 800GHz。新一代芯片的最新加速器由 8x8 个处理元件 (PE) 组成。这些 PE 显着提高了密集计算性能（比 MTIA v1 提高了 3.5 倍）和稀疏计算性能（提高了 7 倍）。这些提高一部分归功于架构的改进，另一部分归功于 PE 网格供电方式：Meta 将本地 PE 存储的大小增加了两倍，将片上 SRAM 增加了一倍，将其带宽增加了 3.5 倍，并将 LPDDR5 的容量增加了一倍。 Reduce的并行加速 CUDA 1. 采用Divergence的支持和Block同步来支持 2. 其他的深度优化：https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf DSA/ASIC 软件框架 为什么Pytorch开始制约AI了 这个想法比较超前了，但是已经有苗头了 pytorch已经显得臃肿，不适合大模型的特定需求 需求变化了，已经不是CNN的时代了，虽然在拼命搞静态图和并行库 一个开源框架统一市场后一段时间，就会开始因为臃肿而不适合新的需求，所以就需要 大胆重构 和numpy等通用接口差异越来越大，学习成本增加 为啥JAX在大模型巨头玩家中的份额这么高呢？ 大模型框架对原来的那种生态，可能要求没那么高，反倒对性能调试调优的效率要求高 JAX的简单，合理的功能抽象（更加符合第一性原理）才是用户的青睐点，pytorch也是靠这个取胜tensorflow的 使用Numpy作为接口，一个计算行业的底层的优秀抽象设计 简介、静态图、自动并行是核心优势 兜兜转转多年，静态图又要风生水起了？ 一个技术成熟到一定程度后，技术路线就会固定化，适合静态计算来提高效率 LLM时代AI加速芯片面临的挑战 算法需求 普遍使用MOE架构降低算力需求 高度定制化的集成度高的大算子 定制化的核心Attention加速算子：FlashAttention KVcache的压缩、加速等： Deepseek的Flash MLA 混合精度及量化 低精度支持及累加精度保证 多卡互联技术，包括快速的分布式all to all的性能，通信异步化，不占用计算核，最大化带宽利用率 更低的latency，更高的throughput 异步通信，动态执行 复杂的存储地址控制，复杂的MMU系统 Atomic的支持，用于复杂算法的reduction运算支持 同步和动态 问题&矛盾 通过传递依赖/同步信息到存储器（memory barrier） 优点：避免调用fence和trigger一个任务的leading和tailing latency 缺点： 为了异步调用和挂起，需要保存抽象的线程信息 占用大量的调度逻辑和流水线存储资源，动态调度的代价 源头统一调度和控制（mailbox 、global sync） 可能方案 通过预调度（prefetch）等预估流水线的latency进行提前trigger，减少latency 专用的异步进行同步的信号，同步信号和存储请求分离，同步信号可以不随流水线传递 把存储相关的程序指令存储到存储器端，控制器直接launch，类似一个可以执行指令的DMA 所有的单元（1D、2D、DMA、NoC）都变成可以执行程序的单元、可同步单元，通过launch和同步来控制 需要解决可执行程序的prefetch/preload需求 可以把整个程序通过计算图来抽象，明确所有launch的时机 通过预编译的“动态计算图”表达所有的launch和join，非常细粒度的调度 专用的“图”执行硬件单元，负责执行图、launch、同步、配置参数计算/生成和传递 标准的launch参数传递和配置方法，负责动态launch参数的配置 标准的标量流水线，通过专用硬件单元（专用指令）进行加速 内存墙 大容量、低延迟（高速）、随机访问 推理阶段，每个对话就必须加载全部的模型参数，只为了计算一个token的输出，算力利用率特别低 为了低延迟，不能使用低速存储替代，容量和速度的需求比例和传统DRAM不一致，可能需要 1T带宽 1T容量 ？ KV Cache的机制对存储容量和带宽的需求 跨界点通信开销 未来算法需求 动态全局随机访问的需求 MOE算法需求 稀疏算法 可能的方法 定义复杂指令，单指令支持操作数的Dequent，GEMM计算 设计很大的L1甚至是L0，用于存储Tensor计算的左右值 可能的方向 超大的dot变成多个小dot的并行 固定尺寸的dot，和输入无关的dot 所有模型使用一种固定的层，同一种组件，不要很多种类，很多奇怪的linear 仿真建模 建模的方法 使用python语法进行module的定义 python语法用来描述module之间的关系和 定义latch的存储器 定义module的算法和输入输出 配置一些固定的规格参数 针对throughput和latency的建模 传递的标准信息都是 inst 定义一个通用的TL（throughput latency）的Module 支持配置 thread 通道数量 处理inst的latency 处理inst的throughput：输入的通道的数量，输出的通道数量 内部stand-by的容量 TL可以被递归的进行组合 TL可以通过简单的python脚本定义输入和输出inst的转换规则 激励是一种特殊的Module TL可以支持自动的生成统计信息和time line图标 基于生产和消费的硬件模型抽象 定义资源，生产，消费的抽象 内置抽象功能的仿真： 容量 ， 带宽throughput ， 规格 ， 延迟latency ， 功耗 支持建立资源的依赖关系 内置的调度算法支持：轮询、贪心、自动调度探索的算法 用于仿真评估的现成运行库、工具库 定义 对象的生产 对象的缓存、处理、转换、调度、流转 对象的消费 flow plan，Place and router 规划 完全抽象后是可以求解一个方案，但是把所有的因素都考虑进去很慢 建模抽象，不能只是一个计算器，而是要能做到抽象的抽象，怎么快速得对问题进行分析？？ 建模过程： 定义工厂，生产和消费的产品以及速度，每个工厂就是一个硬件模块，并行工作 定义产品，产品可以分大小，工厂消费不同大小的产品 总线1：工厂「生产一个数据，消费一个数据」 总线2：工厂「生产一个数据，消费一个数据」 总线3：工厂「生产一个总线1，消费一个总线2」 微架构建模 延迟模型 发送请求，固定延迟，返回请求 硬件资源竞争 存储器的深度/credit/valid 地址相关的bank，读写口竞争，structure harzd RAW、WAW等 Data harzd 数据流模拟：cache line 优先级 算法 调度算法 数字电路仿真的抽象要素Logicsim 状态 Latch 并行 Launch 指令 Status ISA/UISA 作为module之间传递信息的抽象 记录各种状态的信息，状态信息的移动 AI计算需求的抽象 算子角度 信息维度映射，信息过滤，信息选择 ：Dot / GEMM => Join+Reduce 激活 ： ElementWise Broadcast 统计、动态选择、排序 ： Reduce 、G/S 随机数生成 是不是可以把所有的计算都抽象成查找表的方式 操作的方式由操作码+密码数来指定，实际上是定义一个查找表 操作码可以是指令的imm，密码数可以是约定的，也可能是提前载入的 乘法等价？ 信息加工的角度 一堆向量和另外一堆向量 交叉 计算相关度 存内计算只是保持一支数据更靠近计算单元 一堆向量的变换 数字电路的特性 电路 电路模块不能太大，所以要大量的相同功能进行组合 组合带来了复杂的判断逻辑，需要流水化提高throughput 流水化带来的latency需要用复杂的同步机制 微观上说，各种scoreboard，forwarding，多级调度，都是在控制pipeline的精密配合 宏观上说，mailbox，barrier是在更高的维度进行同步 复杂同步对软件和应用有一定的要求 以上几个因素根据不同的设计都有非常大的弹性，整体的方案设计难度大 各种硬件(比如roq，sram读写口等)的数量是有限制的，不像软件视角不感知这个冲突。 各种操作，运算都是有延迟的，都不是实时单拍完成的 加速芯片 存储和传输 大量，大力度的并行搬运数据 各个级别的速度不一样 各个级别的暂存容量大小不一样 并行计算 大量的充分的并行计算 大量的数据复用 卷积，矩阵乘 同步 硬件微架构（电路）的难点和优势 难点/不适合 难以做复杂的算法（调度、分析、统计） 难以做逻辑深度比较深的计算 latency的长和不确定，模块间的时间不确定性 需要大量实现“异步”逻辑,消耗大量的面积 需要很大的buffer存储用于状态footprint,消耗大量的面积 增加外部控制的复杂度， PC需要对所有的操作异步化，降低实时性 为了高频、增加很多级流水线 浪费大量用于同步的逻辑 容易撞到功耗墙 latency的延长和调度不确定性 优势/适合 并行任务非常适合，适合简单的、重复的、大量的计算，比如MMU 带宽足够，可以做大量的并发处理 实时性好，可以做到cycle级别的调度 可以利用数字电路表达的特点加速 对2的幂次的乘除运算 固定算法的hash，mmu等处理 查找表计算：cuda的LOP3.LUT指令 通过软件的静态预先计算，最大化压缩硬件需要处理的表达空间 GMP 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最小化 AI加速芯片上的2D单元 卷积天然的数据复用度是Dot的9倍，对于芯片的压力更小 算力缩放是一个非常重要的问题，涉及架构各代之间的稳定性，保护客户的价值 L1/L2/L0 怎么支持reshape或者swizzel 不同的layout （NHWC等） 以及BPI BPK FF 不同数据精度的支持和混合计算 L0 L1 Fusion的支持 存储的mapping 利于运算和fusion feature*weight vs weightT*featueT 大2D Dot计算的表达 指令表达为小尺寸，对不同的计算需求（1D、2D）的fusion比较友好 但是小尺寸，需要每个单元的L0进行broadcast来复用给其他单元，节省L1的带宽需求，增加了2D的复杂度 小尺寸对load/store的要求也比较高 需要2D的计算表达有一个好的抽象，同时配合非常精密的同步，解决上述问题 为了在各代的代码不发生变化，需要硬件/编译器有自动合并/优化的能力 cuda因为需要兼容simt的编程，复用warp内thread的寄存器，同时为了拼接更多的硬件做一个更大的2D（为了PPA好） 整体上单个thread表达的数据计算尺寸比较小 硬件自动进行合并成大的2D指令 硬件自动处理同步 硬件自动处理数据搬运 因为需要协调多个sm subcore，甚至是block group，灵活性一样是个大问题 XXU4.0利用多个thread之间的手动给的 footprint order信息来组织多个thread的数据共享 左数每次都从L1load重复load，右数缓存在L0，并且各个subcore复用 2D-1D-2D的Fusion 涉及到多次计算后结果的数据精度（位宽）不匹配的问题，需要合并和拆封来保持操作数的宽度一致 8bit int8 需要 64个组成 512bit，而32bit只要16个 不同的操作的输入输出，可能需要有装置的操作 layout的区别：NCHW NHWC 卷积的stride不等1 预测：2D算力会无限的膨胀下去，1D和2D的比例会比较固定 精细化控做到L1，则所有的抗latency的容量就会开在L1，消耗：L1_base x latency，如果是在L0，则L0_base比L1_base小很多。节省很多面积 swizzle和renaming都可以解决L0的bank conflict，renaming会消耗大量的面积 物理限制 算力和存储器（SRAM）都不能做到单体很大 很大的单体有利于软件使用方便 多个小的存储器实例（bank）组成一个大的抽象存储器 物理单元之间很难方便做到cycle by cycle的同步，必须使用额外的同步逻辑 硬件电路虽然可以并行，理论算力很高，但是latency是必然存在的 NV Blackwell 从PTX、SASS可以看到AI面对的问题和很多解决方法都是共同的。 https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#tensorcore-5th-generation-family-instructions https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html#blackwell-instruction-set L0的memory变化 显式的手动管理tensor memory ，是不是2D/1D更像是独立的一个NPU/DSA单元？ Tensor计算直接使用专门的 Tensor memory，做memory to memory的Tensor计算 ，且支持软件动态分配 1D 可以直接load/store这块memory 这样让L0可以动态，目的估计是 破除 2D/1D 每写一种L0 fusion就要写一份code 示例，Fusion算子可以隔离开发 更大的矩阵形状 128/256，降低2D带宽需求 支持 TPC 内 2个SM之间的2D 协同计算，看起来是2个core支持的读数据复用 支持 OCP-MX的micro-scaling 标量thread issue Tensor计算指令，不再是 SIMT style的多个thread 直接支持卷积指令，支持weight-stationary GEMM，带上mask bit 表达哪些是padding 的0 GMP 背景 适应未来的AI计算需求 存储足够量的权重，但是明显的热点内容访问 强动态性，大范围、多次的随机动态访问 节能、低带宽需求，高效率数据流 低延迟 软件定义硬件，硬件尽量简单、透明、底层抽象、灵活 目标 软硬件结合 大量依靠软件优化，发挥最大的物理效率，同工艺下架构效率达2倍 算法和硬件协同优化，同模型精度条件下效率达4倍 DRAM或者多机的不确定数据延迟直接整合到算法处理，硬件不做竞争 硬件采用固定的LUT计算（可能不能等价到矩阵乘法，甚至是乘法本身） 全模型网络级别优化，利用编译器对整个模型进行搜索优化，生成静态计算图 指令控制流水线 指令明确指定指令的调度、L0 Cache的使用、依赖关系的建立和解除 原生支持动态算法： MoE 动态算法 统一的异步通讯管理方案 整个系统有大量的不同的通讯和同步机制 流水线内的credit，L1的数据缓冲 算力核之间的数据交互核同步，NoC的各种协议 L2/L3的复用 分布式栈：网卡的片上调度，网络的延迟不确定性，通路的复用 outstanding/各种缓存的管理和设计 异步的调度 launch控制及指令加载加速 统一的数据流拆分模型 提供统一的编程模型对数据流进行描述 硬件加速的数据流动态计算，减少冗余且高成本的除法/模运算的地址计算 通过自动的 预计算 和特殊硬件加速 自动处理的原子操作，以消除写入全局内存时的warp级串行化 自动进行 乒乓双缓冲 机制 权重和临时数据（KV Cache ..) 分离的架构 权重使用固定的通路 两种数据使用不同的数据流模版进行设计 算力 Binary Lut 方案，减少算力能耗和面积，降低数据搬运量? CIM 专用权重通道，近存架构，存内架构， 存内计算 ? 1D动态算力架构 一维计算阵列，动态性体现在一维的长度上，对应到sequence的长度 二维数据复用，数据通道提供精密的排布和数据复用的调度 编译出向量指令流（LD MUL ST），支持批量动态配置，两块指令流之间流水可以重叠---> DSA？ 例化固定数量的 L1读写计数器，用于同步dma和Mac， dma和Mac，自动从l1 加载指令 问题： 1D指令表达大算力2D VLD VST MLD MST VMUL VMUL_reduce MUL_join 等等 指令 LD/ST的自动mbarrier 软件管理L1的所有bank，每个bank的每个地址都记录一个count，启动初始化的时候初始化count 架构考虑 动态性的表达 硬件竞争的管理 灵活性，扩展性，从edga到集群 自举，所有单元支持自配置，自启动 线程内的依赖都是静态的软件调度，软件直接调度流水线，减少硬件的调度 架构方案 图：编译整个动态计算图，支持 Fork（launch） join sync 平铺：按照可编程硬件单元进行编程，为每个单元生成一定数量的逻辑线程， 支持 sync ... ... 规格 指令流 load和fetch 基于图的信息，和数据流一样得方式，需要发命令和同步 ISA Launch Sync 总结 灵活、低延迟的异步/同步机制 简单、透明的硬件拓扑，软件的深度控制 架构/微架构 设计 标量寄存器和向量寄存器统一，支持自动进行转换 异步单元（SP-PU-L1-DMA）之间都采用异步机制，依赖转移到异步目标 统一的同步机制 静态分配同步资源 原生软硬件支持动态图的执行 LD/ST 避免使用fence功能 Launch：fork 资源初始化（同步资源，各种存储器，状态） launch/signal/wait:join launch pu instrution:  write_back_id  local_id write_back atomic add/sub wait instruction: local_id local_id GE LE counter wait remote instruction: remote_id local_id GE LE counter 灵活性&性能 支持灵活的数据尺寸 支持灵活的算子 混合精度的scale 1D/2D并行 全局Reduce算子，排序算子 L3 Atomic 支持灵活的本地数据复用，尽量减少数据搬运数据 尽量避免算力浪费 leading tailing latency LD/ST inflight delay，DMA Delay ping-pong pipeline kernel启动latency kernel预加载 kernel间的gap kernel间的本地数据复用？ 多个算子的fusion 数据没有ready、需要fence引起的气泡 无缝的同步机制 m-barrier 不对齐、尾数 MAC Vector计算单元支持乱序和并行 软件编译指定明确的VR寄存器依赖，RO WO属性 自动拆解大指令成Vector指令，并行执行 软件静态排布VR寄存器，生成依赖关系，申请和释放管理 Vector指令支持标量对其进行动态配置 Software 使用独立的硬件仿真软件，不依赖硬件仿真运行 ISA 标量 RV64i 向量 VLD VST VMUL VADD REDUCE_ADD REDUCE_MAX REDUCE_MIN VMUL_REDUCE_ADD MLD MST 张量 GEMM 左/右数：128个=7位 + 3位扩展矩阵  = 10位   输出：7位 Opcode fence L1 cache line 计数 VR 软件管理VR data hazard？ VR之间的依赖？ 增加指令的表达信息 软件分配VR 还是 硬件rename ing ，解决bank冲突？ 利用 VR count？软件进行管理依赖关系？ LD ST 计算 的三类指令之间可以并行，通道内部没有必要并行 因为硬件资源没有特殊性，不会因为并行而减少气泡 默认，GEMM指令一定要在前面LD指令之后执行 默认，ST指令一定要在前面计算指令之后执行 L1 软件管理 cache line cache line硬件计数，自动异步等待 针对L1 CacheLine的编程？ 软件指定Load到L1 cache line的位置和有效长度 向量指令按照cache line的粒度和mask来执行指定的计算 这整个流程都是提前编译好，从L1-L1都是提前确定的 针对不同MNK大小需求，可以通过标量指令来快速配置，支持动态性 DMA 附带一片ROM空间，可以在Launch Kernel的时候主动加载一片指令编码，和当前kernel绑定的 支持一个cmd端口，接受其他核心的控制命令 DMA作为对外的接口，通过配置kernel的linear copy kernel作为launch的起点 ？ 简化系统复杂度，通用的功能组件，避免开发专用的launch constructor 增加到128个标量寄存器 支持RV64扩展？ 增加指令流控制 增加配合/加速向量单元的定制指令？ 通过兼容RiscV-V的指令来实现？？ 架构_微架构 架构/微架构 设计 标量寄存器和向量寄存器统一，支持自动进行转换 异步单元（SP-PU-L1-DMA）之间都采用异步机制，依赖转移到异步目标 统一的同步机制 静态分配同步资源 原生软硬件支持动态图的执行 LD/ST 避免使用fence功能 Launch：fork 资源初始化（同步资源，各种存储器，状态） launch/signal/wait:join launch pu instrution: write_back_id local_id write_back atomic add/sub wait instruction: local_id local_id GE LE counter wait remote instruction: remote_id local_id GE LE counter 灵活性&性能 支持灵活的数据尺寸 支持灵活的算子 混合精度的scale 1D/2D并行 全局Reduce算子，排序算子 L3 Atomic 支持灵活的本地数据复用，尽量减少数据搬运数据 尽量避免算力浪费 leading tailing latency LD/ST inflight delay，DMA Delay ping-pong pipeline kernel启动latency kernel预加载 kernel间的gap kernel间的本地数据复用？ 多个算子的fusion 数据没有ready、需要fence引起的气泡 无缝的同步机制 m-barrier 不对齐、尾数 MAC Vector计算单元支持乱序和并行 软件编译指定明确的VR寄存器依赖，RO WO属性 自动拆解大指令成Vector指令，并行执行 软件静态排布VR寄存器，生成依赖关系，申请和释放管理 Vector指令支持标量对其进行动态配置 Software 使用独立的硬件仿真软件，不依赖硬件仿真运行 ISA 标量 RV64i 向量 VLD VST VMUL VADD REDUCE_ADD REDUCE_MAX REDUCE_MIN VMUL_REDUCE_ADD MLD MST 张量 GEMM 左/右数：128个=7位 + 3位扩展矩阵 = 10位 输出：7位 Opcode fence L1 cache line 计数 VR 软件管理VR data hazard？ VR之间的依赖？ 增加指令的表达信息 软件分配VR 还是 硬件rename ing ，解决bank冲突？ 利用 VR count？软件进行管理依赖关系？ LD ST 计算 的三类指令之间可以并行，通道内部没有必要并行 因为硬件资源没有特殊性，不会因为并行而减少气泡 默认，GEMM指令一定要在前面LD指令之后执行 默认，ST指令一定要在前面计算指令之后执行 L1 软件管理 cache line cache line硬件计数，自动异步等待 针对L1 CacheLine的编程？ 软件指定Load到L1 cache line的位置和有效长度 向量指令按照cache line的粒度和mask来执行指定的计算 这整个流程都是提前编译好，从L1-L1都是提前确定的 针对不同MNK大小需求，可以通过标量指令来快速配置，支持动态性 DMA 附带一片ROM空间，可以在Launch Kernel的时候主动加载一片指令编码，和当前kernel绑定的 支持一个cmd端口，接受其他核心的控制命令 DMA作为对外的接口，通过配置kernel的linear copy kernel作为launch的起点**？** 简化系统复杂度，通用的功能组件，避免开发专用的launch constructor 增加到128个标量寄存器 支持RV64扩展？ 增加指令流控制 增加配合/加速向量单元的定制指令？ 通过兼容RiscV-V的指令来实现？？ 算力单元 算力单元详细设计 把 [[Pipe]] 第 6 节 §208 的算力单元抽象展开到 RTL 起手前。每个计算实例 = 一条 engine pipe 实例。 0. 与 DMA 的同构关系 算力单元与 [[DMA]] 结构同构 ：都是"engine pipe 集合 + sync pool + 控制接口 + 共享数据 fabric master"。差别仅在执行通路： DMA 算力单元 执行通路 读 → transform → 写 读 → 计算阵列 → 写 数据端 被动单元（L1、DRAM） 被动单元（L1、DRAM） 内部状态 transform buffer L0 寄存器组 cmd_desc 专属字段 xform_flags, xform_params compute_op, loop, dtype, accum 下文凡是结构与 DMA 一致的，引用 [[DMA]] 不重复展开；只描述差异。 1. 顶层参数 参数 含义 默认 备注 N_CU 计算实例数（每个 = 一条 engine pipe） 4 可异质（见 §13） N_SYNC sync counter 数 32 Q_MAX 单实例命令队列硬件上限 16 2 的幂 K_IN / K_OUT / K_SIG cmd_desc 列表上限 4/4/4 N_L0_VEC 每实例 L0 向量寄存器数 8 N_L0_TEN 每实例 L0 张量寄存器数 4 VEC_LANE 向量阵列 lane 数 16 INT8 lane MAC_M / MAC_K 2D MAC 阵列尺寸 8 / 8 systolic MAC_ACC_W MAC 累加器位宽 32 INT32 OUT_TYPE_SET 可选输出位宽 {8, 16, 32} round/sat 软件控 LOOP_W inner-loop 计数器位宽 16 STRIDE_DIM / STRIDE_W / SHAPE_W 同 [[DMA]] 3 / 32 / 24 其他 COUNT_W / HANDLE_W / ADDR_W / NOC_W 同 [[DMA]] 2. 顶层框图 ┌────────────────── 算力单元 ────────────────┐ ctrl_s ──────────► │ ctrl 解码 + 路由 ─► CU Pool Mgr │ ctrl_s ◄────────── │ ─► Sync Pool │ │ │ │ ┌───────── CU[0] ─────────────┐ │ ctrl_m ◄───────────┤ │ cmd queue (SRAM, Q_MAX) │ │ ctrl_m ──────────► │ │ 完成计数器 │ │ │ │ 执行 FSM │ │ │ │ in/out 地址生成器 │ │ │ │ L0 vreg[] treg[] │ │ │ │ 计算阵列(标量/向量/2D MAC) │ │ │ │ Notify FIFO │ │ │ └─┬─────────────────┬──────────┘ │ │ │ │ │ │ CU[1] ... CU[N_CU-1] │ │ │ │ │ noc_rd ◄───────────┤ NoC 读仲裁 ◄──── 各 CU 的 LOAD 通路 │ noc_wr ◄───────────┤ NoC 写仲裁 ◄──── 各 CU 的 STORE 通路 │ └─────────────────────────────────────────────────┘ 3. 顶层端口 完全复用 [[DMA]] §3 的端口。差别： noc_rd_* / noc_wr_* 可对接任意被动单元（L1、DRAM），目标由 cmd_desc 内的 in/out handle.unit_addr 决定 ctrl_s_* / ctrl_m_* 协议完全一致 4. cmd_desc 字段编码 总宽 ≤ 1024 bit。与 [[DMA]] §4 同构，差别在 op + 专属参数区 。 偏移 位宽 字段 说明 0 4 op 0=SCALAR_ALU, 1=VEC_ALU, 2=VEC_REDUCE, 3=MAC_2D, 4=ELEMENTWISE_ACT(ReLU/GELU), 5=SOFTMAX_PARTIAL, ... 4 4 dtype_in 0=I8, 1=I16, 2=I32, 3=FP8(预留), 4=FP16(预留) 8 4 dtype_out 同上 12 2 accum_mode 0=覆盖, 1=累加（C += A*B）, 2=累加+饱和 14 2 act_kind ELEMENTWISE 时的激活函数选 16 16 loop_count inner-loop 次数；硬件按此自动 step 地址生成器 32 48 src_a_base byte_addr（落在某条被动单元 pipe 内，见 [[L1]] §4.1 / [[DRAM]] 地址映射） 80 96 src_a_stride[3] 176 72 src_a_shape[3] 248 48 src_b_base MAC_2D / VEC_ALU 用 296 96 src_b_stride[3] 368 48 dst_base 416 96 dst_stride[3] 512 72 dst_shape[3] 584 8 l0_a_reg A 操作数落到哪个 L0 寄存器 592 8 l0_b_reg 600 8 l0_c_reg 结果累加目标 608 4+4+4+4 n_in / n_out / n_sig / _ 624 96 in_list[4] 每项 {handle:16, delta:8} 720 96 out_list[4] 816 96 sig_list[4] 912 112 reserved/padding src_b_* 仅对二元算子（MAC / VEC_ALU）有效；一元算子（激活、reduce）写 0。 5. 控制总线协议 完全同 [[DMA]] §5。多 beat PSEND 同样 8 × 128 bit。opcode 表共享。 6. CU Pool 管理器 结构同 [[DMA]] §6。差别： meta[N_CU] 多一个 cu_class 字段（标量 / 向量 / MAC，若异质例化） PALLOC 时 params.cu_class 必须与本实例 class 匹配；不匹配视为软件错（行为未定义） 7. CU 执行通路（×N_CU） 7.1 命令队列 完全同 [[DMA]] §7.1。 7.2 完成计数器 完全同 [[DMA]] §7.2。 7.3 执行 FSM ┌────────────────────────────────────┐ ▼ │ ┌─────────┐ queue 非空 ┌────────┴───┐ │ IDLE ├─────────────────────► │ FETCH │ └─────────┘ └────────┬───┘ ▼ ┌────────────┐ │ WAIT_IN │ └────────┬───┘ all in 到位 │ ▼ ┌────────────────────┐ │ ISSUE 流水 │ │ ├─ LOAD (L1→L0) │ │ ├─ COMPUTE │ │ └─ STORE (L0→L1) │ │ (loop_count 次) │ └────────┬───────────┘ ▼ ┌────────────┐ │ NOTIFY │ └────────┬───┘ ▼ ┌────────────┐ │ DONE │ count++ tail++ └────────┬───┘ └──► IDLE 状态 时长 行为 IDLE 1 cycle FETCH 1 + (STRIDE_DIM-1) cycle 读 cmd_desc + 预算 stride wrap_delta（同 DMA §7.4） WAIT_IN 不定 顺序 COUNT_QUERY in_list ISSUE 流水，见 §7.6 NOTIFY (n_in + n_out + n_sig) cycle 入 Notify FIFO DONE 1 cycle 7.4 L0 寄存器组 项 规格 vreg[N_L0_VEC] 每个 VEC_LANE × dtype_in_bits 宽；存 vector 操作数 / 结果 treg[N_L0_TEN] 每个 MAC_M × MAC_K × dtype_in_bits 宽；存 tile 操作数 端口 vreg: 2R/1W；treg: 2R/1W（FMA 类够用） 复位 不复位 （undefined），软件必须先 LOAD 才用 共享 实例私有，不跨 CU 共享 实现：vreg 用 register file（small），treg 用 SRAM macro（一个 treg ≈ 8×8×8 = 512 bit，N_L0_TEN = 4 时 2 KiB）。treg SRAM 端口配 2R/1W 满足 MAC 2 输入 1 累加。 vreg 与 treg 物理分离 —— 二者访问宽度差很大。 7.5 计算阵列 三类阵列共存（每个实例可裁掉某类，按 cu_class 决定）： 7.5.1 标量 1 lane，宽度 max(dtype_in, dtype_out) ≤ 32 bit op：ADD/SUB/MUL/SHIFT/CMP/SELECT/BIT latency 1 cycle，throughput 1/cycle 7.5.2 向量 VEC_LANE 路并行 INT8 ALU + 16-bit accumulator option op：lane-wise ADD/SUB/MUL/MAC/CMP，跨 lane REDUCE_SUM/MAX/MIN latency 1-2 cycle，throughput 1/cycle elementwise activation：LUT 表（4 KiB ROM/SRAM）实现 ReLU / GELU / sigmoid 等 7.5.3 2D MAC MAC_M × MAC_K systolic 阵列 单元：INT8 × INT8 → INT16 partial → 累加进 MAC_ACC_W 寄存器 输入：A treg (M × K), B treg (K × N，本设计 N = MAC_M)；权重和激活方向可在 cmd_desc 内对调 输出：C treg (M × N)，累加进已有 treg（accum_mode） pipeline depth = MAC_K + 2 cycle 每 cycle 注入一行 A、一列 B；MAC_M = MAC_K = 8 时 8 cycle 填满，第 9 cycle 起每 cycle 一个稳态结果 下溢 / 上溢 / round：cmd_desc.dtype_out 与 accum_mode 决定输出阶段的 saturate + truncate / round_to_nearest。专用硬件块（输出后处理 stage）实现。 7.6 LOAD/COMPUTE/STORE inner-loop 流水 把一条 cmd 拆成 loop_count 次 micro-op，三级流水： cycle: N N+1 N+2 N+3 N+4 N+5 │ │ │ │ │ │ LOAD op0 op1 op2 op3 op4 op5 COMPUTE op0 op1 op2 op3 op4 STORE op0 op1 op2 op3 L0 寄存器组作为级间 buffer：LOAD 写、COMPUTE 读 + 写、STORE 读 软件用 l0_a_reg / l0_b_reg / l0_c_reg 指定不同 op 用哪些 reg，硬件按 op_id % N_L0_REG 做 cyclic 调度避免读写冲突 leading 2 cycle 只 LOAD（COMPUTE/STORE 等数据）；tailing 2 cycle 只 STORE inner-loop 内不可中断 / 不可被外部 stall（除 NoC backpressure） 7.7 in / out 地址生成器 每实例 3 个独立累加器（A、B、C）。复用 [[DMA]] §7.4 的累加器实现。差别： 目标地址落在被动单元（L1 / DRAM）的地址空间内 inner-loop 由 loop_count 驱动，与 shape[3] 一起决定终止 地址生成单次给一个 NoC 请求；burst 长度 = 一拍 L0 寄存器宽度 / NOC_W（典型 1 beat） 7.8 Notify FIFO 完全同 [[DMA]] §7.6。 8. Sync Pipe Pool 完全同 [[DMA]] §8。 9. 共享资源仲裁 9.1 NoC 读 / 写 结构同 [[DMA]] §9.1。NoC 主端可对接 L1 / DRAM；对接 L1 时可按低延迟假设缩 outstanding 深度，对接 DRAM 时按 DRAM 延迟特性配置。 9.2 ctrl_m 主端 完全同 [[DMA]] §9.2。 9.3 ctrl_s 解码路由 完全同 [[DMA]] §9.3。差别：PALLOC 时 params.cu_class 决定走哪个 cu_class 的子池。 10. 跨子模块协议要点 10.1 CU ↔ pending wait CAM 同 [[DMA]] §10.1。 10.2 CU ↔ ctrl_m 同 [[DMA]] §10.2。 10.3 CU ↔ NoC（对接被动单元） LOAD：CU 发 noc_rd_req 给被动单元（L1 / DRAM，master），地址按 §7.7 生成；resp 数据按 req_id 路由到 CU 的 L0 写口 STORE：CU 发 noc_wr_req + wr_data 给被动单元；不等 ack（fire-and-forget），NOTIFY 中的 COUNT_DELTA 由 CU 内部 FSM 在 STORE 流水尾部排出（保证落到目标单元后才发，由 NoC 同 id 同地址保序保证） 10.4 同 cmd_desc 内 in 与 out 包含同一 pipe 参见 [[Pipe]] §289：合法用法。硬件执行顺序： WAIT_IN 阶段查询该 pipe 的 count ≥ in 阈值 LOAD 阶段读被动单元（数据已在） COMPUTE STORE 阶段写被动单元 NOTIFY 阶段：对该 pipe 先发 COUNT_DELTA(-in_delta) ，再发 COUNT_DELTA(+out_delta) 被动单元端按到达顺序应用即可。In-place 累加的语义保证依赖 NoC 同 id 同地址保序。 11. 复位与启动 信号 / 状态 复位值 同 [[DMA]] §11 全部项 L0 vreg / treg 不复位 （undefined） 计算阵列 pipeline 寄存器 复位 MAC 累加器 0 L0 不复位的影响：未 LOAD 就 COMPUTE 是软件 bug，硬件不阻止；结果未定义。 12. 调试可观测点 同 [[DMA]] §12，加： 信号 宽度 用途 dbg_cu_state[N_CU] N_CU × 3 每 CU FSM 状态 dbg_mac_active[N_CU] N_CU MAC 阵列在跑 dbg_inner_loop_idx[N_CU] N_CU × LOOP_W 当前 micro-op 序号 dbg_l0_read_hazard[N_CU] N_CU LOAD/COMPUTE 同 reg 冲突计数（应为 0） dbg_pipeline_stall_cycles[N_CU] N_CU × 16 NoC backpressure 累计 stall 13. 待规格化参数 N_CU 与异质组合（如 1 个 2D MAC 实例 + 2 个向量实例 + 1 个标量实例 vs 同质） VEC_LANE / MAC_M / MAC_K 实际规模 —— 直接决定算力 dtype 支持集合（INT4 / FP8 / FP16 / BF16） MAC 阵列拓扑（systolic 1D vs 2D，weight-stationary vs output-stationary） L0 寄存器组的实际端口数（FMA 用 3R+1W 还是分段 2R+1W） inner-loop 流水级数（3 级 vs 4 级，多一级换 stall 余裕） 输出 round / saturate 硬件块的位宽 是否例化激活函数 LUT 共享还是 per-CU 独立 14. 关联回 Pipe 抽象 抽象层概念 本文档落点 engine pipe 契约 §7.1 队列 + §7.2 count + §7.3 FSM sync pipe 契约 §8 落地 A §6 + §8 各自独立 bitmap 落地 B §3 ctrl vs noc 物理分开 落地 D §8 sync 独立 特性 3（单 pipe 不分叉） §7.3 FSM 单 channel 严格串行 特性 6（硬件零检查） §5 / §6 / §7.4 L0 未初始化不报警 特性 7（参数创建即不可变） §6 alloc 后 cu_class / queue_depth 定死 特性 8（三类 pipe） engine 在 §7，sync 在 §8 §289（in/out 同 pipe） §10.4 NOTIFY 顺序 DynamicGraphMultiProcessor架构 Dynamic Graph Multi Processor 架构 背景 Etched提出，GPU在过去四年间效率并没有变得更好，只是变得更大了：芯片每平方毫米的的TFLOPS几乎持平。 「干净数据+大模型」和「脏数据+大模型」的效果，不会有太大差异。 Etched团队表示，H100有800亿个晶体管，却只有3.3%用于矩阵乘法，这种大模型推理时最常见的运算。只支持Tranformer的Sohu芯片FLOPS有效利用率超过90%（GPU大约是30%） 在前Scaling Law时代，我们强调的是Scale Up，即在数据压缩后争取模型智能的上限；在后Scaling Law时代，需要比拼的则是Scale Down，即谁能训出更具「性价比」的模型。 核心特点 针对AI应用，极高的规格参数 抛弃传统计算机的特性，只支持Int8，加减乘，固定的数据流水pattern，固定的计算管线 算子的输出输入精度更低，计算精度更高，自动混合精度 脉冲阵列，多周期指令（超级大的2D指令）=> 提高算力密度？？ GPU的核心问题就是：管理DRAM的延迟和有限的带宽 所有的硬件单元（包括NoC，Cache等）都在执行一张计算图，通过统一的图关系指令（Fork/Join）来运行，软件可以精细得控制所有的硬件单元，甚至是NoC得flow control行为。 软件控制内存一致性问题，传递信息到NoC/Cache，硬件不做自动化处理 传统上架构，内存一致性（consistency corhenrency） 需要通过fence+sync功能来保证，非常麻烦和易错。 GMP通过显式的同步模型（fork/join）来管理同步关系 线程内的利用静态的依赖关系来尽量避免fence的使用 线程间的编译器自动掺入fence指令 全异步，各种单元之间可以主动同步，每种单元（IP）都可以执行一些指令，都是一个cpu核，执行自己的代码和调度，区别于传统设计，从IP一般都是通过主IP进行控制，比如cache单元被处理核通过cache_hint进行配置 边缘端的越级大模型，极度定制的芯片 固定的计算单元的组合（DSA） 整个SOC级别的动态调度 以NOC为编程中心的调度和数据流 标量单元，向量单元，张量单元灵活组合，操作数（寄存器）可以灵活转换 充分把合理的计算放到合理的单元，节省面积和功耗：一个标量x一个张量 充分利用软件的编译 DSA的效率，GPU的灵活性 动态/静态 动态和静态的范围 全静态（graph）调度逻辑，包括DRAM都静态调 通过LLC隔离DRAM的动态性，LLC以下全静态 L3->LLC->DMA->L1->L0->MALU 静态图 利用fork/join（硬件加速）描述所有的并行性，比如 L1/L0的服用，DMA的操作后同步 有并行就有调度器，有调度就有缺陷 动态 Root 调度 & 核内调度 & 流水线调度 数据流、NoC、Cache Fork Join 充分利用CPU领域的灵活性，RISCV多核，加速的硬件调度指令（地址计算，IP配置等），同步指令，ForkJoin等 静态 NPU流水线，NPU Kernel，1D 2D指令及组合 DSA Confige Atomic Reduce专用加速硬件，近存储计算，压缩解压缩 LD/ST+Relu专用硬件 特点 NPU的流水线延迟固定，包括LD/ST（编译器已知），都是从L2/L1 <-> L2/L1 都是MLI指令，可以打包成HLI指令，不支持跳转等 不支持不对齐？ 不对齐是不是延迟就未知了 同时只能执行一个任务（一堆MLI指令） 指令堆会有 Leading/Tainling 时间，scheduler感知得到，可以进行调度，前后两条可以自动pipeline，NPU本身不考虑依赖 强大的软件仿真和验证平台，减少硬件的debug需求，大大简化硬件的调试/检查/报错功能 和主流ASIC区别 MLI不是硬件拆分？由编译器拆分，为了减少硬件的难度 scheduler控制NPU的指令搬运、 明确latency 支持打包成HLI 方案 多GM之间共享数据及同步 scheduler可以做及其动态的交互和同步 内存墙 自动精度缩放，权重存储的是<8bit，DMA搬运的时候变成8bit，每个参数都可以进行不同bit的缩放 Launch/Sync/DataFlow 高效的实现Launch和Sync，抽象出Fork/Join指令 主动Fork和Join的不一定是同一个对象，可以A fork B，C join B fork等价与launch，join等价与wait 可以一次性fork多个，也可以一次性join多个 硬件单元存储一个fork出来的任务的队列进行调度执行，不同单元的并行度不一样 支持片上动态的Launch和Sync 在NoC里面同时实现Launch控制，同步，数据流 统一在一个地方调度，所有单元并行，一起执行一个大的Graph Kernel 所有单元的指令特点：latency短且固定，只存在一级调度，调度后不可阻塞 精确的调用（launch）和同步带来的好处 突破内存墙：大算力->大延迟->大的in flight存储，大的pipeline存储，通过精密的同步和异步最小化存储器的使用 突破数据墙：越低的精度，需要的带宽比例越大，低精度算力面积呈指数下降，但是带宽需求是线性下降。只有精确的同步和调度才能最优化带宽资源的使用。 提升算力密度：算力单元的简化，静态化 大幅提升随机存取和计算的能力，是智能的重要指标 通过高效的控制和同步，最大化提高片上SRAM的复用率 总的片上的SRAM的容量很大，但是分散在很多小的L1，L1之间可以快速的同步数据，没有L2、LLC 分散的小的L1等价一个很大的cache，缓存下GEMM的整个右支 在有限的面积下，既满足了容量的需求，又满足了带宽的需求 支持非常细粒度的同步，比如 L1的各个bank之间的读写的同步，计算核的各个thread内的指令间同步，计算核内的各种Engine（ALU，RegRead）的同步，MALU的ld Cal St的同步。 动态graph，动态体现在动态launch，子图的动态的调度，循环调度，动态高效的同步支持，子图是完全静态的。 不同的算子需求下，都能写出高性能的算子，DSA的架构可能高性能，但是通用性差。完备指令架构很难做到优秀的PPA。gpu有强大的gs来适应大部分的不同数据流需求，simt的灵活调度提供充足的指令流，simi的指令尽量避免（软件解决）各种data和structure的hazard，减少流水线的复杂度。 通过紧密的同步控制，结合编译器，可以在通用性的前提下做到DSA的性能。通过减少流水线的fence需求，和硬件资源空泡的概率，特别是TensorALU、LD、ST等长latency的单元。传统的解决办法是通过多线程来提供充足的standby的微指令给硬件单元自己调度，避免空泡。 问题 支持多线程，为了充分利用硬件，不空泡 编程复杂 需要复杂的同步pipeline控制 需要复杂的编程L2L1复用 不支持多线程 通过graph指令，配合编译器，单线程实现并行 专用的加速指令，加速DMA操作，地址计算，同步/异步等 DMA和NPU等外设，可以直接读取Scheduler的Inst SRAM的方式执行，避免配置参数拷贝 相对于scheduler的外设，通过fork和join到对应的外设来支持异步编程，避免复杂的多线程语义 物理上可以有超线程并行的硬件，共用一个线程的寄存器和堆栈，通过fork，join软件控制表达并行和依赖 VMM怎么和VLD VST并行，pipeline kernel种类太多，体积大怎么办？？ 各个IP同步的开销太大 HWSync硬件指令，同步资源集成到RISCV CPU里面去，独立的中断信号线 各个IP之间通过mailbox相互trigger，mailbox可以动态配置，但是像pipeline那样太复杂了，可以支持forkjoin抽象 数据的写透（fence），流水线都很长，导致leading tailing太大 每个NPU的任务都有一样的（明确的）leading和tailing时间，所以可以在scheduler直接流水发射 Graph Multiprocessor = GM 同一片代码多个并行传递参数问题 fork支持参数的计算，批量fork自动累加传递的参数，类似于for功能的批量launch， 一个kernel最小的fork出来的执行单位表达类似于一个gemm指令 架构 32bit指令宽度 标量+向量+2D混合指令 体系架构 支持未来的超动态网络？加速动态网络？ 通过静态图的自动动态调度执行，能高效进行基于数据的动态网络 动态的launch/fork和join 编程的抽象 SIMT的编程对象是一个thread，并且要求所有的任务拆解成一堆一样的thread GMP更像是MGMT，Multi Graph Multi Thread，编程的基本对象是Graph，Graph可大可小，可以组合，可以动态执行（Launch），可以被多个计算核并行加速 专用的fork/join指令和硬件记录graph的launch和结束 单个处理单元内部的不是多个线程free竞争run，而是经过编译器严格规划的，graph是processor运行的基本单元，graph内部的多个分支的并行需要精密的同步， 每个fork出来的subgraph对应的处理级别都是明确的，通过PU执行的graph，fork出来的subgraph，就是hw thread级别。 同步指令只能在同级别的graph之间进行？ fork和join支持“无主模式”，fork出来之后主graph消失，每个子graph记住自己和哪些graph是并行的，利用这个信息进行同步，直到执行join之后，回到上一级graph，此时的join类似于barrier。 counter : wait 指令总是等待自己的一个counter值，signal远程的一个counter值 ，怎么实现灵活动态的同步需求？ 预先的pull一次，配置上动态wait信息，然后再wait signal 通过launch来指定动态同步信息 动态性是通过launch来体现的，动态信息是launch的时候统一生成的 一次launch出来的并行sub graph可以有一个graph group的信息资源，可以实现“signal下一个”功能，达到按序执行的调度目的 同步的指令都可以附带一个delay信息，因为静态的流水线，可以提前预估未来发生的同步点 同步指令不需要有状态信息，不需要读通用寄存器，不需要和其他有data hzd 大量的短的非stall的流水线 替代 少量的长的流水线，更多的非stall流水线有利于静态优化 架构 Fork Join fork指令附带的信息 配置的代码段的参数 需要被join的信息 代码段的index 当前launch的编号ID：编号ID能直接解析出硬件的IP和memory mapping fork：专用于指令load的指令 join指令的附带信息 launch的编号ID 向量寄存器VR L0.5 ，向量寄存器作为L0.5存在，所有的软件线程公用 向量寄存器VR L0，软件线程独立的寄存器，每个线程只有4个，直接设计在流水线上 张量寄存器TR L0.5 , 张量寄存器作为L0.5存在，所有的软件线程公用 DTE支持transpose pad slice deslice PU单元内部的线程之间支持快速同步 线程内支持微架构的流水线操作和控制 Reuse flags 用于软件主动控制4个register的cache 6个barrier硬件资源用于软件主动创建thread之间的异步依赖关系（cuda的关键调度流水线长度是6），而不需要浪费硬件面积 async_group 对历史的异步指令进行分组，以便灵活的进行同步 async copy bulk 一个指令拷贝多个数据 mbarrier：数据接受端支持主动维护数据传输状态，可以避免数据发起端频繁和接受端同步和fence来保证信号的前后顺序 所有的流水线的信息传递都因为有竞争，不能做无阻塞，常用的利用credit和valid/ready的阻塞方式，浪费太多的面积 本质上是硬件实现了某个程度的自动化的功能 如果能从软件角度就静态化，能节省大量的面积 从源头上（调度流水级）软件通过特殊指令控制发射的带宽来做到软流水无阻塞执行 产品架构 单个芯片die规模 4T int8算力 集成256MB DDR3 DRAM 1866 16bit 约3.7GB/s带宽 小规模的芯片，便于开发，降低风险和成本，便于仿真 支持无缝的拼接，多个芯片能容易组合成一个大的芯片，板级集成 降低流片成本，单个die非常小 降低封装成本，无需高级的芯片整合封装方案 die to die 通过低成本的低速链接，不需要高速的serdes方案 ISA ctrl Fork(code_index) launch Join() Scalar scalar 计算 Vector Tensor 微架构文档规范 适用于 design/ 下所有微架构 .md 文档。参考实现：[[LLC]] + LLC.pipeline.html 。 0. 定位 微架构文档是 人面向 RTL / [[logix]] 模型实现的合同 ：写完 .md 就能据此写 RTL，不再回头补 每个单元一份主 .md （如 L1.md 、 DMA.md 、 LLC.md ），描述功能与流水线 拍级流水线 强制 作为独立章节，不和功能描述混写 可选的交互式可视化 HTML 文件挂在同目录（ .pipeline.html ），不写入 git LFS、不依赖外部资源 1. 主文档骨架 主 .md 按以下顺序组织章节（编号用 #### N. 标题 ）： 编号 章节 内容 §1 顶层参数 容量 / 宽度 / 深度等常量的表 §2 顶层框图 ASCII 框图或 mermaid，标出对外端口与内部主要模块 §3 顶层端口 控制 / 数据 / 出对方向各一节，沿用本仓库已有端口约定 §4 地址映射 若是存储类单元；否则按需 §5 控制总线协议子集 opcode 表、异步操作约定、不做合法性检查的边界 §6 - §M 各功能模块详述 每个功能模块独立小节（阵列、命中/缺失、出口、维护等） §M+1 流水线 拍级时序契约，见 §2 §M+2 - Pin / Lock / Bypass 等可选语义、与其他单元关系、复位、调试、待规格化参数、DGMP 关联 数字章节用 H4 #### N. 、子节用 H5 ##### N.M 、子子节用 H6 ###### N.M.K 。三级到顶。 2. 流水线章节 每个微架构文档必须有"流水线"章节。结构如下，编号 N 取主文档下一个未用数字（LLC.md 取 §11）： 2.1 引言 一段话说明：本章把哪些功能章节（§a-§b）在全局时钟上展开到拍级 工艺前提（如 SRAM 端口形态 1RW vs 1R1W）必须显式 配套交互式视图： .pipeline.html 的位置和用途 渲染说明：Mermaid / WaveDrom / YAML 在哪些环境能渲染 2.2 §N.1 资源清单 按拍管理的资源表。一个三列表足够： 资源 端口形态 来源 §x.y 资源是 硬件实体 ——一个仲裁器、一组 SRAM 端口、一个 FF 寄存器组、一个表。不要把功能动作写进资源清单。 2.3 §N.2 路径一览 mermaid flowchart LR 显示所有路径的分叉、合流、异步链。每条路径用一个 stage 块（不展开内部）。 2.4 §N.3 ~ §N.X 每条路径 每条路径一个子节，内部固定四件套（用 H6 ###### 编号 §N.X.Y）： 拓扑 ：mermaid flowchart LR ，stage 节点链 + 分支条件 拍级展开 ：WaveDrom JSON 代码块，横轴 cycle，纵轴 stage / latch / 资源占用，跨信号依赖用 edge 字段 latch 字段表 ：每个 stage→stage latch 携带的字段（含字段名 + 简短解释） 资源占用表 ：拍 × 资源 的矩阵 跨路径关系（与主路径的冲突、回填唤醒等）单独成一个子子节。 2.5 §N.{倒数 2} 资源占用全景与仲裁规则 资源 × 占用方表 ：每个资源列出主要占用方 / 次要占用方 / 仲裁规则 stall 触发表 ：什么条件下哪条路径会 stall、解除条件 全 bank 吞吐分析 （如果适用） 2.6 §N.{倒数 1} 嵌入式 schema YAML 代码块作为 机读真相源 ，给工具消费（未来生成 logix / RTL 骨架、跑静态检查）。schema 由四部分构成： unit: clock: resources: : { kind: ..., ports/slots/width/banks: ... } paths: : stages: - { id: , in: [...], uses: [...], out: [...] } arbitration: - { resource: , contenders: [...], rule: "" } stall_conditions: - { src: , effect: "" } 写出哪几条路径作为样本即可，完整版可旁挂 .pipeline.yaml 文件。 3. 拍号命名约定 模式 拍号 例 主路径 S0 / S1 / S2 / S3 / S4 读命中 主路径分支 S3' / S4'（带撇号） 读缺失从 S2 分叉到 mshr 异步流水线 F0 / F1 / F2 回填（fill） 独立 FSM W0 / W1 / W2 / W3 维护扫描 walker 特殊短路径 B0 / B1 旁路 bypass 重发 S2'（带撇号，复用主路径名 + 撇号） 回填后 master 重发到 S2 不同路径的拍号空间不冲突，方便资源占用全景表交叉引用。 4. 引用与跨文档链接 同文档内 ： §N 、 §N.M 、 §N.M.K 跨文档 ： [[文件名]] §N （Obsidian wiki link 风格），文件名不带 .md 后缀 大段引用别人的协议时，直接说"完全沿用 [[L1]] §3.2 的 NoC slave 信号集"，不要重复贴信号表 5. 三种载体的分工 载体 表达 谁读 Mermaid flowchart 拓扑结构（节点 + 箭头） 人（GitHub / Obsidian / VS Code 原生渲染） WaveDrom JSON 拍级时序（波形 + 跨拍依赖 + 资源冲突） 人（ wavedrom.com/editor.html 、Obsidian + WaveDrom 插件） 嵌入式 YAML 机读 schema（resources / paths / stages / arbitration） 工具 三者 信息互补 ，不互相替代。同一份事实在多处出现时，YAML 是真相源——Mermaid / WaveDrom 是渲染产物（手写阶段允许双源，工具成熟后单源生成）。 不要使用：drawio（不好 diff、信息无法被工具消费）、PNG/SVG 截图（同前 + 改一处要重导）、Word/PPT（同前）。 6. 交互式 HTML（可选但推荐） 文件名约定： .pipeline.html ，与主 .md 平级。 单文件、零外部依赖（CSS / JS 内嵌），打开即用。数据来源理想是同目录的 .pipeline.yaml ；起步阶段允许把数据写在