MCHA: A Memory-Centric Hierarchical Architecture for Parallel-Sequential Computing

1

MCHA: A Memory-Centric Hierarchical Architecture for Parallel-Sequential Computing Daijing Shi , Hongxiao Zhao , Yihan Fu , Zhan Chen , Jiayi Li , Yihang Zhu , Anjunyi Fan , Yaoyu Tao , Yuchao Yang , Bonan Yan Peking University, Beijing 100871, China bonanyan@pku.edu.cn

arXiv:2608.04443v1 [cs.AR] 5 Aug 2026

Abstract Emerging workloads, such as Multi-Agent Reinforcement Learning (MARL), large-scale neuromorphic computing, and probabilistic graphical models, intrinsically exhibit parallel-sequential computing patterns. While these tasks demand massive parallelism to achieve high throughput, they are severely bottlenecked by irregular data access patterns centralized to main memory. Consequently, conventional architectures face fundamental limitations when executing these workloads, primarily manifesting as global buffer saturation and memory-bound bottlenecks. To address these challenges, we propose the Memory-Centric Hierarchical Architecture (MCHA), a reconfigurable hardware solution tailored for parallel-sequential execution. MCHA leverages a hierarchical communication strategy that facilitates distributed, inter-core data routing, thereby significantly reducing the bandwidth burden on the global memory. Complementing the hardware, MCHA introduces a novel parallel-sequential programming model that utilizes event-driven conditional triggers to effectively hide data transmission latency within the execution pipeline. We benchmark MCHA against a diverse suite of parallelsequential tasks, including MARL, motor variable control, and Markov random fields. Validated through our opensource, cycle-accurate simulator, MCHA demonstrates performance speedups ranging from 153.06× to 2456.96× over NVIDIA A100 GPUs on MARL workloads, while maintaining robust programming flexibility across other application domains. Furthermore, the architecture successfully reduces main memory access from 96% to 5.44%. When synthesized in a 28 nm process, the MCHA implementation occupies an area footprint of 2.92mm2 and consumes 115.36 mW of power at 200 MHz. MCHA is open-sourced at https://github.com/carabdis/MCHA. Keywords: Reconfigurable Architecture, Memory-Centric, Scalable Architecture, Bulk Synchronous Parallel, DataDriven Programming

This paper has been accepted for publication at the 59th IEEE/ACM International Symposium on Microarchitecture (MICRO). © IEEE. Personal use of this material is permitted. Permission from IEEE must be obtained for all other uses, in any current or future media, including reprinting/republishing this material for advertising or promotional purposes, creating new collective works, for resale or redistribution to servers or lists, or reuse of any copyrighted component of this work in other works.

(a) Throughput / GFLOPS 104 100

Extra Memory Access

10-1

Compute Bound

MPE Simple Spread Switch Riddle Starcraft STORM

Measured Throughput 100 102 104 Arithmetic Density (FLOPs/Byte) (c) Memory Utilization (%) GPU Utilization (%) (b) Time Cost (s) 40 Global Memory Ops Switch Starcraft STORM 72 MPE Riddle Local Memory Ops 1.5 30 Compute 96% 68 1.0 20 10

-2

10-2

64 0.5 95% 92% 94% 83% 89% 60 0.0 20-1280 20-1280 20-1280 20-1280 128 256 512 1024 2048 4096 Number of Agents Number of Agents Memory GPU Kernel (d) Intensity (%) Compute 21.1% Transmitted Data: 1410.6 KB 47.66× Redundant Data Essential Data: 29.6 KB 0 Memory 78.9% 10 10 0

10-1 0

25

50

70 100 125 Cumulative Active Time (ms)

150

175

200

Figure 1: (a) Roofline model of the NVIDIA RTX 3090 with PSC applications; (b) Simulated access operation breakdown of the GPU architecture; (c) Measured memory account and computing intensity; (d) Measured memory access and computing ratio during the active time.

1 Introduction Parallel-Sequential Computing (PSC) is a hybrid paradigm that combines concurrent execution of multiple tasks with stepwise sequential operations. This approach seeks to balance computational efficiency, system scalability, and adherence to logical dependencies within a problem. In contrast to purely parallel architectures (e.g., Tensor Processing Unit [1] and the Loihi family [2]) and strictly sequential processing, PSC decomposes complex problems into independent subproblems and interdependent phases. This allows it to leverage distributed hardware while preserving data consistency, control flow, and algorithmic correctness. Such a design mitigates key limitations of fully parallel execution, including synchronization overhead, data races, and bottlenecks described by Amdahl’s law [3]. Simultaneously, it avoids the inefficiency of purely sequential processing when applied to large-scale or computationally intensive workloads. By dynamically allocating parallelizable segments and enforcing sequential ordering where dependencies exist, the PSC execution model optimizes resource utilization and accelerates computation for data-intensive and algorithmically complex tasks. Consequently, it is widely adopted in high-performance computing, scientific simulation, and related domains.

2

PSC encompasses a broad spectrum of computing tasks, including Multi-Agent Reinforcement Learning (MARL) [4– 8], large-scale neuromorphic computing [9–12], and probabilistic graphical models [13–15]. Additionally, the Bulk Synchronous Parallel (BSP) paradigm, which underpins industrial data-intensive frameworks such as MapReduce [16], Pregel [17], Spark [18], and PolyGraph [19], is also integral to PSC. These tasks demand massive parallelism to manage frequent, fine-grained interactions among thousands of entities (e.g. agents in MARL), with strict synchronization required at each time-step boundary. In conventional compute-centric architectures (CPUs & GPUs), these repeated interactions induce irregular data access patterns that saturate centralized main memory. Consequently, as depicted in Figure 1(a), deploying such workloads on GPUs or many-core CPUs results in execution throttled by the global memory bandwidth (i.e. the “memory wall” [20–22]). The memory wall is exacerbated by the mismatch between irregular working sets and rigid hardware hierarchies. When on-chip storage fails to capture unpredictable access patterns, the system incurs frequent, high-latency DRAM fetches, leading to severe execution stalls [12, 23, 24]. Our analysis of a grid-world MARL task on an NVIDIA RTX 3090 (Figure 1(b)) reveals a significant inefficiency: data movement overhead accounts for 83% to 96% of total processing operations, depending on the simulation scale. Figures 1(c) and (d) further demonstrate that the majority of execution cycles and energy are consumed by redundant memory traffic rather than productive computation. Even recent reconfigurable hardware with data parallelism, such as Plasticine [25, 26] or Leviathan [27], remains tethered to data-centric or compute-centric paradigms. Plasticine employs a coarse-grained reconfigurable architecture (CGRA) with dataflow compilation [28–30] to decouple memory access from task processing. However, the lack of transparency between memory and processing units still necessitates redundant data fetches. Leviathan [27] utilizes in-cache near-data computing [31–34] but remains burdened by frequent cache coherence management overhead for PSC tasks. Other Processing-In-Memory (PIM) or BSP domain-specific accelerators (DSAs) are either highly application-specific [35–38] or bottlenecked by global buffer (GB) access for synchronization [39–43]. Consequently, existing works suffer from a ping-pong effect in the sequential part of PSC, where data is repeatedly transmitted between main memory and multi-level caches to satisfy rigid cache coherence protocols [44–47]. To address this challenge, we introduce MCHA, a reconfigurable and memory-centric multi-chip computing architecture. By devising a multi-tiered inter-core communication network and a data-driven programming model, MCHA effectively overcomes the data-transfer bottlenecks that limit the performance of GPUs, FPGAs, and CGRAs across a broad range of PSC tasks. Our key contributions are summarized as follows:

MARL

Neuro. Model

Parameter Gathering

Update Transformation f

State Write-Back

f

BSP f

Figure 2: The three-stage processing of PSC tasks.

⋄ MCHA Architecture: After quantifying the inefficiency of traditional architectures for PSC workloads, we propose a novel reconfigurable hardware architecture. It replaces centralized global buffers with a distributed memory fabric interconnected via a Network-on-Chip (NoC), substantially lowering cache coherence overhead. ⋄ Data-Driven Programming Model: We develop a formal programming model and an associated workflow to efficiently map PSC tasks onto the proposed reconfigurable hardware platform. ⋄ Comprehensive Architectural Analysis Framework: We release an open-source, cycle-accurate simulator (available at https://github.com/carabdis/MCHA) to evaluate the throughput and scalability of MCHA under various multi-chip configurations. Through extensive experiments, we derive new insights into hierarchical memory bandwidth limitations and propose effective mitigation strategies. For evaluation, we implement MCHA at the register-transfer level (RTL) using the TSMC 28nm HPC+ process design kit (PDK). Experiments across a suite of MARL benchmarks [48–52] show that a 4-chip MCHA system delivers a speedup of 153.06× to 2456.96× compared to an NVIDIA A100 GPU. We further validate the generality of MCHA using Motor Variable Control and BSP applications, where a 32chip configuration achieves speedups of 1.17× and 3.90× over state-of-the-art DSAs. These results underscore the efficiency of MCHA in handling irregular, memory-intensive dataflow patterns.

2 Target Workload: Parallel-Sequential Patterns PSC Formulation As depicted in Figure 2 and Algorithm 1, PSC follows parallel-sequential patterns, which can be characterized by an evolving state space composed of entityspecific data. To formalize the hardware requirements of these tasks, assume the global system state at any temporal index (“time” or “steps”) t as St . Each entity i within this state is represented by a multi-dimensional feature vecn tor xn i (t) ∈ R . The persistent structural constraints of the task, such as graph topology, environmental boundaries, or synaptic weights, are encapsulated in the parameter set p. The state transition logic is governed by an update function fi , which determines the feature vector of entity i at state t + ∆t based on features of its neighborhood Ni and itself: ) ( ⃗n ⃗n x⃗n i (t + ∆t) = fi xi (t), {xj (t)}j∈Ni , p , i = 1, 2, . . . , m

(1)

3

Algorithm 1 Parallel-Sequential Computing Tasks Input: neighborhood N , entity number m, persistent structural parameter p Output: Final entity state x⃗n (tend ), total state number tend . for each t do for each i ∈ [0, m] do ▷ Parameter Gathering n ⃗ Ni = FindNeighbor(xi (t)) Di = GBAccess(Ni ) end for parfor each i ∈ ( [0, m] do ) ▷ Update Transformation y⃗n (t) = f x⃗n i i i (t), Di , p Entities x⃗n (t),

end for for each i ∈ [0, m] do ▷ State Write-Back ⃗n x⃗n i (t + ∆t) = yi (t) ▷ Multicore Cache Coherence Sync. end for end forreturn tend where ∆t represents the step size between different states; Ni denotes the set of indices for entities influencing the update function of the ith entity; m is the total entity population count. From an architectural perspective, the primary overhead originates from the execution of f , which we decompose into three stages called Parameter Gathering, Update Transformation, and State Write-Back. Bottleneck Analysis We revisit Gustafson-Barsis’s law [53] through a memory-centric analysis. According to Gustafson’s formulation, the theoretical speedup S achieved via parallel computing is S = s + (1 − s)N = N − (N − 1)s,

(2)

where N denotes the number of processors and s represents the sequential fraction of the total execution time. Consider an ideal PSC task described with Equation 1, in which m independent entities are processed concurrently compute across N cores, each requiring compute latency T1-step . The sequential overhead arises from synchronization operations, specifically, gathering and writing back results. It inC +CBack , where curs a data-transfer latency Tdata-transfer = Gather BW CGather and CBack denote the data sizes for synchronization among the m cores, and BW is the global memory bandwidth. Hence, the sequential fraction s in Equation (2) can Tdata-transfer be approximated as s = compute . Tdata-transfer +T1-step

Substituting this into Equation (2) yields the memory-aware speedup approximation: S = N − (N − 1) =N−

1 T compute 1 + T 1-step data-transfer

N −1 compute T1-step ·BW

1+ C

(3)

.

Gather +CBack

When memory access becomes the dominant bottleneck, T

compute

i.e., T 1-step

data-transfer

≪ 1 due to relatively low bandwidth, the over-

all speedup S collapses to “1”, indicating no effective parallelism gain. Challenges. Based on the preceding analysis, memory access readily becomes the dominant performance factor due to two interconnected architectural challenges. Challenge 1: Global Buffer Saturation As shown in Figure 1(b), memory access dominates the execution profile, highlighting the significant communication overhead from CGather and CBack . The absence of direct peer-topeer (P2P) communication primitives between processing elements in conventional accelerators means that scaling m forces multiple cores to contend for global memory resources. This contention transforms the centralized GB into a serial bottleneck. Furthermore, in compute-centric architectures, the runtime dependency of Ni on entity states incurs substantial over-fetch overhead. The hardware must fetch entire cache lines to access non-contiguous entity data, resulting in severe bandwidth underutilization. Although the Update Transformation phase can be parallelized for SIMD/SIMT acceleration [54, 55], its low arithmetic intensity, often comprising only a few linear layers or element-wise nonlinearities, prevents effective utilization of computational units. Consequently, processing elements frequently stall waiting for data from saturated memory operations, leading to the second challenge. Challenge 2: The Memory-Bound Trap Figure 1(c–d) illustrates that regardless of workload complexity, GPU compute utilization remains low in PSC tasks. This indicates the bottleneck is an architectural latency issue, not a lack of raw computational throughput. The inefficiency originates from the rigid coupling of data indexing and processing in conventional architectures, which enforces sequential execution for irregular memory accesses. As a result, most of the procompute cessing time T1−step in Equation 3 is spent stalled, awaiting data fetches from the centralized global memory. Motivation Building on the challenges outlined above, we introduce two key solutions in MCHA: Multi-Tiered Peer-to-Peer Interconnect (Section 3): Our core idea is to parallelize the original sequential part Tdata−transf er with a novel programming model and reconfigurable hardware substrate. MCHA employs a multitiered peer-to-peer communication fabric which filters the data communication according to the transmission distance, thereby eliminating contention for higher-level transmission paths. By removing this centralized bottleneck, MCHA avoids memory-access serialization and supports near-linear scalability. Data-Driven Programming Model (Section 4–5): The proposed data-driven model replaces sequential control flow with fine-grained, data-driven memory operations. Communication between cores is realized through memory-mapped I/O (MMIO) FIFOs, decomposing inter-core data transfers

4

into discrete asynchronous events across the fabric. Consequently, these operations are hidden from the pipelines of each core, improving data independence of these cores and thus sustaining high computational throughput.

3 MCHA Architecture The MCHA architecture employs a computing layer of Memory-Centric Cores (MCCs) interconnected by a multitiered NoC. This design provides high-bandwidth data transmission while bypassing traditional multi-level cache structures to minimize multi-core cache coherence overhead. The NoC fabric is organized based on spatial access frequencies to maximize data locality. Compared with Cerebras [56], which implements data locality by holding everything on-chip, a three-tier implementation is illustrated in Figure 3(a), providing the basis for our architectural analysis and performance evaluation. The fabricated chip die photo of our implementation is embedded in Figure 3(e). 3.1 Computing Layer: Memory-Centric Cores The fundamental component of MCHA is the MCC, which augments on-chip memory with localized processing capabilities to minimize data movement overhead (Figure 3(b)). To ensure deterministic pipeline control, each MCC integrates a RISC-V core as an embedded orchestrator, supported by dual SRAM banks dedicated to instruction and data storage, respectively. Consequently, MCCs serve as the primary repositories for entity states xi , facilitating direct inter-core communication and bypassing the centralized GB and broader memory hierarchy. Inter-core communication is managed via MMIO FIFOs with explicit programs. This approach simplifies the programming abstraction while effectively masking data transmission latency within the processing pipeline. 3.2 Communication Layer-1: Processing Block In MCHA, MCCs are grouped into Processing Blocks (PBs) to handle resource-intensive update functions f . When a parallel-computed function (e.g., a full matrix multiplication) exceeds single-core capabilities, a PB facilitates collaborative execution through MCHA Communication Layer-1. This layer ensures data coherence between neighboring cores and manages the frequent data traffic resulting from multicore collaboration. Intra-block communication is supported by high-bandwidth FIFOs that enable near-instantaneous transfer of intermediate results. Thus, during multi-core collaboration, entity state xi remains resident in local memory, maximizing data locality. Additionally, P2P interconnection between PBs and an independent DRAM Fabric is embedded in this layer to facilitate multi-block collaboration and efficient DRAM access (Figure 3(c)).

3.3 Communication Layer-2 & 3: 2D-Mesh The Communication Layer-2 and 3 of MCHA consists of a 2D-mesh topology formed by P2P interconnection linking neighboring PBs (Figure 3(d–e)) and high-speed chipto-chip (C2C) interface across multi-chip boundaries (e.g., using UCIe or PCIe protocol [57, 58]). This distributed mesh serves as an active buffer, replacing multi-level caches to avoid redundant data movement caused by cache coherence across the memory hierarchy in multicore systems. The mesh enables P2P communication to resolve neighborhood Ni memory queries and other essential inter-core data access without the serialization overhead of centralized memory controllers. These links ensure system-wide data transparency, permitting the spatial distribution of complex functions f (Algorithm 1) across multiple PBs or chips when necessary. A dedicated DRAM Fabric handles off-chip traffic independently of the inter-core mesh. This decoupling isolates the processing fabric from high-latency off-chip memory access and avoids DRAM data from being transmitted across the whole multi-tiered network to improve DRAM access efficiency.

4 Data-Driven Programming Model Beyond conventional cache coherence management overhead, the performance of PSC workloads is limited by stalls arising from the tight coupling of data indexing and processing, which is a byproduct of the explicit data-fetch from memory to the central execution unit in the compute-centric model. We address this through a Data-Driven Programming Model that achieves latency hiding by breaking data access into a series of asynchronous events. By spreading these events across the distributed memory fabric, MCHA allows data transmission to be overlapped with local processing, significantly reducing the impact of memory-related stalls. Figure 4 shows the programming model for MCHA. To enable data-driven processing, we categorize architectural operations of functions into four primary triggers: 1. Memory Conditional Trigger: Initiates functions when local SRAM variables meet logical predicates, enabling state-dependent data processing. 2. Phase Conditional Trigger: Switches between triggering configurations at synchronization points to facilitate dynamic function deployment and smooth state transformations. 3. Input Trigger: Activates functions upon message arrival in hardware FIFOs, decoupling execution from NoC latency and ensuring data-ready activation. 4. Default Memory Scan: Performs an autonomous scan of local SRAM to trigger execution based on resident data in the absence of external FIFO inputs. This enables the

5

Comm. Layer-3 Inter-Chip

Computing Layer: MCC RISC-V Core Program Defined Address

Low

Define Instruction SRAM

Dataflow

Data Storage SRAM Output FIFO

(c)

(d)

Communication Layer-1 MCC

Edge I/O

MCC

Edge I/O

Local Decoder

Edge I/O

Edge DRAM MCC I/O I/O MCC

Input FIFO

DRAM Fabric

(e) Communication Layer-3

Communication Layer-2 DDR & C2C Interface PB

PB

PB

PB

PB

PB

PB

PB

PB

DDR & C2C Interface

DDR & C2C Interface

Comm. Layer-1 Intra-PB Comm. Layer-2 Inter-PB

(b)

DDR & C2C Interface

Comp Layer MCC

Data Access Frequency High

P2P Interconnect

(a)

DRAM Fabric

PB #1

PB #2

PB #3

PB #4

Fabricated Chip

Figure 3: (a) MCHA Communication “Pyramid”; Architecture of (b) Computing Layer; (c,d,e) Communication Layer-1, -2, & -3. Embedded figure in (e): MCHA die photo for silicon verification. GPU (b) 1 for i in [0, m]: Data Memory Instruction 2 N[i] = FindNeighbor(x[i]) Memory Data Phase 3 D[i] = GBAccess(N[i]) Log Log Jump 4 parfor i in [0, m]: 5 y[i] = f(x[i], D[i], p) PC > > 6 for i in [0, m]: Serialization Cause Condition Data List 7 x[i] = y[i] Severe Stall Trigger Ops MUX 1 parfor i in [0, CoreNum]: MCHA 2 if Core[i].MemReq(): (c) 3 # Mem Condition Trigger Instruction Ni 4 Core[i].MemFunc() Input Trigger Memory 5 elif Core[i].PhaseReq(): Type 6 # Phase Condition Trigger Ops 7 Core[i].PhaseChange() RISC-V Identify xj 8 elif Input is not None: Core Responder 9 # Input Trigger 10 Core[i].ProcFunc(Input) (d) 11 elif Core[i].DefaultScan(): Data Memory 12 # Default Memory Scan i-th Memory RISC-V Scan 13 Core[i].ScanMem() Core 14 # Do Default Function Default Instruction Ni 15 # FindNeighbour() Trigger Memory Conditional Trigger Enable High Core Utilization Rate (e) Conventional Pipeline: Core A FindNeighbor GB IN DRAM I/O GB OUT f Wait Write Back Wait Sync

(a)

Core B FindNeighbor Wait GB IN DRAM I/O W.GB OUT f Wait Write Back Sync Initial Excecution Write Back MCHA Pipeline Core A DRAM I/O FindNeighbor Reception Retrieval f Sync NoC Inter-Core Transmission MCHA Saving Core B DRAM I/O FindNeighbor Reception Retrieval f Sync

Figure 4: (a) Compute-centric and data-driven programs; (b) Dataflow for Conditional Trigger and Phase Conditional Trigger; (c) Dataflow for Input Trigger; (d) Dataflow for Default Memory Scan; (e) Pipeline Comparison of Conventional architecture and MCHA.

MCC to initiate the computing loop without external signaling. The corresponding dataflows of these triggers within MCCs are presented in Figure 4(b–d). A more detailed explanation and realization of our design is provided in Section 6.2.2. Furthermore, within this framework, to decouple data indexing from the processing pipeline, inter-core communication is decomposed into a three-stage asynchronous trigger sequence: ⋄ Initial Execution, which triggers the generation and transmission of a data index, demonstrated by any trigger; ⋄ Index Reception, which initiates a local memory search and subsequent message return, demonstrated by the Input Trigger; ⋄ Data Retrieval, where the arrival of the indexed data resumes the suspended computation, demonstrated by the Input Trigger. By partitioning communication into these discrete asynchronous triggers, MCHA effectively masks data transmission latency by overlapping it with the active processing pipeline.

Figure 4(e) presents the conceptual dataflow of programs in Figure 4(a). In a conventional pipeline, the centralized GB obstructs the execution flow, leading to severe memorybound performance degradation. In contrast, by utilizing the Default Memory Scan trigger, MCCs initiate parallel offchip data loading to maximize DRAM bandwidth utilization through hardware-level concurrency. Computation of the updating function f will be triggered by Memory Conditional Trigger. Detecting the completion of the asynchronous data indexing process and stage transmission of the pipeline is handled through the Phase Conditional Trigger. The implementation of the trigger-based data-driven programming model isolates the computing pipeline of each MCC, reducing Tdata−transf er in PSC workloads thanks to the elimination of stalls.

5 Communication Protocols The design target of the programming model in Section 4 is to enhance data independence of MCCs for isolation of the local pipeline from being obstructed by others. To realize asynchronous communication for this purpose, we formalize the communication protocols of MCHA with the help of MMIO FIFOs embedded in the MCCs instead of relying on a centralized memory controller like NVIDIA Hopper GPU [59]. The direct mapping of the NoC interface into the local address space of each core eliminates the control-plane distinction between computation and communication, isolating the data access in each MCC from that of others. This design hides communication latency within the execution pipeline, which is particularly advantageous for PSC patterns. Furthermore, the multi-tiered communication hierarchy mitigates bandwidth contention for long-distance transfers (e.g., inter-chip and off-chip data access) by isolating local traffic from global data movement. 5.1 Core Message Management In MCHA, each MCC functions as a self-contained execution environment where operations are triggered asynchronously by incoming messages or local memory states, consistent with the principles detailed in Section 4. Workloads are deployed to an MCC’s local independent instruction memory with optimized execution flow from compilation. As illustrated in Figure 5(a), communication FIFOs bypass complex network drivers and are instead mapped directly to the local address space. Consequently, transmitting data to a neigh-

6

boring MCC is compiled to a standard memory-store (store) operation. When the RISC-V core writes to a designated output address, the data is enqueued into the output FIFO. On this other side of receiving, data reception is handled via memory-load (load) operations, where the input FIFO functions as the core’s event queue for incoming messages awaiting processing. The implementation of the MMIO FIFOs masks data transmission latency within the MCC execution pipeline, isolating the MCC local processing from the global NoC data transmission. 5.2 Transmission Message Definition To shield the local MCC pipeline from unexpected stalls induced by the multi-tiered NoC latency, we isolate MCC execution from the global dataflow through a latency-hiding communication framework. This framework stratifies data traffic based on transmission frequency and bandwidth availability, orchestrated by a streamlined 32-bit packet format. By supporting two hardware-optimized message types that match the spatial locality patterns in MCHA (Figure 5(b)), the design effectively minimizes routing overhead while ensuring high flexibility for intertwined PSC execution. Type 1 serves high-throughput, one-hop communication, corresponding to Communication Layer-1. It constitutes the dominant traffic during the Gather phase of Algorithm 1, where MCCs exchange entity states with adjacent PBs or neighbors in the 2D-mesh. Each Type 1 message is a single self-contained 32-bit packet. The first 3 bits serve as the Type ID and specify the hop direction (North, South, East, West, or Local). The subsequent 5 bits in the same byte define the target (Cores or DRAM), leaving a 3-byte payload for user-defined data. By restricting communication to a single hop, the hardware avoids complex routing headers and multi-cycle state tracking, enabling near-instantaneous data-residency updates. Type 2 supports lower-frequency global communication using variable-length bursts, corresponding to Communication Layer-2 and 3. Each burst comprises a single header followed by a specified number of data packets. As with Type 1, the initial 3 bits identify the packet type. A dedicated length field (bits 27–28 in Figure 5(b)) specifies the total number of subsequent packets in the burst. The remainder of the header carries the routing information required to traverse the 2D-mesh and reach distant chip-level coordinates. To maximize effective bandwidth, only this initial packet contains routing metadata; all following packets are devoted almost entirely to user-defined content, thereby enabling efficient transmission of weights, environmental parameters, or extended entity histories. The entire transmission process is software-configurable via the RISC-V ISA. This dual-tier protocol ensures that local neighborhood exchanges are resolved with minimal latency, while Type 2 messaging provides the necessary transparency for global MCC interactions across a large-scale multi-chip fabric.

5.3 Intra-Block Transmission Multi-core collaboration within the PB triggers intense P2P data sharing, characterized by high-frequency, fine-grained message inquiries. To support this behavior, we deploy the optimized Type 1 transmission format in Section 5.2. To realize stall-free coordination, PBs incorporate a highbandwidth bidirectional interconnect linking the MCCs, as shown in Figure 5(c), enabling intra-block communication without traversing the global 2D-mesh network. This localized routing effectively unburdens the global mesh and prevents cascaded core stalls. In place of the powerintensive crossbars in general-purpose NoCs, the design employs a FIFO-based queuing and arbitration scheme optimized for PSC workloads. At the output stage of each MCC or PB boundary, a hardware message decoder (Figure 5(c)) parses the Type ID and routing metadata of outgoing packets (Section 5.2). This decoder translates high-level message definitions into physical control signals that drive internal multiplexers: for Type 1 messages, it identifies the single-hop destination port; for Type 2 messages, it locks the path for the duration of a packet burst, ensuring contiguous transmission without interleaving jitter. Resource contention arises when multiple MCCs attempt to transmit to the same destination concurrently. Following the first- come-first-served principle inherent to FIFO-based designs [60], each input interface is equipped with a dedicated counter. If an incoming message request is not accepted by the target FIFO, its counter increments whenever the target FIFO attempts to fetch messages. During each arbitration cycle, the selector evaluates all active requests and grants priority to the message with the largest counter value—indicating the longest waiting time—pushing it into the input FIFO. Once transmission begins, the corresponding counter is reset. This policy serves as a critical antistarvation mechanism, preventing indefinite delays in state updates. An ablation study comparing the proposed arbitration policy with conventional Round-Robin and random-selection schemes is presented in Figure 5(d). The simulation configures 8 input-output FIFO pairs, each with 1,000 packets destined for randomly chosen targets. For the evaluated PSC tasks like MARL, long latency tails directly impact sequential processing speed and MCC stalling time; thus, we compare the latency at the 99th percentile (P99 latency) across the three policies. With the simulated packet sending rate (x-axis) increase, the P99 latency (y-axis) increases to saturation. This demonstrates that the first-come-first-served policy reduces the tail latency by 6.26% to 41.8% depending on packet generation rate, confirming the effectiveness of the selected design. 5.4 Global Neighboring Connection Following the same design philosophy of Section 5.2 and 5.3, we decouple timing constraints across the MCHA

7

Read Write RISC-V Core

Data Storage SRAM Input FIFO Output FIFO

Data Address

Output FIFO

FIFO Address

(b) Type 1: 29 28 27 2524 0 31 TYPE ID 1 Target Payload 000~101 3b 24b Type 2: Header Packet 2928 2726 1312 0 31 TYPE ID Length Dest. Payload 110~111 2b Addr. 14b 13b Data Payload Packets 32b

(c)

Pop Signal (d)

Payload Payload 32b 32b

7 6 5 4 3 2 1 0

Message Decoder Input FIFO

MUX

(a)

Activate Target Counter

6 System Deployment and Mapping Strategy

Waiting Counter

MAX INDEX P99 Latency (103 Cycles) Ours Round Robin Random

0.2 0.3 0.4 0.5 0.6 0.7 0.8 Packet Sending Rate (# / Cycle)

Figure 5: (a) MCC Message management method; (b) Message Packet Definition; (c) PB transmission strategy; (d) Ablation study comparing the P99 latency with other routing methods. (a) Sender HST_S HSR_S

HST_S Request HST_R Acknowledge

(b) Receiver HST_R HSR_R

Phase 1

Phase 2

Phase 3

Phase 4

HST_S EndTime

Start HST_R

Time

Figure 6: (a) Port definition for 1 internal connector pair (b) Timing diagram of the connectors in (a)

with an asynchronous network between neighboring PBs to achieve independent, self-governing execution. This allows each PB to operate within its own independent clock domain, reducing global clock distribution complexity and mitigating peak power surges. As shown in Fig. 6, inter-block data transfers (e.g., Block A to Block B) are governed by a robust 4-phase asynchronous handshake protocol: Phase 1. Request: The protocol begins with Block A driving its handshake sending port (HST_S) high to signal a request. Phase 2. Acknowledgement: Upon detecting this request, Block B responds by driving its acknowledgement port (HST_R /HSR_S) high. Phase 3. Clear Request: Block A then clears its request signal and reset HST_S to “0”. Phase 4. Clear Acknowledgement: After which Block B clears its acknowledgement, completing the cycle and freeing the channel for the next transfer. This lightweight handshake acts as a hardware-level flow control, ensuring data residency and integrity across the mesh even when neighboring blocks operate at different frequencies or states. When scaling to multi-chip systems, MCHA extends its 2D-mesh communication protocol across chip boundaries using high-bandwidth interconnects such as PCIe [57] or UCIe [58]. These system-level links maintain information transparency, enabling a distant core on a remote chip to be accessed via the same asynchronous handshake strategy. Consequently, MCHA provides a modular and scalable fabric capable of supporting PSC across multiple chips.

This section elaborates how to maximize the communication efficiency of the multi-tiered NoC and the parallelism of MCCs through the formulation of algorithms along with their deployment strategies on the MCHA platform. Figure 7(a– b) illustrates the programming and deployment workflow of MCHA with a representative program example, matrix processing A · (prod(B)) − C, where prod() computes the product of all numerical elements in a matrix. 6.1 Collective Usage of MCC Buffers Since MCHA enables high independence to the MCCs, the management of the on-chip memory and workload varies from the conventional systems. Instead of allocating specific tasks to different cores, MCHA maps computational entities to cores based on the memory footprint of their state data (xn i in Algorithm 1). Only cores that control this data can modify it; others can either read the data or send write requests to these cores. Compared with threads in a GPU that are assigned to specific operations each time (Figure 7(c–d)), cores in MCHA perform different operations based on memory types to parallelize the irregular data access, namely Tdata−transf er in PSC workloads. As presented in Figure 7(a), Core 0 and 2 are responsible for loading the data of matrices A and C, acting as the entity buffer. When the state data size surpasses the capacity of on-chip SRAM, only the necessary indices for accessing off-chip DRAM are maintained in SRAM, as Core 0 indexing data of vector A in Phase 1. By organizing these index structures in alignment with the underlying hardware topology (e.g., arrange the elements of A and C alternately in the DRAM), we convert the serialized and irregular memory access characteristics of traditional paradigms into predictable, parallelized local transactions. Consequently, the most frequently accessed variables remain physically colocated with the corresponding processing logic, thereby eliminating the “memory wall” at the architectural level. 6.1.1 Locality-Aware Entity Clustering To filter local data transfers through the multi-tiered NoC, thereby minimizing communication overhead in the 2Dmesh, the mapping process prioritizes computational locality by assigning neighboring entities to the same MCC or adjacent cores within a PB based on their interaction frequency. For PSC workloads, we can use clustering [61] or graphpartitioning [62] algorithms to group entities with overlapping neighbor sets (Ni in Algorithm 1), thus maximizing data reuse within local SRAM banks. In the provided example, this process is represented as placing the three cores within the same PB. As the neighbor sets Ni evolve during task execution, e.g., due to agent movement necessitating intercore communication, MCHA supports periodic entity migration, with frequency and strategy tailored to the task. In complex environments such as Multi-Agent Particle Environ-

8

ments (MPE) [63, 64] or StarCraft maps [49], the MARL simulation space is initially partitioned into blocks (Figure 7(e)). Throughout sequential steps in PSC patterns, these block boundaries are dynamically adjusted based on local agent density, rebalancing computational load and ensuring consistent architecture-wide performance. 6.1.2 Systolic Mapping for Complicated Update Functions When processing computationally intensive update functions f , such as deep neural network inference, Tdata−transf er is embedded in the data sharing between collaborating MCCs. Workloads requiring frequent communication, for example, the lookup-table of non-linear functions and corresponding input generatos, are implemented within the same blocks. Others are spread across the system according to the dataflow structure, formulating a systolic array. For example, in Figure 7(a–b), Core 1 acts as the buffer of intermediate results for computing A · (prod(B)). Applying an input-stationary systolic mapping [65], these cores first store the input data (A in this example) and then, triggered by the completion of input transmission, network weights (B in this example) are streamed through these MCCs while previously prepared input data remain stationary in local SRAM. By sequencing weights according to input channels, the architecture can skip null-value operations and support sparsity-aware computation. In detail, MCHA MCCs demonstrate this process through 2 continuous phases, one for temporal buffering (Phase 1 of Core 1), the other for data indexing and computing (Phase 2 of Core 1). The completion of the input transmission acts as the Phase Change Trigger of Phase 1, changing the core to the computing phase, where data indexing of B is default. The core then performs operations according to pre-defined instructions (prod in this example, multiply-and-accumulate for neural networks), triggered by the arrival of indexed data. Figure 7(f) illustrates the systolic mapping of a 2-layer Neural Network (NN) in MCHA. The blue cores (labeled as NN-L1) store the data generated in the Parameter Gather process, namely the input data of the first NN layer, acting as intermediate buffers (a similar role as Core 1 in Figure 7(b)). This spatial unrolling prevents computationally dense transformation phases from becoming bottlenecks, thereby maintaining the efficiency of the data-driven computing flow. 6.2 Data-Driven Programming Model Implementation by RISC-V ISA The reconfigurability of MCHA is demonstrated through the programmable triggering logic based on standard RISC-V instruction set (RV32I as an example). The ISA execution of each MCC is governed by the state of the FIFOs and local SRAM conditions through 4 different triggers in the programming model at each stage.

(a) MCHA Core 0 Core 1 Core 2 (b) A Core 0 In = ReadFIFO() A1 A2 1-3 P1 P2 E1-3 C1 C2 F1-3 Phase 1: MCHA if MemSize()==3: A3 P3 C3 goto Phase 2 # Comp B1-3 elif In is None: (c) cudaMemcpy(DRAM, B) WriteFIFO(A_addr, DRAM) vecProd<<<1,3>>>(B, D) else: WriteMem(In) # Memory GPU cudaMemcpy(DRAM, [A, C]) vecScalProd<<<1,3>>>(D, A, E) Phase 2: vecMinus<<<1,3>>>(E, C, F) if HandShake(): cudaMemcpy(F, DRAM) WriteFIFO(A, Core1) (d) goto Phase 3 # Halt Cache Line GPU elif In is not None: RecvHandShake(In) B1 B B B1 A C B1 A C Phase 1: Core 1 Thread 0 B2 1 2 B2 1 1 B2 1 1 if MemSize()==3: B3 D B3 B3 D E1 B3 F1 E3 goto Phase 2 # Comp B1 B2 A1 A2 C2 A1 A2 C2 elif In is not None: A A Thread 1 WriteMem(In) # Buffer D B3 A2 D E2 A2 F2 E3 3 3 Phase 2: if Count==3: B1 B2 C1 A3 C3 C1 A3 C3 C2 C2 Thread 2 WriteFIFO(Shake, Core0) D B D E F3 E3 C C goto Phase 1 # Cycle 3 3 3 3 elif In is not None: (f) (e) WriteMem(P * In, Core2) Inital MCC Count += 1 Partition else: PB WriteFIFO(B_addr, DRAM) ... t Steps Systolic Core 2 Phase 2: Flow ... Border if In is not None: Gather & Update Migration WriteFIFO(In - C, DRAM) Count += 1 NN-L1 NN-L2

Figure 7: (a) A program for MCHA; (b) Conceptual dataflow of (a); (c) Program using CUDA; (d) Conceptual dataflow of (c); (e) Inprocess update enabled by 2D-mesh; (f) On-chip systolic function demonstration;

6.2.1 MMIO-Based FIFO Interaction By leveraging the MMIO FIFOs described in Section 5.1, sending and receiving data in MCCs reduces to standard RISC-V LW (Load Word) and SW (Store Word) instructions. This decouples communication from computation, allowing each MCC’s processing pipeline to remain uninterrupted and transmission latency to be effectively hidden. (a) When a compiled SW instruction targets the output FIFO address (WriteFIFO), the hardware pushes the 32-bit payload into the transmission queue. To prevent data loss, if the FIFO is full, the hardware interface stalls the pipeline or maintains the value in the register file until a slot becomes available, similar to the behavior of memory handshake in conventional SW operations. Crucially, the ISAs ensure atomic transmission; once a multi-packet message (Type 2) begins, the hardware locks the port to ensure the packets are transmitted contiguously without interleaving jitter. (b) An LW instruction targeting the input FIFO (ReadFIFO) pops the frontmost 32-bit packet. To enable high-speed triggering, the hardware utilizes a zero-signature mechanism: if the FIFO is empty, the LW returns a 32-bit “0”, which triggers the Default Memory Scan if it exists. Conversely, if a valid packet is present, the protocol guarantees that at least one bit is non-zero (bit 28 for Type 1 or bit 31 for Type 2, as shown in Figure 5(b)).

6.2.2 Control Flow Instruction Mapping As presented in Figure 7(a), each time MCC is idle, it will read the address corresponding to the input FIFO if Input Trigger is not empty and check all the other triggers. Thus, the basic structure of the program at each phase is a loop, which repeats the data detection process until jumping into other phases. The 4 triggering conditions are accomplished by RISC-V branching instructions, as demonstrated by the

9

Table 1: Design Parameters for Large-Scale Simulation

Part MCC

PB

MCHA

Parameter Memory Size (KB) Pipeline Parallelism FIFO Buffer Processing Throughput (GFLOPS/Core) Edge Connector DRAM Connector Number of Cores Connector Bandwidth (GB/s) External Connector DRAM Access Parallelism Number of Blocks External Connector Bandwidth (GB/s) DRAM Access Bandwidth (GB/s)

Value 1 1 8 1.6 4 1 4 1024 4 16 16 64 409.6

if statement in Figure 7(a), allowing the core to react to the memory fabric in real-time. (a) Memory & Phase Conditional Triggers: These two kinds of triggers rely on reconfigurable data statistics stored at pre-programmed addresses within the local SRAM. By periodically reading these status variables, the RISC-V branch logic determines if the core should handle special operations or switch functional phases. (b) Input Trigger: By executing an LW from the input FIFO followed by a conditional branch, such as BNE, the core can immediately detect incoming data. If the result is non-zero, the Program Counter (PC) jumps to the corresponding processing function. (c) Default Memory Scan: If an LW from the input FIFO returns zero (empty FIFO), the core can default to a memory scan. This is managed via a scan pointer maintained in the register file, which incrementally traverses the local SRAM to check for a response. Ultimately, every functional trigger in the MCHA model is represented as a PC jump to a specialized firmware routine. By mapping the programming model components into existing RISC-V instructions, MCHA introduces minimal efforts for ISA construction. 6.3 Dataflow-Defined Workload Assignment To maximize MCHA throughput, the assignment of the workloads should involve identifying essential reactive elements and mathematically modeling throughput requirements to determine the optimal pipeline depth for the target computing tasks.

6.3.1 Functional Requirements for Data-Driven Programs Similar to object-oriented programming, the MCHA requires programmers to break the target algorithm into multiple statically defined phases. The 4 types of triggers of a specific phase act as the member functions of an object. The programmer should define Phase Conditional Trigger, which

serves as the termination condition for a computational loop, analogous to loop-exit logic in conventional paradigms. Furthermore, each phase must include at least one definition of an Input Trigger or a Default Memory Scan to dictate the specific operations dominating that phase (WriteFIFO and WriteMem in Figure 7(a) for Core 0). If neither is defined, the MCC becomes a “zombie” core, because it lacks both external stimulus responses and autonomous self-activation. The transmission from the conventional programs to MCHA ones can be accomplished with the help of Large Language Models (LLMs). We provide an example that transforms a neural network to MCHA programs in C++ with the help of Deepseek-R1 [66] in thinking mode in our github (https: //github.com/carabdis/MCHA/LLM.md). The program must explicitly include handshake signal generation to manage the flow of data across the pipeline, ensuring that residency is maintained without causing meshwide stalls. As presented in Figure 7(a), Phase 2 of Core 0 and 1, the explicit handshake acts as the data barrier for the serialized part in PSC workloads. 6.3.2 Analytical Deployment and Throughput Matching After decoupling the PSC workloads into serialized phases, the deployment of the algorithm can be defined as a flow matching problem. We provide a 3-step analytical process to deploy a program onto MCHA fabric to balance the workload across the system. Step-1 Throughput Extraction: The processing characteristics of MCC can be described through two critical metrics, consumption rate Ripop and production rate Ripush . Rpop is defined as the average cycles required for a core to pop and parse a 32-bit input packet from the input FIFO. Rpush is defined as the average cycles required for the core to compute the assigned workload and push the resulting packet to the output FIFO. Thus, the goal is to match the ingress and egress rates of each stage in the pipeline. For example, neglecting the DRAM access latency, Core 0 in Figure 7(a) generates 1 output packet per cycle in Phase 2, namely R0push = 1. Core 1 consumes the 1 input packet from Core 0 every 3 cycles to compute the intermediate results of A · (prod(B)), thus R1pop = 1/3. Step-2 Ideal Pipeline Derivation: Under the assumption of infinite hardware resources, the ideal parallelized core count for each phase (Niideal ) required to achieve a non-blocking pipeline without hardware waste is derived from the ratio of these rates. To maintain a steady state where FIFOs neither overflow nor starve, the flow balance between serialized phases is defined as: ideal Niideal ∗ Ripush = Ni−1 ∗ Ripop

(4)

In our example, since R0push = 1 and R1pop = 1/3, the ideal ratio N0ideal : N1ideal is 1 : 3. This ratio determines the quantitative relationship between MCCs operating in neighboring phases, minimizing unnecessary data movement of interme-

10

Table 2: The power and area breakdown of MCHA

Part

MCC

PB

MCHA

Component SRAM FIFO RISC-V Core Control Total (MCC) MCC Control Total (PB) PB Ext. Connector Int. Connector Total

Area Area (mm2 ) (%) 0.0303 82.78 0.0006 1.64 0.0044 12.02 0.0013 3.55 0.0366 0.1464 94.57 0.0084 5.42 0.1548 2.4768 84.80

Power Power (mW) (%) 0.4968 41.40 0.0706 5.88 0.4913 40.94 0.1412 11.78 1.2000 4.8000 76.78 1.4514 23.21 6.2514 100.0224 86.70

0.0560 1.92

1.9352 1.68

0.3878 13.28

13.401811.62

2.9206

115.3594 -

-

diate results. Step-3 Realistic Constraint Scaling: Given a fixed physical scale, the deployment logic determines the maximum nummax ber of concurrently active phases Mpara . Due to memory structural hazards—where buffering and computation cannot occupy the same SRAM bank simultaneously—parallel phases must be interleaved with idle or buffering phases or employ twice the memory capacity via ping-pong buffering. max The algorithm computes Mpara ≥ 1 based on the minimum buffer cores required per phase. In the former example, if phases require {1, 2, 1} memory cores for Mpara = 2, but max only 3 cores are available, then Mpara = 1. Starting from max Mpara , the system iteratively reduces the number of active phases until memory capacity constraints are satisfied. This scaling ensures robustness across implementations, from single-chip edge devices to large-scale multi-chip fabrics.

7 Evaluation 7.1 Experimental Setup To evaluate the performance and efficiency of large-scale MCHA architecture, we developed a cycle-accurate simulator. This simulator is directly extracted from RTL Verilog HDL implementation, including MCCs, PBs, on-chip 2D-mesh design, and the chip-wise connectors. The chip-wise connectors demonstrated in the Verilog code are SPI connectors as a research prototype, which are replaced by parameters of commercial PCIe 4.0 in the simulator. Other hardware parameters are modeled based on logic synthesis results using TSMC 28nm PDK at 200MHz. Detailed parameters of one MCHA chip using in the simulator are in Table 1. The corresponding area and power overhead of our implementation are given in Table 2. To verify MCHA, we compare the performance of the MCHA system composed of 4 (MCHA-4) and 32 (MCHA-32) chips to various kinds of DSAs and generalized GPUs. NVIDIA

A100 GPU is used as the baseline for massive-scale parallel workloads with flexibility. PEARL [67], ActiveN [68], MC2 A [69], Dalorex [70], and PolyGraph [19] are chosen as the specialized accelerators for different benchmarks. The evaluation covers three distinct domains, highlighting MCHA’s ability to handle various types of programs: (1) MARL: a comprehensive suite including MPE Simple Spread [63], StarCraft II (SMAX) [49], Switch Riddle [50], and STORM [51, 52]. Performance is compared against JAX-MARL [49] running on A100 or RTX 3090 and PEARL [67]. For each environment, we measure the average processing speed of the baselines computing 1,000 timesteps. The number of agents existing in these environments is fixed as 1,000, to manifest the efficiency in large-scale MARL computation. To ensure a fair comparison, the MCHA implementation utilizes the same hyperparameters as the JAX-MARL demonstrations, including neighborhood size Ni , action space, and update functionsfi . (2) Large-Scale Neuromorphic Computing: The benchmark is a 2% MVC (Motor Variable Control) nervous system model [71], which follows the setup established by ActiveN [68], testing the architecture’s capability to handle sparse, high-frequency spikes and deterministic timing. (3) Typical BSP for Graph Processing: We evaluate MCHA ’s performance on regular data access patterns using PageRank [72], BFS [73], and Markov Random Field (MRF) image segmentation [74], compared with existing DSAs (MC2 A [69], PolyGraph [19], and Dalorex [70]) and GPUs [19, 69, 75, 76]. The evaluation uses the orkut graph [77] when comparing PageRank and BFS against PolyGraph [19], and the RMAT-26 graph [78] when comparing against Dalorex [70]. MRF image segmentation input data follows [74]. For the above workloads, MCHA matches the algorithmic setup of every corresponding baseline [19, 49, 68–70] in numerical precision, task size, step definition, and convergence condition. 7.2 End-to-End Computing Throughput 7.2.1 Multi-Agent Reinforcement Learning (MARL) Figure 8(a) compares the absolute execution time (y-axis) of different hardware (x-axis), including MCHA (M1, M2), GPUs (G1–G2), and FPGA (F1). The presented throughput time is measured end-to-end including data off-loading and data processing. MCHA-4 delivers 153.06× to 2456.96× speedups versus GPU-accelerated frameworks. MCHA-32 further achieves 1.22∼4.25× speedup over MCHA-4. This performance leap stems from the elimination of redundant global memory copies due to frequent cache coherence and overfetch in GPUs. Furthermore, the high data locality

11

7.2.3 Typical BSP for Graph Processing Figure 8(a), (b)&(d) presents the results. MCHA presents comparable results versus various DSAs for these applications. Compared with general-purpose GPUs for BSP tasks, MCHA-32 achieves 86.69×, 3.93×, and 4.56× speedup, respectively, while MCHA-4 also presents comparable performance. We further compare MCHA with Dalorex[70], whose 256 cores match the compute-unit count of MCHA-4 (“M2” in Figure8). M1L and M2L denote M1 (MCHA-32) and M2 (MCHA-4) with on-chip memory scaled to 4.2,MB/core to match Dalorex’s configuration, respectively. In their original setup (M1 and M2), MCHA shows little advantage over Dalorex; however, M1L achieves a 5.66∼7.28× throughput improvement, and M2L reaches throughput comparable to Dalorex. This gap is explained by DRAM access overhead: Dalorex’s model excludes DRAM latency, whereas MCHA’s includes it. On the RMAT-26 [78] dataset, DRAM accesses consume 81.42% of M1’s execution time under the original 1KB/core configuration — an overhead largely eliminated by the enlarged on-chip memory in M1L. The lower speedup of BSP implementations relative to MARL and MVC stems from differences in arithmetic intensity across the workloads. The most informative metric is the bandwidth utilization rate. In the BSP-MRF task, the performance ratio between MCHA-4 and the GPU (0.61) is close to the ratio of their respective DRAM bandwidths (0.67). This correlation proves that MCHA-4 maintains

(a) 105 104 103 102 101

DRAM Access (MB)

146.152 67.318 0.000

0.092

64.161 30.685 0.000

0.035

461.526 141.667 0.854 107.375 40.645 4.378

0.139 0.077

Steps / s MARL MPE Simple Switch Riddle Starcraft Spread

InterBlock (MB)

STORM

BSP PageRank

GTEPS 2 10 BFS 101 100

BFS

101 100

(c) Speed Up (d) Speed Up 101 4.56 102 113.6 96.6 1.17 1 15.2 12.2 101 100 0.61 2.8 0.20 1 0.78 100 -1 0.13 0.06 0.29 10

(e) Speed Up 104 M1 MO1 102 100

MO2

MO3

MO4

M 1 M M 2 C2 A G 3 AI PG A M A XC

(b) GTEPS 102 PageRank

M 1 M 2 AN N Lo E ih i G 1 G 2 G 4

Figure 8(c) shows the comparison results benchmarking MCHA (M1, M2), neuromorphic accelerators (AN, NE, Loihi) [2, 68, 79] and GPUs (G1, G2, G4). MCHA-4 shows 5.36× improvement over the neuromorphic-specific accelerators [2, 68], completing a single task in 0.7842s, while the MCHA-32 presents 39.996× speedup. When compared with ActiveN, MCHA-32 performs 1.175× speedup with 9.36% power consumption, presenting the high efficiency of MCHA originating from the data-driven nature. Furthermore, the near-linear performance improvement between MCHA-4 and MCHA-32 originates from the high spatial locality achieved by the communication strategy (more details are investigated in Section 7.3.3).

InterChip (MB)

IntraBlock (MB)

M 1 M 2 F1 G 1 G 2 M 1 M 2 F1 G 1 G 2 M 1 M 2 F1 G 1 G 2 M 1 M 2 F1 G 1 G 2 M 1 M 2 P1 G 1 G 2 M 1 M 2 P1 G 1 G 2

7.2.2 Large-Scale Neuromorphic Computing

Core Memory Benchmark Ops (MB) MPE Simple 211.900 Spread Switch 102.355 Riddle Starcraft 1044.362 STORM 158.167

M1L M1 M2L M2 DA T M1L M1 M2L M2 DA T

Compared to PEARL (F1), MCHA-4 achieves a speedup of 2.28× to 4.78×. This performance gain is primarily attributed to the data-driven adaptivity inherent to the MCHA architecture, which masks data transmission latency across the multi-tiered NoC. Furthermore, MCHA-4 eliminates the significant communication overhead typically associated with data transfers between external GPUs and the FPGA fabric, streamlining the execution of PSC workloads.

Table 3: The memory footprint of MCHA-4

(f) Contribution Percentage (%) G1 G3 80 35.5 29.5 36.3 31.9

SM PR MVC MRF SM PR MVC MRF SM PR MVC MRF SM PR MVC MRF Ba SM se PR lin MVC e MRF SM PR MVC MRF

achieved in MCHA in turn removes most of the irregular data access to DRAM, thus improving the overall performance.

60 40 20

25.6 30.2

18.9 39.4

20.2 36.5

26.0 35.1

SMAX PR MVC MRF MO1 MO2 MO3 MO4 M1: MCHA-32, M2: MCHA-4, M1L: MCHA-32 w/ more SRAM, M2L: MCHA-4 w/ more SRAM, F1: FPGA-Gym, P1: PolyGraph, AN: ActiveN, NE: Neuro Engine, G1: A100, G2: RTX3090, G3: RTX 2090 Super, G4: TitanX, XC: Xeon CPU, DA: Dalorex, T: Tesseract, MO1: MCHA w/o asynchronous communication, MO2: MCHA w/o mulit-tiered NoC, MO3: MCHA w/o data-driven model, MO4: MCHA w/o 2D-Mesh, PR: PageRank, SM:SMAX

Figure 8: The speedup compared with the existing (a) MARL and BSP; (b) more BSP; (c) MVC; (d) BSP-MRF accelerators. (e) The ablation study of MCHA-4. (f) The detailed contribution of different characteristics of MCHA.

a near-optimal bandwidth utilization rate in regular data processing tasks, while also presenting the drawbacks of MCHA on computing-intensive tasks because of the limited scale. Figure 8(e) and (f) present an ablation study of MCHA-32, starting from a bus-based architecture and incrementally adding one feature at a time until reaching the full MCHA design. Each feature’s contribution is quantified by the latency reduction it yields over the bus-based baseline on a given workload. The results show that the bulk of the speedup over GPUs stems from the multi-tiered NoC and the data-driven model proposed by MCHA, with the asynchronous communication strategy providing additional gains. This breakdown confirms the effectiveness of MCHA’s core tenets, i.e., data localization and data-driven programming. 7.3 Architectural Analysis To elucidate the fundamental reasons for the observed MCHA performance gains, we explore the architectural design space and profile the benchmark. We first examine the simulation results of the end-to-end throughput measure-

12

To further demonstrate the efficacy of the MCHA architecture, we compare the throughput of MCHA-32, MCHA-4, and an NVIDIA A100 on MARL workloads after factoring out GPU software framework overhead. Even after this normalization, MCHA-32 still achieves 4.01∼21.93× speedup. This overhead originates predominantly from XLA [80] justin-time compilation, which, though intended to optimize

Inter-Block Comm.

(e)

0.02 0.04 0.06 0.08 0.10 Active Cumulative Time (ms)

Intensity (10-3 %)

0.4 0.3 0.2 0.1 0.0

Intra-Block Comm.

STORM

M 1 M G 2 O 1 G 1

M 1 M G 2 O 1 G 1

M 1 M G 2 O 1 G 1 M 1 M G 2 O 1 G 1

0.8 0.6 0.4 0.2 0.0 0.0

Inter-Chip Comm.

(d) Intensity (10-3 %) 2.0 1.5 1.0 0.5 0.0 0.01 0.02 0.03 0.04 0.05 Active Cumulative Time (ms) (f) Intensity (10-3 %) 1.2 0.8 0.4 0.0

0.0 0.05 0.10 0.15 0.20 0.25 0.30 0.35 Active Cumulative Time (ms)

0.01 0.02 0.03 0.04 0.05 0.06 Active Cumulative Time (ms)

Figure 9: (a) The total data movement of MCHA-4 and RTX3090. (b) The throughput of MCHA-32, MCHA-4 and A100 GPU with and without the software framework time costs. The computing intensity of (c) MPE Simple Spread, (d) Switch Riddle, (e) Starcraft, and (f) STORM of MCHA. 104 103 102 101 100 10-1

Throughput (GFLOPS)

MPE Simple Spread Starcraft

STORM Switch Riddle MCHA-4 Peak A100 Throughput Intra-Block Comm. Inter-Block Comm. DRAM Access Inter-Chip Comm.

10-2 10-1 10-0 101 102 103 104 Arithmetic Density (FLOP/Byte)

Figure 10: Roofline model analysis for MCHA-4 and A100. (a)

MCHA-4

Throughput (Steps / s)

MCHA-8

MCHA-16

MCHA-32

104 103 10 (b)

20

40

80

Throughput (Steps / s)

160 320 640 Number of Agents

1280 2560 5120 10240

0

0 92

81

0

96

48

40

0 24

20

20

10

60

Number of Neurons

51

0

80

25

12

0

64

32

16

0

104 103 80

A key insight from this breakdown is that these workloads are overwhelmingly communication-dominant. Intra-block and inter-block data movements constitute 57.3%∼69.1% of operational cycles, which underscores why conventional GPUs struggle with PSC tasks with the detailed memory footprints of MCHA in Table 3. Lacking an efficient P2P fabric for inter-core data exchange, GPUs force high-frequency transmissions through the global memory hierarchy. This can be illustrated by profiling the MPE Simple Spread benchmark on an RTX3090 with Nsight Systems. The essential data per agent amounts to 8 bytes, comprising the weight vector (hidden dimension 64) and the environment state. As shown in Figure9(a), the total off-chip memory traffic of the RTX 3090 running MPE Simple Spread with 1,000 agents over 1,000 steps is 14.88× larger than that of MCHA-4, a disparity driven by cache-coherence traffic across L2 cache lines. The reason MCHA-4’s data access still exceeds the application’s essential working set is the four-fold replication of the neural network, stored once per chip. Nevertheless, MCHA-4’s multi-tiered NoC preserves data locality and bypasses off-chip memory and the global buffer. Compared to the GPU computing intensity in Figure 1(d) in Section 1, where DRAM access accounts for 78.9% of the total operations, DRAM traffic in MCHA occupies only 2.26%∼5.44%. This significant reduction validates that MCHA successfully transforms global memory bottlenecks into localized on-chip data traffic.

Compute

Intensity (10-3 %)

40

Figure 9 shows the results of profiling runtime operational types for MCHA-4. The x-axis is the execution time. The y-axis shows the operational intensity in percentage: “compute” for computation, “memory” for memory I/O access operations, “Comm.” for the three types of NoC operations.

Memory (c)

20

7.3.1 Computing Intensity Analysis

RTX 3090 (b) Throughput (Steps/s) 1410.6KB 107 MPE Simple Switch Starcraft MCHA-4 5 Spread Riddle 10 94.8 KB Extra Data 103 68.8% 101 97.9% 10-1 Essential Data 29.6KB GO1: GPU exclude software framework

(a)

10

ment and find the components of the processing time of MCHA-4. We also provide the memory footprint of MCHA4 on MARL workloads. Based on these analyses, we formulate a customized roofline model, establishing the operational upper bounds for MCHA. Utilizing this model, we examine the scalability characteristics of MCHA to identify the structural features that enable near-linear performance growth across increasing core counts and the variation of speedup factor between MCHA-4 and MCHA-32 in different cases. Further, we conduct a sensitivity analysis, quantifying how variations in multi-level interconnect bandwidths and RISC-V compute intensity, as detailed in Table 1, impact the end-to-end throughput of PSC applications.

Figure 11: The scale-up effect of (a) SMAX (b) MVC.

memory-bound workloads, consumes 97.5%∼99.79% of total GPU execution time. Such a stark imbalance underscores the advantage of the proposed data-driven programming model: it statically determines the instruction schedule from the base algorithm while sustaining high hardware utilization. The overlap of computation and data movement further validates the model’s ability to hide transmission latency within the pipeline. 7.3.2 Roofline Model Analysis Figure 10 shows the roofline models for MCHA-4 and A100. Unlike conventional GPUs, which are typically bounded by a single global memory bandwidth (as shown in Figure 1(a)), MCHA exhibits a tiered execution profile characterized by 4 bandwidth limits. While the hierarchy of these bandwidth limits is fixed by the hardware design, the actual throughput bottleneck is dynamically determined by its specific algorithmic dataflow. For example, STORM and Switch Rid-

13

dle workloads are constrained by the Intra-Block Communication bandwidth. Though the limitation is slight, this suggests that their high-frequency agent interactions saturate the local communication bandwidth before reaching the MCC’s compute limits. The remaining workloads are limited by the MCC computing throughput, operating within the bandwidth headroom provided by MCHA-4. This divergence emphasizes the key role of the memory-centric nature of the proposed data-driven programming principles. By strategically organizing tasks to maximize data residency, MCHA minimizes high-latency transmissions and shifts the operational point toward the compute-bound region for higher throughput.

MCC Throughput Computing Latency

Intra-Block Bandwidth Inter-Block Bandwidth

(a) Throughput (GFLOPS) 104

(b)

103 102

Changing Direction

101 10

0

10-2

10-1 100 101 102 103 Arithmetic Density (FLOP/Byte) (c) Throughput (GFLOPS) 104 103

104

Changing Direction

102 101 100

10-1 100 101 102 103 10-2 Arithmetic Density (FLOP/Byte) (e) Throughput (GFLOPS) 104

104

103 102

Changing Direction

101 100

7.3.3 Scaling-Up Effect The architectural scalability of MCHA is fundamentally driven by its high spatial and temporal locality. Workloads with spatially organized neighborhood definitions, such as large-scale maps containing thousands of agents, are partitioned using a dynamic border strategy, as illustrated in Figure 7(e). This allows the workload to be distributed across MCCs while maintaining data residency. Conversely, for graph-based neighborhood topologies, a min-cut partitioning algorithm can be implemented to decompose the global graph into localized sub-graphs. In this configuration, intersub-graph edge data is temporarily buffered in DRAM, ensuring that the high-frequency internal processing remains localized within the on-chip mesh. Figure 11 presents the scaling-up characteristics of the MCHA hardware across varying workload sizes and hardware scales, using SMAX and MVC as the benchmark. The persistent parameters and corresponding entities are spread evenly on the hardware for specific comparison. Figure 11(a) shows that at small agent counts (10 to 40 agents in SMAX), the system throughput remains relatively constant. This phenomenon occurs because the primary overhead of the execution at this time is not the compute intensity or communication volume, but the pipeline latency required to complete the initial dataflow sequence. As the agent population grows, the dominant bottleneck shifts from pipeline latency to compute and transmission throughput. In this regime, the processing time increases proportionally with the number of agents. Additionally, Figure 11(a) also shows that the accelerating rates of MCHA-32 compared with MCHA-4 vary when the benchmark scale is different. This is because the even distribution of the agents in the hardware leads to performance limits originating from the Inter-Chip Communication. Similar trends can be observed from Figure 11(b) for the MVC benchmark. Comparing Figure 11(a) with (b), for 2% MVC benchmark, MCHA-32 achieves 7.47× speedup compared with MCHA-4, while this ratio decreases to 3.85× for SMAX, including 10240 agents. The reason for this phenomenon is that the arithmetic density of SMAX is higher than that of the MVC with a reinforce-

10-1 100 101 102 103 10-2 Arithmetic Density (FLOP/Byte) (g) Throughput (GFLOPS) 104

104

103 102

Changing Direction

101 10

0

10-2

(i) 10

10-1 100 101 102 103 Arithmetic Density (FLOP/Byte) Throughput (GFLOPS) 4

104

103 102

Changing Direction

101 100 10-2

10-1 100 101 102 103 Arithmetic Density (FLOP/Byte)

104

80 60 40 20 0

global DRAM Bandwidth Inter-Chip Bandwidth Latency (103 Cycles / Step) Lower is Better Changing Direction

102.4 51.2 25.6 12.8 6.4 MCC Throughput (GFLOPS) (d) Latency (103 Cycles / Step) 600 Lower is Better 500 Changing 400 Direction 300 200 100 0 2304 576 144 36 9 Intra-Block Bandwidth (GB/s) (f) Latency (103 Cycles / Step) 600 Lower is Better 500 Changing 400 Direction 300 200 100 0 1024 256 64 16 4 Inter-Block Bandwidth(GB/s) (h) Latency (103 Cycles / Step) 600 500 Lower is Better 400 Changing Direction 300 200 100 0 64 32 16 8 4 Inter-Chip Bandwidth (GB/s) (j) 3 Latency (10 Cycles / Step) 600 500 Lower is Better 400 Changing Direction 300 200 100 0 409.6102.4 25.6 6.4 1.6 global DRAM Bandwidth (GB/s)

Figure 12: The roofline model and throughput comparison based on different (a–b) MCC throughput; (c–d) Intra-block bandwidth; (e–f) Inter-block bandwidth; (g–h) Inter-chip bandwidth; (i–j) global DRAM bandwidth.

ment neural network included as the updating function.

7.3.4 Design Space Exploration According to the roofline model presented in Section 4, the design space of MCHA can be described by multiple bandwidths and computing limits. Figure 12 shows the influence of varying parameters in MCHA-4 on the roofline models. The varying design parameters cover MCC throughput (a– b) and the bandwidths of Intra-block (c–d), Inter-block (e– f), Inter-chip (g–h), and global DRAM (i–j). Inter-chip bandwidth and global DRAM bandwidth have the largest impact on the overall effective throughput. This phenomenon can be observed through the slight growth of the processing time before arriving at the turning point of the roofline model, as the decrease of inter-chip bandwidth and global DRAM bandwidth. This indicates that the communication between different chips or off-chip memory should be controlled in this MCHA architecture, which is straightforward for our design. In other cases, the influence of the parameters on the final throughput is similar to the conventional roofline model.

14

7.4 Related Works MCHA vs. Compute-Centric Architectures and Frameworks. While many-core CPUs and GPUs are the current mainstream hardware solution executing PSC tasks, they are doomed to hit the “memory wall” [20–22]. Software frameworks like JAX-MARL [49] attempt to alleviate this through JIT-compiled kernel fusion, yet they are still throttled by rigid hardware cache hierarchies. For MARL workloads, this leads to massive cache coherence overhead, which can consume up to 96% of total operations as shown in Figure 1(b). In contrast, MCHA abandons GBs, instead granting programmers explicit control over a reconfigurable memory fabric through the multi-tiered NoC. MCHA vs. Near-Data Processing. While NDP and Processing-In-Memory (PIM) designs like Tesseract [81] and UPMEM [82] successfully reduce data movement, they are generally tailored for static irregularity, such as PageRank [72] or BFS [73] on static graphs. Conversely, PSC is defined by dynamic irregularity with an additional time dimension, where neighbor relationships change as the algorithm progresses. Existing PIM implementations, such as StreamPIM [39], frequently lack the inter-core bandwidth and architectural flexibility needed for the complex phase transitions of these tasks. MCHA bridges this gap by utilizing a multi-layer communication strategy for P2P entity interaction that releases the burden of off-chip memory access. MCHA vs. Specialized Accelerators. Accelerators for specific irregular primitives can achieve high efficiency for their targeted kernels, such as pointer-chasing in graph processing (Minnow [83], GraphR [84]) or sparse-tensor algebra (ExTensor [85]). However, they often lack support for general PSC tasks, making commercialization challenging with only supporting highly niche computing applications. Furthermore, many DSAs still rely on a centralized global buffer, which becomes a serializing bottleneck as the computing scale increases. MCHA differentiates itself by proposing the data-driven programming model, enabling a reconfigurable fabric to support a diverse suite of workloads.

8 Conclusion We propose MCHA architecture, which accelerates PSC workloads, a critical paradigm spanning MARL, SNN, and BSP, by enabling P2P communication to relieve the global bandwidth pressure that bottlenecks conventional systems. A data-driven programming model is provided for algorithm deployment. The architecture demonstrates strong data and processing locality, achieving speedups of 153.06× to 2456.96× over A100 GPUs in MARL tasks using a 4-chip system. Its generality and near-linear scalability are further validated through multi-application evaluation and roofline analysis. Two limitations merit mention: (a) MCHA exploits locality through fixed topologies, leaving highly dynamic graphs unsupported; (b) porting CPU/GPU code to MCHA’s programming model currently requires manual effort. We envi-

sion that advances in workload abstraction and agentic-AIdriven code-conversion toolkits will together unlock MCHA’s full potential.

References [1] Norman P Jouppi, Cliff Young, Nishant Patil, David Patterson, Gaurav Agrawal, Raminder Bajwa, Sarah Bates, Suresh Bhatia, Nan Boden, Al Borchers, Rick Boyle, Pierre-luc Cantin, Clifford Chao, Chris Clark, Jeremy Coriell, Mike Daley, Matt Dau, Jeffrey Dean, Ben Gelb, Tara Vazir, Ghaemmaghami, Rajendra Gottipati, William Gulland, Robert Hagmann, C. Richard Ho, Doug Hogberg, John Hu, Robert Hundt, Dan Hurt, Julian Ibarz, Aaron Jaffey, Alek Jaworski, Alexander Kaplan, Harshit Khaitan, Daniel Killebrew, Andy Koch, Naveen Kumar, Steve Lacy, James Laudon, James Law, Diemthu Le, Chris Leary, Zhuyuan Liu, Kyle Lucke, Alan Lundin, Gordon MacKean, Adriana Maggiore, Maire Mahony, Kieran Miller, Rahul Nagarajan, Ravi Narayanaswami, Ray Ni, Kathy Nix, Thomas Norrie, Mark Omernick, Narayana Penukonda, Andy Phelps, Jonathan Ross, Matt Ross, Amir Salek, Emad Samadiani, Chris Severn, Gregory Sizikov, Matthew Snelham, Jed Souter, Dan Steinberg, Andy Swing, Mercedes Tan, Gregory Thorson, Bo Tian, Horia Toma, Erick Tuttle, Vijay Vasudevan, Richard Walter, Walter Wang, Eric Wilcox, and Doe Hyun Yoon. In-datacenter performance analysis of a tensor processing unit. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 1–12. IEEE, 2017. [2] Mike Davies, Narayan Srinivasa, Tsung-Han Lin, Gautham Chinya, Yongqiang Cao, Sri Harsha Choday, Georgios Dimou, Prasad Joshi, Nabil Imam, Shweta Jain, Yuyun Liao, Chit-Kwan Lin, Andrew Lines, Ruokun Liu, Deepak Mathaikutty, Steven McCoy, Arnab Paul, Jonathan Tse, Guruguhanathan Venkataramanan, Yi-Hsin Weng, Andreas Wild, Yoonseok Yang, and Hong Wang. Loihi: A neuromorphic manycore processor with on-chip learning. IEEE Micro, 38(1):82–99, 2018. [3] Gene M. Amdahl. Validity of the single processor approach to achieving large scale computing capabilities. In Spring Joint Computer Conference (SJCC), AFIPS ’67 (Spring), page 483–485, New York, NY, USA, 1967. Association for Computing Machinery. ISBN 9781450378956. doi: 10.1145/1465482.1465560. URL https://doi.org/10.1145/1465482.1465560. [4] Peng Chen, Xinyu Zhao, Lina Zeng, Luxinyu Liu, Shengjie Liu, Li Sun, Zaijin Li, Hao Chen, Guojun Liu, Zhongliang Qiao, Yi Qu, Dongxin Xu, Lianhe Li, and Lin Li. A review of research on SLAM technology based on the fusion of LiDAR and vision. Sensors, 25(5):1447, 2025.

15

[5] Yiming Gao, Chao Jiang, Wesley Piard, Xiangru Chen, Bhavesh Patel, and Herman Lam. HgPCN: A heterogeneous architecture for E2E embedded point cloud inference. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1588–1600, 2024. doi: 10.1109/MICRO61859.2024.00116. [6] Jianye Hao, Tianpei Yang, Hongyao Tang, Chenjia Bai, Jinyi Liu, Zhaopeng Meng, Peng Liu, and Zhen Wang. Exploration in deep reinforcement learning: From single-agent to multiagent domain. IEEE Transactions on Neural Networks and Learning Systems, 35 (7):8762–8782, 2023. [7] Thanh Thi Nguyen, Ngoc Duy Nguyen, and Saeid Nahavandi. Deep reinforcement learning for multiagent systems: A review of challenges, solutions, and applications. IEEE Transactions on Cybernetics, 50(9):3826– 3839, 2020. [8] Zepeng Ning and Lihua Xie. A survey on multi-agent reinforcement learning and its application. Journal of Automation and Intelligence, 3(2):73–91, 2024. [9] David Beniaguev, Idan Segev, and Michael London. Single cortical neurons as deep artificial neural networks. Neuron, 109(17):2727–2739, 2021. [10] Carver Mead. Neuromorphic electronic systems. Proceedings of the IEEE, 78(10):1629–1636, 2002. [11] Kaushik Roy, Akhilesh Jaiswal, and Priyadarshini Panda. Towards spike-based machine intelligence with neuromorphic computing. Nature, 575(7784): 607–617, 2019. [12] Youhui Zhang, Peng Qu, Yu Ji, Weihao Zhang, Guangrong Gao, Guanrui Wang, Sen Song, Guoqi Li, Wenguang Chen, Weimin Zheng, Feng Chen, Jing Pei, Rong Zhao, Mingguo Zhao, and Luping Shi. A system hierarchy for brain-inspired computing. Nature, 586 (7829):378–384, 2020. [13] Boyao Li, Alexander J Thomson, Houssam Nassif, Matthew M Engelhard, and David Page. On neural networks as infinite tree-structured probabilistic graphical models. In Conference on Neural Information Processing Systems (NeurIPS), volume 37, pages 4598–4628, 2024. [14] Marcel Kollovieh, Bertrand Charpentier, Daniel Zügner, and Stephan Günnemann. Expected probabilistic hierarchies. In Conference on Neural Information Processing Systems (NeurIPS), volume 37, pages 13818– 13850, 2024. [15] Shivvrat Arya, Tahrima Rahman, and Vibhav Giridhar Gogate. SINE: Scalable MPE inference for probabilistic graphical models using advanced neural embeddings. In International Conference on Artificial Intelligence and Statistics (AISTATS), 2025.

[16] Jeffrey Dean and Sanjay Ghemawat. MapReduce: simplified data processing on large clusters. Communications of the ACM, 51(1):107–113, 2008. ISSN 00010782. doi: 10.1145/1327452.1327492. URL https: //doi.org/10.1145/1327452.1327492. [17] Grzegorz Malewicz, Matthew H. Austern, Aart J.C Bik, James C. Dehnert, Ilan Horn, Naty Leiser, and Grzegorz Czajkowski. Pregel: a system for large-scale graph processing. In ACM SIGMOD International Conference on Management of Data (SIGMOD), SIGMOD ’10, page 135–146, New York, NY, USA, 2010. Association for Computing Machinery. ISBN 9781450300322. doi: 10.1145/1807167.1807184. URL https://doi. org/10.1145/1807167.1807184. [18] Matei Zaharia, Mosharaf Chowdhury, Tathagata Das, Ankur Dave, Justin Ma, Murphy McCauly, Michael J. Franklin, Scott Shenker, and Ion Stoica. Resilient distributed datasets: A Fault-Tolerant abstraction for In-Memory cluster computing. In USENIX Symposium on Networked Systems Design and Implementation (NSDI), pages 15–28, San Jose, CA, 2012. USENIX Association. ISBN 978-931971-92-8. URL https://www.usenix.org/conference/nsdi12/ technical-sessions/presentation/zaharia. [19] Vidushi Dadu, Sihao Liu, and Tony Nowatzki. Polygraph: Exposing the value of flexibility for graph processing accelerators. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 595–608. IEEE, 2021. [20] Krste Asanovic, Ras Bodik, Bryan Christopher Catanzaro, Joseph James Gebis, Parry Husbands, Kurt Keutzer, David A Patterson, William Lester Plishker, John Shalf, Samuel Webb Williams, and Katherine A. Yelick. The landscape of parallel computing research: A view from Berkeley. Technical report, Technical Report UCB/EECS-2006-183, EECS Department, University of California at Berkeley, 2006. [21] Brian M Rogers, Anil Krishna, Gordon B Bell, Ken Vu, Xiaowei Jiang, and Yan Solihin. Scaling the bandwidth wall: challenges in and avenues for CMP scaling. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 371–382, 2009. [22] Andrei Ivanov, Nikoli Dryden, Tal Ben-Nun, Shigang Li, and Torsten Hoefler. Data movement is all you need: A case study on optimizing transformers. Proceedings of Machine Learning and Systems, 3:711–732, 2021. [23] Martin Burtscher, Rupesh Nasre, and Keshav Pingali. A quantitative study of irregular programs on GPUs. In IEEE International Symposium on Workload Characterization (IISWC), pages 141–151. IEEE, 2012. [24] Seunghee Shin, Guilherme Cox, Mark Oskin, Gabriel H Loh, Yan Solihin, Abhishek Bhattacharjee, and

16

Arkaprava Basu. Scheduling page table walks for irregular gpu applications. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 180–192. IEEE Computer Society, 2018. [25] Raghu Prabhakar, Yaqi Zhang, David Koeplinger, Matt Feldman, Tian Zhao, Stefan Hadjis, Ardavan Pedram, Christos Kozyrakis, and Kunle Olukotun. Plasticine: A reconfigurable architecture for parallel paterns. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 389–402, 2017. [26] Raghu Prabhakar, Ram Sivaramakrishnan, Darshan Gandhi, Yun Du, Mingran Wang, Xiangyu Song, Kejie Zhang, Tianren Gao, Angela Wang, Xiaoyan Li, Yongning Sheng, Joshua Brot, Denis Sokolov, Apurv Vivek, Calvin Leung, Arjun Sabnis, Jiayu Bai, Tuowen Zhao, Mark Gottscho, David Jackson, Mark Luttrell, Manish K. Shah, Zhengyu Chen, Kaizhao Liang, Swayambhoo Jain, Urmish Thakker, Dawei Huang, Sumti Jairath, Kevin J. Brown, and Kunle Olukotun. Sambanova SN40L: Scaling the AI memory wall with dataflow and composition of experts. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1353–1366. IEEE, 2024. [27] Brian C Schwedock and Nathan Beckmann. Leviathan: A unified system for general-purpose near-data computing. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1278–1294. IEEE, 2024. [28] S Alexander Chin, Noriaki Sakamoto, Allan Rui, Jim Zhao, Jin Hee Kim, Yuko Hara-Azumi, and Jason Anderson. CGRA-ME: A unified framework for CGRA modelling and exploration. In IEEE International Conference on Application-Specific Systems, Architectures and Processors (ASAP), pages 184–189. IEEE, 2017. [29] Zhaoying Li, Pranav Dangi, Chenyang Yin, Thilini Kaushalya Bandara, Rohan Juneja, Cheng Tan, Zhenyu Bai, and Tulika Mitra. Enhancing cgra efficiency through aligned compute and communication provisioning. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS), pages 410–425, 2025. [30] Jiajun Qin, Tianhua Xia, Cheng Tan, Jeff Zhang, and Sai Qian Zhang. PICACHU: Plug-in CGRA handling upcoming nonlinear operations in LLMs. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems, pages 845–861, 2025. [31] Shaizeen Aga, Supreet Jeloka, Arun Subramaniyan, Satish Narayanasamy, David Blaauw, and Reetuparna Das. Compute caches. In IEEE International Symposium on High Performance Computer Architecture (HPCA), pages 481–492. IEEE, 2017.

[32] Anant Agarwal, Ricardo Bianchini, David Chaiken, Kirk L Johnson, David Kranz, John Kubiatowicz, BengHong Lim, Kenneth Mackenzie, and Donald Yeung. The MIT Alewife machine: Architecture and performance. ACM SIGARCH Computer Architecture News, 23(2):2–13, 1995. [33] Sam Ainsworth and Timothy M Jones. Graph prefetching using data structure knowledge. In International Conference on Supercomputing (ICS), pages 1–11, 2016. [34] Sam Ainsworth and Timothy M Jones. An eventtriggered programmable prefetcher for irregular workloads. ACM Sigplan Notices, 53(2):578–592, 2018. [35] Ruiyang Chen, Zhuoran Song, Yicheng Zheng, Zeyu Zhu, Gang Li, Naifeng Jing, Xiaoyao Liang, and Haibing Guan. HEAT: NPU-NDP heterogeneous architecture for transformer-empowered graph neural networks. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 263–276, 10 2025. doi: 10.1145/3725843.3756117. [36] Mohammadreza Saed, Prashant Nair, and Tor Aamodt. RayN: Ray tracing acceleration with near-memory computing. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 277–291, 10 2025. doi: 10.1145/3725843.3756067. [37] Je-Woo Jang, Junyong Oh, Youngbae Kong, JaeYoun Hong, Sung-Hyuk Cho, Jeongyeol Lee, Hoeseok Yang, and Joon-Sung Yang. Accelerating retrieval augmented language model via PIM and PNM integration. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 246–262, 10 2025. doi: 10. 1145/3725843.3756020. [38] Wonung Kim, Yubin Lee, Yoonsung Kim, Jinwoo Hwang, Seongryong Oh, Jiyong Jung, Aziz Huseynov, Woong Park, Chang Park, Divya Mahajan, and Jongse Park. Pimba: A processing-in-memory acceleration for post-transformer large language model serving. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 292–307, 07 2025. doi: 10. 48550/arXiv.2507.10178. [39] Yuda An, Yunxiao Tang, Shushu Yi, Li Peng, Xiurui Pan, Guangyu Sun, Zhaochu Luo, Qiao Li, and Jie Zhang. StreamPIM: Streaming matrix computation in racetrack memory. In IEEE International Symposium on High-Performance Computer Architecture (HPCA), pages 297–311, 2024. doi: 10.1109/HPCA57654. 2024.00031. [40] Seah Kim, Hasan Genc, Vadim Vadimovich Nikiforov, Krste Asanović, Borivoje Nikolić, and Yakun Sophia Shao. MoCA: Memory-centric, adaptive execution for

17

multi-tenant deep neural networks. In IEEE International Symposium on High-Performance Computer Architecture (HPCA), pages 828–841, 2023. doi: 10. 1109/HPCA56546.2023.10071035. [41] Jeehyun Kim, Donghyeon Kim, Seokwon Kang, Bongjoon Hyun, Inho Lee, and Yongjun Park. PIMCCA: An efficient PIM architecture with optimized integration of configurable functional units. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 63–77, 10 2025. doi: 10.1145/3725843.3756034. [42] Hyojun Son, Gilbert Jonatan, Xiangyu Wu, Haeyoon Cho, Kaustubh Shivdikar, José L. Abellán, Ajay Joshi, David Kaeli, and John Kim. PIMnet: A domain-specific network for efficient collective communication in scalable PIM. In IEEE International Symposium on HighPerformance Computer Architecture (HPCA), pages 1557–1572, 2025. doi: 10.1109/HPCA61900.2025. 00116. [43] Seunghyuk Yu, Hyeonu Kim, Kyoungho Jeun, Sunyoung Hwang, Seongmin Cho, and Eojin Lee. ComPASS: A compatible PIM protocol architecture and scheduling solution for processor-PIM collaboration. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 49–62, 10 2025. doi: 10.1145/ 3725843.3756017. [44] Hwanjun Lee, Minho Kim, Yeji Jung, Seonmu Oh, KiDong Kang, Seunghak Lee, and Daehoon Kim. Beyond page migration: Enhancing tiered memory performance via integrated last-level cache management and page migration. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1763– 1776, 10 2025. doi: 10.1145/3725843.3756063. [45] Houxiang Ji, Yifan Yuan, Yang Zhou, Ipoom Jeong, Ren Wang, Saksham Agarwal, and Nam Kim. Rearchitecting end-host networking with CXL: Coherence, memory, and offloading. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1809–1823, 10 2025. doi: 10.1145/3725843. 3756102. [46] Jinshu Liu, Hamid Hadian, Yuyue Wang, Daniel Berger, Marie Nguyen, Xun Jian, Sam Noh, and Huaicheng Li. Systematic CXL memory characterization and performance analysis at scale. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS), pages 1203–1217, 03 2025. doi: 10.1145/ 3676641.3715987. [47] Yan Sun, Jongyul Kim, Zeduo Yu, Jiyuan Zhang, Siyuan Chai, Michael Kim, Hwayong Nam, Jaehyun Park, Eojin Na, Yifan Yuan, Ren Wang, Jung Ho Ahn, Tianyin Xu, and Nam Kim. M5: Mastering page migration and

memory management for CXL-based tiered memory systems. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS), pages 604–621, 03 2025. doi: 10.1145/3676641.3711999. [48] J Terry, Benjamin Black, Nathaniel Grammel, Mario Jayakumar, Ananth Hari, Ryan Sullivan, Luis S Santos, Clemens Dieffendahl, Caroline Horsch, Rodrigo PerezVicente, Caroline Horsch, Clemens Dieffendahl, Niall L. Williams, Yashas Lokesh, and Praveen Ravi. Pettingzoo: Gym for multi-agent reinforcement learning. In Conference on Neural Information Processing Systems (NeurIPS), volume 34, pages 15032–15043, 2021. [49] Alexander Rutherford, Benjamin Ellis, Matteo Gallici, Jonathan Cook, Andrei Lupu, Garðar Ingvarsson, Timon Willi, Ravi Hammond, Akbir Khan, Christian Schroeder de Witt, Alexandra Souly, Saptarashmi Bandyopadhyay, Mikayel Samvelyan, Minqi Jiang, Robert Tjarko Lange, Shimon Whiteson, Bruno Lacerda, Nick Hawes, Tim Rocktäschel, Chris Lu, and Jakob Nicolaus Foerster. JaxMARL: Multi-agent rl environments and algorithms in JAX. In Conference on Neural Information Processing Systems (NeurIPS), 2024. [50] Jakob Foerster, Yannis M Assael, Nando de Freitas, and Shimon Whiteson. Learning to communicate with deep multi-agent reinforcement learning. In Conference on Neural Information Processing Systems (NeurIPS), pages 2137–2145, 2016. [51] Christopher Lu, Timon Willi, Christian A Schroeder De Witt, and Jakob Foerster. Model-free opponent shaping. In International Conference on Machine Learning (ICML), pages 14398–14411. PMLR, 2022. [52] Jakob Foerster, Richard Chen, Maruan Al-Shedivat, Shimon Whiteson, Pieter Abbeel, and Igor Mordatch. Learning with opponent-learning awareness. In International Conference on Autonomous Agents and MultiAgent Systems (AAMAS), pages 122–130, 07 2018. doi: 10.65109/HGWA8807. [53] John L Gustafson. Reevaluating Amdahl’s law. Communications of the ACM, 31(5):532–533, 1988. [54] Erik Lindholm, John Nickolls, Stuart Oberman, and John Montrym. NVIDIA tesla: A unified graphics and computing architecture. IEEE Micro, 28(2):39–55, 2008. [55] Michael J Flynn. Some computer organizations and their effectiveness. IEEE Transactions on Computers, 100(9):948–960, 2009. [56] Lie, Sean. Cerebras architecture deep dive: First look inside the hardware/software co-design for deep learning. IEEE Micro, 43(3):18–30, 2023.

18

[57] David Mayhew and Venkata Krishnan. PCI express and advanced switching: Evolutionary path to building next generation interconnects. In IEEE Symposium on High Performance Interconnects (HOTI), pages 21–29. IEEE, 2003.

[68] Xiaoyi Liu, Zhongzhu Pu, Peng Qu, Weimin Zheng, and Youhui Zhang. ActiveN: a scalable and flexiblyprogrammable event-driven neuromorphic processor. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 1122–1137. IEEE, 2024.

[58] Debendra Das Sharma, Gerald Pasdast, Zhiguo Qian, and Kemal Aygun. Universal chiplet interconnect express (UCIe): An open industry standard for innovations with chiplets at package level. IEEE Transactions on Components, Packaging and Manufacturing Technology, 12(9):1423–1431, 2022.

[69] Shirui Zhao, Jun Yin, Lingyun Yao, Martin Andraud, Wannes Meert, and Marian Verhelst. MC2 A: Enabling algorithm-hardware co-design for efficient markov chain monte carlo acceleration. arXiv preprint arXiv:2507.12935, 2025.

[59] Choquette, Jack. NVIDIA Hopper H100 GPU: Scaling performance. IEEE Micro, 43(3):9–17, 2023. [60] William James Dally and Brian Patrick Towles. Principles and practices of interconnection networks. Elsevier, 2004. [61] Abiodun M Ikotun, Absalom E Ezugwu, Laith Abualigah, Belal Abuhaija, and Jia Heming. K-means clustering algorithms: A comprehensive review, variants analysis, and advances in the era of big data. Information Sciences, 622:178–210, 2023. [62] Ümit Çatalyürek, Karen Devine, Marcelo Faraj, Lars Gottesbüren, Tobias Heuer, Henning Meyerhenke, Peter Sanders, Sebastian Schlag, Christian Schulz, Daniel Seemaier, and Dorothea Wagner. More recent advances in (Hyper)Graph partitioning. ACM Computing Surveys, 55(12):1–38, 2023. [63] Igor Mordatch and Pieter Abbeel. Emergence of grounded compositional language in multi-agent populations. In AAAI Conference on Artificial Intelligence (AAAI), 2017. [64] Bei Peng, Tabish Rashid, Christian Schroeder de Witt, Pierre-Alexandre Kamienny, Philip Torr, Wendelin Böhmer, and Shimon Whiteson. FACMAC: Factored multiagent centralised policy gradients. In Conference on Neural Information Processing Systems (NeurIPS), volume 34, pages 12208–12221, 2021. [65] Ching-Jui Lee and Tsung Tai Yeh. ReSA: Reconfigurable systolic array for multiple tiny dnn tensors. ACM Transactions on Architecture and Code Optimization, 21(3):1–24, 2024. [66] DeepSeek-AI. DeepSeek-R1 incentivizes reasoning in LLMs through reinforcement learning. Nature, 645 (8081):633–638, 2025. [67] Jiayi Li, Hongxiao Zhao, Wenshuo Yue, Yihan Fu, Daijing Shi, Anjunyi Fan, Yuchao Yang, and Bonan Yan. PEARL: FPGA-based reinforcement learning acceleration with pipelined parallel environments. In Design, Automation & Test in Europe Conference (DATE), pages 1–7. IEEE, 2025.

[70] Marcelo Orenes-Vera, Esin Tureci, David Wentzlaff, Margaret Martonosi. Dalorex: A data-local program execution and architecture for memory-bound applications. In IEEE International Symposium on HighPerformance Computer Architecture (HPCA), pages 718–730. IEEE, 2023. [71] Maximilian Schmidt, Rembrandt Bakker, Kelly Shen, Gleb Bezgin, Markus Diesmann, and Sacha Jennifer van Albada. A multi-scale layer-resolved spiking network model of resting-state dynamics in macaque visual cortical areas. PLOS Computational Biology, 14 (10):e1006359, 2018. [72] Lawrence Page, Sergey Brin, Rajeev Motwani, and Terry Winograd. The PageRank citation ranking: Bringing order to the web. Technical report, Stanford infolab, 1999. [73] Richard C Murphy, Kyle B Wheeler, Brian W Barrett, and James A Ang. Introducing the graph 500. Cray Users Group (CUG), 19(45-74):22, 2010. [74] Xiangyu Zhang, Ramin Bashizade, Yicheng Wang, Sayan Mukherjee, and Alvin R Lebeck. Statistical robustness of markov chain monte carlo accelerators. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS), pages 959–974, 2021. [75] Shirui Zhao, Nimish Shah, Wannes Meert, and Marian Verhelst. AIA: A 16nm multicore SoC for approximate inference acceleration exploiting nonnormalized knuth-yao sampling and inter-core register sharing. In IEEE European Solid-State Electronics Research Conference (ESSERC), pages 29–32. IEEE, 2024. [76] Glenn G Ko, Yuji Chai, Marco Donato, Paul N Whatmough, Thierry Tambe, Rob A Rutenbar, David Brooks, and Gu-Yeon Wei. A 3mm2 programmable bayesian inference accelerator for unsupervised machine perception using parallel gibbs sampling in 16nm. In IEEE Symposium on VLSI Circuits (VLSI), pages 1–2. IEEE, 2020. [77] Alan Mislove, Massimiliano Marcon, Krishna P Gummadi, Peter Druschel, and Bobby Bhattacharjee. Measurement and analysis of online social networks. In

19

ACM SIGCOMM Conference on Internet Measurement (IMC), pages 29–42, 2007. [78] Jure Leskovec, Deepayan Chakrabarti, Jon Kleinberg, Christos Faloutsos, Zoubin Ghahramani. Kronecker graphs: an approach to modeling networks. Journal of Machine Learning Research, 11(2), 2010. [79] Hunjun Lee, Chanmyeong Kim, Yujin Chung, and Jangwoo Kim. NeuroEngine: A hardware-based eventdriven simulation system for advanced brain-inspired computing. In ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS), pages 975–989, 2021. [80] Amit Sabne. XLA : Compiling Machine Learning for Peak Performance, 2020. [81] Junwhan Ahn, Sungpack Hong, Sungjoo Yoo, Onur Mutlu, and Kiyoung Choi. A scalable processing-inmemory accelerator for parallel graph processing. In ACM/IEEE Annual International Symposium on Computer Architecture (ISCA), pages 105–117. IEEE, 2015. [82] Birte Friesel, Marcel Lütke Dreimann, and Olaf Spinczyk. A full-system perspective on UPMEM performance. In The Workshop on Disruptive Memory Systems (DIMES), pages 1–7, 2023. [83] Dan Zhang, Xiaoyu Ma, Michael Thomson, and Derek Chiou. Minnow: Lightweight offload engines for worklist management and worklist-directed prefetching. ACM SIGPLAN Notices, 53(2):593–607, 2018. [84] Linghao Song, Youwei Zhuo, Xuehai Qian, Hai Li, and Yiran Chen. GraphR: Accelerating graph processing using ReRAM. In IEEE International Symposium on HighPerformance Computer Architecture (HPCA), pages 531–543. IEEE, 2018. [85] Kartik Hegde, Hadi Asghari-Moghaddam, Michael Pellauer, Neal Crago, Aamer Jaleel, Edgar Solomonik, Joel Emer, and Christopher W Fletcher. ExTensor: An accelerator for sparse tensor algebra. In ACM/IEEE International Symposium on Microarchitecture (MICRO), pages 319–333, 2019.

Related documents

• Reshaping the SDLC for Data- and AI-Centric Systems#463285
• ReMP: Low-Downtime Runtime Model-Parallelism Reconfiguration for LLM Serving#618115
• ReMP: Low-Downtime Runtime Model-Parallelism Reconfiguration for LLM Serving#617615
• ReMP: Low-Downtime Runtime Model-Parallelism Reconfiguration for LLM Serving#287100
• TierBPF: Page Migration Admission Control for Tiered Memory via eBPF#1052125