arXiv CSopen access

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

arXiv:2606.12638v1 [cs.DC] 10 Jun 2026

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads Ranganath Selagamsetty

Matthew Poremba

Bradford Beckmann

[email protected] University of Wisconsin-Madison Madison, Wisconsin, USA

[email protected] AMD Research and Advanced Development Bellevue, Washington, USA

[email protected] AMD Research and Advanced Development Bellevue, Washington, USA

Joshua San Miguel

Mikko Lipasti

[email protected] University of Wisconsin-Madison Madison, Wisconsin, USA

[email protected] University of Wisconsin-Madison Madison, Wisconsin, USA

Abstract

overheads, inter-GPU synchronization bottlenecks, and networkinduced latencies. Simply adding more GPUs does not guarantee proportional speedups, highlighting the need for architectural and software optimizations that can better exploit available parallelism. One promising strategy to mitigate these overheads is kernel fusion, where consecutive compute and communication operations are combined to reduce scheduling delays and improve resource utilization [30, 41, 46, 48, 49]. A recent example of this approach is the fused GEMV+AllReduce kernel [30], which merges a general matrix-vector multiplication (GEMV) with an AllReduce collective communication operation. This technique reduces launch overheads by overlapping computation with communication. Figure 1 shows a timeline profile of the fused GEMV+AllReduce kernel’s execution on a four-GPU system. In the ideal scenario, all GPUs spend the majority of the kernel execution performing arithmetic computations needed to calculate output tile results, noted by the timelines on all GPUs being predominantly green. This is ideal as no GPU spends excessive portions of their execution spin-waiting on a flag update, noted by the short red segments highlighted in Figure 1c). However, despite workload symmetry across GPUs, variability in runtime behavior still leads to imbalanced performance due to transient network traffic and contention. Figure 2 shows another instance of the fused GEMV+AllReduce kernel launch to the same four-GPU system. Despite running the same kernel on the same hardware, the distribution of time spent in regions of the kernel code greatly vary across the GPUs. Like in Figure 1, Figure 2b) shows the global view of the timeline profiles on all four GPUs. However, unlike Figure 1, GPUs two and three no longer exhibit ideal behavior. The majority of the fused GEMV+AllReduce kernel’s execution time is spent spin-waiting on a flag write update from a peer GPU, noted by the large red periods in the profiles. Due to variability from runtime characteristics, the multi-GPU system experiences load imbalance across the devices, leading to prolonged kernel execution times [11, 43]. As fused kernels become more prevalent [30, 41, 46, 48, 49], there is an urgent need for architectural tools that can accurately model and evaluate these behaviors. Existing GPU simulators, including gem5 [10] and gem5-gpu [29], offer cycle-level accuracy for microarchitectural research but lack native support for modeling multi-GPU systems and their associated communication dynamics.

As distributed AI workloads grow in scale, multi-GPU systems have become essential for training large models. Although techniques like kernel fusion and overlapping communication with computation help reduce delays, they also introduce irregular and transient traffic patterns that are difficult to model using existing tools. These techniques rely heavily on fine-grained synchronization and peer-to-peer communication, which place significant pressure on interconnect bandwidth and latency. In this work, we introduce Eidola, a scalable extension to the gem5 simulation framework that enables detailed modeling of interGPU communication traffic. The extension is scalable as our GPU model serves as a succinct eidolon, emulating the minimal characteristics needed for traffic modeling. Eidola uses annotated timing profiles from real applications to emulate peer-to-peer GPU writes with cycle-level precision. This allows researchers to simulate and analyze synchronization behavior across large multi-GPU configurations. The simulator supports configurable per-GPU traffic patterns and enables isolated performance analysis under different communication scenarios. We demonstrate Eidola’s effectiveness by reproducing variability in fused kernel execution and by implementing a SyncMoninspired synchronization mechanism, confirming reductions in polling-related memory traffic. Our results show that Eidola provides a flexible and scalable platform for studying inter-GPU communication and supports architectural exploration in modern distributed GPU systems.

Keywords Computer architecture research, gem5, instrumentation, multi-gpu systems

1

Introduction

The growing demand for large-scale machine learning models has ushered in an era of multiscale graphics processing unit (GPU) computing [14, 34], pushing system design toward exascale resources. Transformers and large language models (LLMs) continue to grow in parameter count, driving up the computational and communication demands required for training [26]. While scaling out with additional GPU resources can help alleviate the training burden, diminishing returns are often observed due to kernel launch 1

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

a) Zoomed in View of Flag Write Section

b) Broad, Multi-Gpu View of Ideal GEMV+AllReduce Timeline GPU 0

GPU 1

GPU 2

GPU 3

c) Zoomed-in View of Spin Wait Section Final Reduction

Workgroups

Workgroups

Output Tile Computation

Issue xGMI Write

Short Spin Wait Segment

Flag Write Time

Time

Figure 1: An ideal timing profile of the fused GEMV+AllReduce kernel on a four-GPU system. Figure 1b) shows a global view of the kernel’s execution on all four GPUs. In all four GPU timing profiles, each row represents a specific workgroup’s execution timeline, and time is shown on the horizontal axis. Figure 1a) highlights a period of the kernel’s execution where partial output tiles needed by remote GPUs are computed, noted by the successive patterns of green segments capped by brown markers. Once all partial output tiles needed by remote GPUs are completed, a write to the a peer GPU’s flag variable is issued, noted by the blue segment in Figure 1a). The green-brown segments between the blue and red marker indicate a GPU computing partial output tiles needed by the local GPU. Finally, Figure 1c) magnifies the final portion of the kernel. The red segment indicates the time a workgroup spins on a flag variable. The final green-brown segment indicates the final output tile reduction computation. The colors and code segments are consistent with those shown in Figure 3. a) Variability from hardware can lead to excessive spinning and prolonged total execution time

b) Broad, Multi-Gpu View of Non-Ideal GEMV+AllReduce Timeline GPU 0

GPU 1

GPU 2

GPU 3

Time

Time

Workgroups

Workgroups

Workgroups

GPU 2

Time

Figure 2: A non-ideal timing profile of the fused GEMV+AllReduce kernel on a four-GPU system. Like Figure 1b), Figure 2b) shows a global view of the kernel’s execution on all four GPUs, Figure 2a) shows a zoomed in view of GPU2’s execution of the fused GEMV+AllReduce kernel. The non-ideality is apparent via the long red portions in the timeline, where all the workgroups from GPU2 spend a majority of their time spin-waiting on peer GPU flag updates. This gap limits the ability of researchers to explore GPU design trade-offs in the context of distributed training workloads. In this work, we make the following contributions:

• Scalability studies that highlight the sub-linear simulation times to evaluate multi-GPU kernels that execute on systems with tens to hundreds of GPUs.

• A characterization exposing performance variability in fused kernel execution on mutli-GPU systems. • Eidola, a scalable extension to the gem5 simulator to support multi-GPU configurations with configurable communication topologies and traffic patterns. • A case study that implements the salient features of a recent work in GPU microarchitecture, SyncMon.

Our extension, Eidola, supports per-GPU analysis by incorporating timing and traffic profiles. In our study, these profiles were provided from real applications, but our framework can be used with synthetically generated profiles from probabilistic models. This enables researchers to evaluate the architectural impact of multi-GPU traffic on individual GPU performance in a controlled, cycle-level environment. This work provides a foundation for future 2

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

architectural explorations in distributed GPU systems, especially as machine learning training pipelines continue to scale. The rest of this paper is organized as follows. Section 2 provides background on the gem5 simulator, target workloads, and the training and instrumentation cycle for ML systems. Section 3 describes the design and implementation of Eidola. Section 4 presents experimental results, followed by a case study in Section 5. Section 6 discusses related work, and Section 7 identifies directions for future research. Section 8 concludes with a summary of our work.

2

novel fused kernel that combines a general matrix–vector multiplication (GEMV) operation with an AllReduce collective communication step. The GEMV kernel performs a dense linear algebra operation that multiplies a matrix by a vector and accumulates the result. The AllReduce operation, in contrast, is a distributed communication primitive that aggregates values across multiple GPUs and then broadcasts the aggregated result back to all participants. These kernels are instrumental in auto-regressive transformer models such as ChatGPT, LLaMA2 [44], and Megatron-LM [40], which rely heavily on repeated matrix operations and global synchronization of model parameters. The proposed GEMV+AllReduce fusion reduces the scheduling and synchronization delays typically encountered when these operations are launched as separate kernels. The fused kernel is especially beneficial in deep learning workloads such as transformer training, where fine-grained operations and frequent synchronizations across GPUs are common.

Background

This section provides the necessary background to contextualize our work. In Section 2.1, we describe the gem5 simulation framework, which serves as the foundation for our extensions. Section 2.2 introduces the target workload characteristics, focusing on fused computation–communication kernels commonly used in multiGPU systems, and discuss how these kernels rely on fine-grained inter-GPU synchronization. Finally, Section 2.3 outlines the typical training and profiling workflow in GPU clusters, highlighting how runtime instrumentation produces the timing information leveraged by our methodology.

2.1

1 __device__ void gemv_allreduce (...){ 2 for (tile : remote tiles to compute){ 3 compute partial tile results 4 xGMI write result to peer GPUS 5 } 6 7 xGMI write to flags[my_gpu] ∀ rgpus 8 9 for (tile : local tiles to compute){ 10 compute partial tile results 11 local write result to own GPUS 12 } 13 14 for (rgpu : remote gpus) 15 while (!flags[rgpu]); 16 17 reduce tiles 18 broadcast results 19 }

The gem5 Simulator

The gem5-gpu simulator [29] originally extended the gem5 simulation infrastructure [10] to model heterogeneous systems that combine CPUs and GPUs. Built by coupling gem5 with GPGPU-Sim [9], the tool enables full-system simulation of tightly integrated CPU-GPU architectures, supporting detailed modeling of memory hierarchies, coherence protocols, and system-level interactions. Over the past decade, gem5 has become a foundational platform for architectural research across domains including cache coherence [18], memory systems [13], interconnects [24, 28], and emerging accelerators [12, 20], due to its cycle-level accuracy, modularity, and community support. Upstream gem5 provides high-fidelity models of modern AMD GPU architectures, incorporating detailed execution pipelines, SIMT scheduling, memory subsystem behavior, and interconnect integration. These models have been rigorously validated and corroborated against empirical performance measurements from real AMD hardware, giving researchers confidence in its accuracy [31, 32, 36]. As such, gem5 enables researchers to evaluate architectural modifications in a realistic simulation context, making it a powerful tool for studying the performance implications of both low-level microarchitectural features and system-level interactions in heterogeneous systems. We build on gem5 to extend its capabilities for multi-GPU simulation and to better model the complex communication patterns observed in modern distributed machine learning workloads.

2.2

Compute remote tile Issue xGMI data write

Compute remote tile Issue xGMI data write

Issue xGMI flag write

Compute local tile Issue local data write

Compute local tile Issue local data write

Wait for peer xGMI flag write

Reduce tiles Broadcast results

Figure 3: Pseudocode for fused GEMV+AllReduce kernel. Colored blocks on the right match color coordination from Figures 1 and 2. A key architectural insight from the paper is the use of AMD’s external global memory interconnect (xGMI) to implement peerto-peer communication between GPUs. In this fused kernel design, inter-GPU synchronization is managed through xGMI writes, where one GPU directly writes to the memory space of a peer GPU. These writes target specific flag variables, which are then polled in a lightweight control loop to gate the control flow of subsequent operations. This eliminates the need for costly round-trip synchronization protocols by enabling fine-grained, low-latency signaling between GPUs. As a result, GPUs can coordinate progress through the fused kernel with minimal interruption, enabling better overlap of computation and communication and reducing idle time across the cluster. The key steps in the fused GEMV+AllReduce kernel are shown in Figure 3. Lines 2-5 note that the kernel first computes partial sums for output tiles needed by remote GPU. Once complete, each workgroup writes a flag update to peer GPUs to indicate completion, noted by line 7. Then, each workgroup proceeds to compute partial sums that will be reduced by the executing GPU. The results of these computation are stored locally on the executing

Fused Kernels for Multi-GPU Systems

To mitigate the performance bottlenecks arising from kernel launch overheads and inter-GPU synchronization in distributed machine learning workloads, recent work has explored the fusion of computation and communication primitives. Notably, [30] proposes a 3

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

GPU. The local GPU must now wait for all remote GPUs to compute the partial tiles needed by the local GPU, and for the flag updates to indicate so, noted by lines 3. Once signaled, the local GPU can perform a reduction to determine the final result for the output tile (line 17) and broadcasts the results to all GPUs (line 18). Inter-GPU communication through xGMI operates within the GPU memory hierarchy, which is organized into several caching layers and implements directory-based cache coherence. Each compute unit (CU) contains private L1 caches that service local memory accesses and instructions. These L1 caches are connected to a shared L2 cache, which may be partially or fully shared among multiple CUs depending on the GPU architecture. The L2 cache maintains coherence among the L1 caches and interfaces with the directory, which serves as the global coherence manager for the device. The directory tracks ownership and sharing states for cache lines, coordinating memory requests between the GPU’s internal caches, device memory, and external agents such as peer GPUs or CPUs. In [30], the synchronization elements, or flag variables used for inter-GPU coordination, are designated as non-cacheable memory locations. Treating these regions as non-cacheable ensures that updates are observed immediately across devices without requiring explicit invalidation or coherence traffic through the L1 or L2 caches. When a peer GPU issues an xGMI write to a polling GPU’s memory space, the request is directed through the xGMI fabric to the target GPU’s cache directory. The directory records the update atomically with respect to any pending polling reads from the local GPU, ensuring that subsequent accesses observe the most recent flag value. This behavior is facilitated by rocSHMEM, a GPU-resident communication runtime developed by AMD and AMD Research to provide GPU-centric networking through an OpenSHMEM-like interface [15]. rocSHMEM exposes one-sided communication primitives such as put, get, and atomic operations directly from GPU kernels, allowing threads to perform fine-grained data movement and synchronization without host intervention. It implements these primitives atop the xGMI interconnect, enabling intra-kernel communication and synchronization across GPUs at device speed. The library allocates a single symmetric heap across all participating GPUs, residing in each device’s memory space, which ensures a uniform address layout and simplifies pointer arithmetic in distributed kernels. By integrating communication into the GPU execution model, rocSHMEM reduces code complexity and allows for tighter coupling of computation and communication phases, where progress signaling and data exchange occur continuously within a single launch. This mechanism provides an efficient and coherent signaling pathway between GPUs. By avoiding round-trip synchronization through the CPU or host memory subsystem, xGMI writes enable fine-grained coordination at device speed. The write completes transparently at the cache directory level, simplifying coherence management in the polling GPU’s memory hierarchy while maintaining the atomicity and ordering guarantees required for synchronization. As a result, GPUs can exchange progress signals through low-latency peer-to-peer writes, supporting tightly coupled fused kernels that overlap computation and communication effectively.

While this mechanism provides programmer-transparent communication between multiple GPUs, it also introduces new architectural complexity that is difficult to analyze using existing tools. The precise timing of xGMI transactions, cache directory updates, and polling behaviors directly influence GPU utilization, synchronization efficiency, and the ability to hide communication latency. Small variations in these interactions can lead to measurable differences in overall throughput and scaling behavior, especially in tightly coupled workloads such as fused GEMV+AllReduce [30]. Capturing these nuances requires simulation frameworks that model both intra-GPU microarchitectural events and inter-GPU communication traffic at fine granularity. This tightly integrated approach to fused computation and synchronization highlights the importance of modeling both microarchitectural behavior and interconnect traffic in modern multi-GPU systems. However, cycle-level simulators have yet to fully support such fused communication patterns, motivating our extension of gem5 to capture these interactions in a realistic and configurable simulation environment.

2.3

Training Cycle on GPU Clusters

Training modern machine learning and artificial intelligence models at scale follows an iterative life-cycle that combines model design, large-scale deployment, and continuous performance tuning. A typical workflow begins with model construction and data preparation, followed by distributed training across multi-GPU or multi-node clusters. Each training iteration executes a forward pass to compute activations and a backward pass to propagate gradients, both of which are highly parallelized across many GPUs. Distributed data-parallel training frameworks such as PyTorch Distributed [22], Horovod [39], and DeepSpeed [33] manage the synchronization of gradients through collective communication operations, which aggregate partial results across devices. Large-scale systems such as NVIDIA DGX SuperPODs [3] or AMD Instinct™ MI300 clusters [4] rely on high-bandwidth, lowlatency interconnects to mitigate the inter-GPU communication as model sizes and datasets complexities grow. However, scaling inefficiencies persist due to factors such as kernel launch latency, load imbalance across devices, and variable communication delays introduced by the underlying network topology [43]. Even in seemingly homogeneous environments, contention within the interconnect fabric and subtle runtime differences among GPUs can lead to performance variability across nodes [37, 51]. These sources of non-determinism affect convergence time, reduce throughput, and complicate scheduling and performance prediction. Profiling is an essential step in addressing these challenges [6, 35]. During distributed training, runtime profiling tools such as NVIDIA Nsight™ Systems [5], ROCm Profiler [2], and PyTorch Profiler [1] are used to measure GPU utilization, communication latency, and kernel execution timing [19, 21, 25, 50]. These measurements help developers identify inefficiencies such as unbalanced workloads, excessive synchronization waits, or communication bottlenecks. The collected data also inform strategies for hyperparameter tuning, data partitioning, and kernel fusion, which can significantly improve hardware utilization and reduce total training time. Profiling further provides the empirical foundation for simulation-based 4

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

models are critical for evaluating architectural proposals that target inter-GPU communication, memory hierarchy design, and collective operation efficiency. As GPU clusters continue to scale toward exascale performance, profiling and modeling remain central to understanding and optimizing the interaction between computation and communication at every layer of the AI training stack.

ML Model Training and Infrastructure Optimization Cycle (1) Measurement and Profiling (4) Deploy Enhancements Across Compute Cluster Improvements were helpful

Portions of the Cycle Facilitated by Eidola

(2) Instrumentation and Perturbation on Characteristic Workloads Modifications

3

not beneficial

Eidola

This section describes the design and implementation of Eidola within the gem5 simulation framework. We begin in Section 3.1 by introducing the use of a setup kernel, which is launched prior to the kernel under study to prime simulator-side data structures with timing and communication information. We then discuss additional design considerations in Section 3.2, outlining alternative approaches we explored during development and highlighting key limitations that future users of gem5 should be aware of. Together, these sections detail both the core mechanisms of Eidola and the practical insights gained during its integration.

(3) Data Collection from Modifications

Figure 4: Iterative optimization cycle for large-scale GPU cluster training. The workflow consists of (1) measurement and profiling of baseline workloads, (2) instrumentation and controlled perturbation of representative kernels, (3) data collection and analysis of resulting performance characteristics, and (4) deployment of validated optimizations across the cluster. The stages highlighted in blue correspond to phases where Eidola facilitates analysis by enabling controlled replay and modeling of inter-GPU communication behavior.

3.1

Registering Writes via Setup Kernel

In the typical training life-cycle of large-scale GPU workloads, runtime profiling is an integral part of performance tuning and optimization. The timing profiles gathered during this process, or synthetically generated patterns, can be directly leveraged to inform simulation studies. Specifically, these profiles capture inter-GPU communication behavior, including peer-to-peer writes that occur during fused computation–communication kernels. By annotating each write operation with its corresponding timestamp, the collected traces provide detailed insight into when and how GPUs exchange data. For the fused GEMV+AllReduce kernel examined in this work, only the timestamps of peer-to-peer write operations are required to reconstruct realistic inter-GPU synchronization behavior. For other distributed workloads, users can choose to collect timing information with greater granularity, including additional communication events or memory operations, depending on the desired level of modeling fidelity. This flexibility allows the simulator to accommodate a broad range of use cases, from fast, approximate modeling to highly detailed, cycle-level analysis of multi-GPU communication dynamics. To inject these communication events into the simulation, we introduce the new GPU pseudo op, register_write, shown in Figure 5. This function was inserted into the application before the launch of the main computation kernel and recognized by the simulator as part of a setup phase. During simulation, this setup kernel executes in functional mode, not in detailed timing mode. The sole purpose of the setup kernel (or kernels if multiple writes to register) is to preload the simulator with information about upcoming inter-GPU writes that will be enacted during detailed simulation of the main kernel. The register_write function takes four parameters: addr, data, size, and wakeupTime. The addr parameter specifies the destination address for the emulated write, data holds the value to be written, and size indicates the width of the write in bytes (ranging from one to eight bytes). The wakeupTime parameter defines the time offset, in nanoseconds, after kernel launch at which the

performance modeling, allowing researchers to explore new architectural optimizations before hardware implementation. This can be in the form of tuning a synthetic traffic generator during simulation [8, 17, 27, 47], or, as done in this work, replaying real-life traffic patterns [7, 23]. Figure 4 illustrates this iterative optimization cycle. After an initial training deployment, developers first perform measurement and profiling to identify performance bottlenecks at both the kernel and system levels. These observations motivate targeted instrumentation of characteristic workloads, where specific communication or synchronization behaviors are isolated and, in some cases, deliberately perturbed to study their impact. The resulting execution traces are then collected and analyzed to understand how low-level effects—such as inter-GPU communication latency or synchronization delay—propagate to overall training performance. Once validated, optimizations derived from this process are deployed across the full compute cluster, completing the cycle. Eidola is designed to facilitate the first three stages of this workflow. By leveraging lightweight timing profiles and enabling controlled replay of inter-GPU communication events, Eidola allows researchers to systematically study the impact of communication timing and synchronization behavior without requiring repeated execution on large-scale hardware. In particular, the ability to perturb communication timing in a controlled, cycle-level simulation environment enables fine-grained analysis that complements traditional profiling tools. This makes Eidola a practical bridge between empirical measurement and architectural exploration, supporting rapid iteration on communication-centric optimizations before deployment at scale. By accurately capturing the timing and synchronization patterns observed in real multi-GPU workloads, profiling enables researchers to construct realistic models of distributed training behavior. These 5

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

1 __device__ void register_write (uint64_t addr, 2 uint64_t data, uint32_t size, uint64_t wakeupTime){ 3 asm volatile( 4 "\n\ v_mov_b32 v10, %0\n\ v_mov_b32 v11, %1\n\ 5 v_mov_b32 v12, %2\n\ v_mov_b32 v13, %3\n\ 6 v_mov_b32 v14, %4\n\ 7 gem5_register_write" 8 : 9 : "v"((addr>>32) & 0xFFFFFFFF), 10 "v"((addr>> 0) & 0xFFFFFFFF), 11 "v"(data), "v"(size),"v"(wakeupTime) 12 : "v10", "v11", "v12", "v13", "v14"); 13 }

direct xGMI writes rather than host-mediated messaging. By modeling these events at cycle granularity and maintaining an efficient WTT structure, the simulator achieves both temporal fidelity and scalability, enabling the study of tightly coupled multi-GPU workloads that overlap computation and communication within a single fused kernel execution. The flexibility of Eidola’s write scheduling methodology allows this framework to be useful for a variety of memory traffic generation schemes, including synthetic, [8, 17, 27, 47], as well as replicatory [7, 23].

3.2 Figure 5: Code snippet shown implementation details of proposed setup kernel.

Other Design Considerations

3.2.1 CPU Orchestration vs Setup Kernel. Before adopting the GPU pseudo op approach, we initially explored using CPU pseudo ops to emulate GPU writes. This method involved instrumenting the application with CPU threads that issued pseudo ops during the execution of the main GPU kernel. It was appealing due to its simplicity, minimal simulator modifications, and consistency with gem5’s original heterogeneous CPU-GPU simulation model. However, this approach proved impractical. In gem5, CPU and GPU components are serialized, not concurrent as in real hardware. Enabling CPU threads to issue writes while the GPU kernel runs would require frequent simulator context switches. These switches are configured based on host machine parameters such as core count, clock frequency, and microarchitecture, which are not part of the simulated system. This made the timing behavior unreliable and hard to reproduce for cycle-level studies. We advise against using CPU pseudo ops for emulating GPU behavior in gem5. While simpler, this method introduces nondeterminism in multi-GPU simulation contexts1 . Prior work examined gem5 simulation scalability in multi-CPU systems [38], but we avoided this implementation as naively partitioning GPUs between functional and timing models did not improve simulation time due to their larger, high-fidelity design.

write should be issued. At runtime, these parameters are read by the simulator and stored in a queue, sorted by wakeupTime. These timestamps are converted into cycles based on the device clock frequency defined in the gem5 configuration. We extend the simulator with a data structure to maintain all pending writes in a write tracking table (WTT), implemented as a priority queue sorted by wakeupTime. When a write is registered, it is inserted into the WTT at a position determined by its wakeupTime, ensuring that the earliest writes are processed first. During simulation of the main kernel, the head of the WTT is polled on every simulated cycle tick. If the current simulated time is less than the wakeupTime of the head entry, the simulator simply advances to the next cycle, incurring negligible runtime overhead (only a single comparison per cycle). When the simulated time meets or exceeds the wakeupTime, all write entries at the head of the queue with the same timestamp are popped and enacted as xGMI writes. These events are processed in constant time, as the number of simultaneous writes is typically small and bounded by the number of participating GPUs. This priority queue-based design ensures both temporal accuracy and efficiency, allowing users to register writes in arbitrary order. The means sequential calls to register_write() need not correspond to the chronological order of their execution. By decoupling registration from enactment, this mechanism allows flexible modeling of inter-GPU communication patterns with minimal simulation overhead. As mentioned above, when the main kernel enters detailed timing mode, the simulator begins polling the WTT on every simulated cycle. Once the current simulation time reaches the wakeupTime associated with the next pending event, the simulator issues the corresponding xGMI write on behalf of the emulated GPU. The write transaction completes at the cache directory level, accurately reflecting the hardware behavior described in Section 2.2, where peer-to-peer writes update non-cacheable memory regions atomically and without round-trip acknowledgment. Upon completion, the memory state of the receiving GPU is updated to reflect the new flag value. When the receiving GPU subsequently polls this address the updated data is detected and propagated to the requesting CU. This implementation provides a faithful simulation of the lowlatency, one-sided communication semantics employed in rocSHMEM [15], where intra-kernel synchronization primitives rely on

3.2.2 Event Queues vs Write Tracking Table. Earlier in Section 3.1, we described our modifications to the gem5 simulator that introduced a new GPU pseudo-op to interact with the custom WTT data structure. The WTT provides an intuitive interface for scheduling and emulating timed inter-GPU writes, implemented as a priority queue whose head is polled each simulation cycle. While this design offers transparency and ease of debugging (all pending and enacted writes can be directly inspected during runtime), it does incur a polling overhead due to the per-cycle check on the queue head. This cost is minute, since the poll reduces to a constant-time comparison between the current simulation tick and the wakeup time of the head entry. In the common case, the scheduled write lies in the future, so the simulator performs only this O(1) check without additional processing. In the native gem5 simulation framework, timing-sensitive events such as memory transactions, DMA operations, and synchronization events are typically scheduled using event queues. These queues allow the simulator to efficiently manage future events without active polling, as each event is inserted with an associated 1 gem5’s GPU model typically executes on a KVM-backed CPU, the CPU-side portions of

the application simulation depend on the host machine’s hardware characteristics (e.g. core count, clock behavior, scheduling policy, etc.). As a result, GPU–CPU interaction can exhibit nondeterministic timing and context-switch behavior that does not reflect the modeled system. 6

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

timestamp and automatically triggered when its scheduled time arrives. Leveraging this existing event-driven infrastructure would eliminate the need for explicit polling of the WTT, thereby improving runtime efficiency and scalability for large-scale multi-GPU simulations. Integrating the WTT functionality into gem5’s event queue subsystem represents a promising direction for future work. This approach would maintain the same temporal fidelity while reducing the simulator’s runtime overhead, particularly in workloads that register thousands of inter-GPU writes per kernel. Moreover, using the native event queues would align the GPU pseudo-op mechanism with gem5’s established scheduling paradigm, simplifying maintenance and improving interoperability with other timing events in the simulation pipeline. While the current WTT-based implementation prioritizes debuggability and transparency of simulator state, future releases may incorporate an event-driven backend to further enhance performance and reproducibility.

4

This experiment validates that our instrumentation and simulator extensions provide a transparent and flexible way to explore the performance impact of inter-GPU communication in distributed workloads. Table 1 summarizes the simulator configuration and application parameters used in Figures 6 and 9.

5

To demonstrate the utility and flexibility of our simulation methodology, we implement the key features of a recent GPU microarchitecture [16]. This work introduced the synchronization monitor (SyncMon), which enables fine-grained synchronization support in GPU hardware. The SyncMon mechanism, along with associated hardware logic, allows GPU threads to use monitor/wait semantics similar to the monitor/mwait instruction pair in CPUs. Instead of continuously polling in a tight spin-wait loop, threads can now yield GPU compute resources while waiting, adopting a spin-yield synchronization pattern. We extend our simulator to incorporate the core concepts of SyncMon, namely the Monitor Log, which we implement as a simulator-side data structure rather than allocating it within the device’s memory space. This design decision reflects the purpose of this case study: not to propose new architectural mechanisms or evaluate hardware feasibility, but to demonstrate the flexibility and extensibility of our simulation infrastructure. The gem5 framework serves primarily as a research exploration tool, and in that spirit, our goal is to show that with the proposed enhancements to gem5, complex microarchitectural concepts such as SyncMon can be replicated, explored, and scaled across multi-GPU systems with minimal effort. By modeling the Monitor Log within the simulator state, we can easily control and observe its parameters, validate synchronization behavior, and characterize the effects of scaling out to larger GPU networks. The intent is not to reproduce a facsimile of the full SyncMon microarchitecture at hardware fidelity, but rather to illustrate that key ideas in GPU interconnect and synchronization research can be rapidly prototyped within our framework. Through this implementation, GPU threads can register monitored memory locations and suspend execution until a peer GPU write is detected at that address. Once a write is observed, modeled as an xGMI write as described in Section 3.1, the simulator triggers the corresponding monitor entry, resuming the waiting wavefronts. This mechanism reduces unnecessary read traffic and more accurately captures the performance characteristics of spinyield synchronization in [30]. To simulate this behavior, we introduce two new GPU pseudo operations, monitor() and mwait(), which extend the gem5 instruction interface in a manner similar to the register_write() function described earlier in Section 3.1. Together, these pseudo ops emulate the semantics of the x86 monitor and mwait instructions, allowing a GPU thread to register interest in a specific memory location and suspend execution until an update is detected. Our monitor() GPU pseudo op takes three arguments: addr, numBytes, and wakeValue. The addr argument specifies the starting address of the memory region to be monitored, while numBytes indicates the size of the flag or synchronization variable. This size flexibility accommodates padded flags used to prevent false sharing. The wakeValue argument defines the expected value of the

Results

Our simulation framework enables full instrumentation of interGPU network traffic patterns through annotated, timestamped write instructions. This design provides researchers with precise control over when communication events occur relative to GPU kernel execution. To demonstrate this capability, we evaluate how varying the wakeupTime parameter of our registered write impacts observed memory traffic in simulation. Figure 6 illustrates this behavior by sweeping the wake-up time of the emulated write from 0 to 40 microseconds. A wake-up time of 0 indicates the emulated write is issued at the start of the main kernel execution, while higher values delay the emulated write correspondingly. The vertical axis shows the total number of memory read requests observed during simulation. Each bar is divided into two categories: non-flag reads (shown in blue), which represent general memory traffic, and flag reads (shown in red), which result from spin-wait loops polling for the flag update. These polling reads correspond to the flag-waiting behavior on lines 14-15 of the fused kernel code shown in Figure 3. Table 1: Simulator and Application Configurations Simulation Configurations Number of CUs in simulated GPUs Number of emulated GPUs Workgroups launched per GPU Application Configurations Input Matrix dimension M Input Matrix dimension K Input Matrix dimension N

Case Study: SyncMon

4 3 208 256 8192 1

As expected, the number of flag reads increases linearly with the delay introduced by wakeupTime. This confirms that the simulator faithfully reproduces fine-grained timing behavior consistent with GPU spin-wait synchronization. The clear separation between non-flag and flag traffic also highlights the utility of our method in isolating and analyzing communication-driven memory access patterns. 7

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

Memory Traffic Scales with Output Tile Computation Time

800

Non-Flag

700 500 400 300 200 13422 26956

100

40489 53975

67470 80871

94545 108034

121518

162090 135050 148558

175615 189039

202539 216026

229645 243165

283715 256669 270197

297220 310732

324255 337768

351277

378326

418829 391826 405234

445835 459385

472896 486409

499864 513436

40

39

38

37

36

35

34

33

31

30

29

28

26

Modelled Output Tile Computation Time ( sec)

25

24

23

22

21

20

19

18

17

16

15

14

13

12

11

9

10

8

7

6

5

4

1

2

65940 65970 66112 66061 66182 66366 66169 66289 65943 66006 66191 66283 66049 65975 66008 66059 66176 66138 66032 66103 66103 66045 66097 65943 66183 65948 65926 65922 66064 66352 66028 66084 66172 65920 66023 65985 66170 66125

0

3

Kilo Reads (#)

600

Flag 526829 540460

Figure 6: Fine-grained control of simulated multi-GPU communication traffic. Horizontal axis shows sweep of wakeupTime parameter for registered writes in setup kernel. Vertical axis shows number of read requests issued by the fused GEMV+AllReduce kernel application. Non-flag related read requests are shown in blue (patterned with forward slashes), while those issued from the spin-wait are shown in orange (patterned with back slashes). Monitor, Mwait, and xGMI Interactions with Monitor Log

1

monitor(0x3FD005008, 8, 0x1)

it demonstrates that complex synchronization schemes can be prototyped and evaluated in gem5 with minimal instrumentation and simulator overhead, providing researchers with a flexible environment to explore new ideas. 1 initializes an entry in the simulatorThe monitor() operation ○ side Monitor Log, shown in Figure 7, with the address of the monitored variable, the size of the monitored region (stored in the monitor mask), and the expected value that will trigger a wake-up. When the instrumented application issues an mwait() operation 2 the wavefront ID of the calling thread is for the same address ○, recorded in the corresponding Monitor Log entry, and the wavefront is descheduled by the GPU scheduler. Multiple wavefronts may register to the same table entry when different synchronization flags share a cache line, as the instrumentation does not impose any ordering or exclusivity among flags. Execution resumes only when an emulated xGMI write to the 3 Upon such a write, the simulator persame cache line occurs ○. forms a masked comparison between the written data and the stored wake-up value. If a match is detected, all waiting wavefronts are marked as schedulable. Figure 8 presents a timeline of the instrumented fused GEMV+AllReduce kernel. The simulation begins in fast functional mode to initialize the gem5 structures and any setup kernels, and then transitions to detailed timing mode to collect fine-grained synchronization and communication statistics. Eidola leverages common profiling timing data from the standard training lifecycle of distributed machine learning workloads. Lightweight timing profiles are collected from real executions to capture inter-GPU communication behavior. For the fused GEMV+AllReduce kernel and the other kernels from [30], peer-to-peer GPU write timestamps fully encapsulate the network traffic. These timestamps are registered into the simulator before kernel execution using a GPU pseudo-op that schedules future emulated inter-GPU writes. During gem5’s execution-driven simulation of a single target GPU, executing in detailed timing mode, all other GPUs are abstracted as lightweight eidolons (eGPUs) that replay recorded communication events. The framework allows researchers to explore different wake-up granularities: at a coarse level, all waiting wavefronts can be woken following Mesa-style synchronization semantics (mwait must be

2

mwait(0x3FD005008)

Monitor Log Line Address Compare Value Monitor Mask Waiting WFs 0x3FD004FC0 0x0000...000100000000000000000001 0x00000000FFFFFFFF [31, 746] 0x3FD005000 0x0000...000100000000000000000000 0x00000000FFFF0000 [912] 0x3FD005040 0x0000...000000000000000000000001 0x000000000000FFFF [351] 0x3FD005080 0x0000...000100000000000000000001 0xFFFFFFFF00000000 [267, 431, 638, 855]

3

Complete write, read latest value

xGMI(0x3FD005008, 8, 0x1)

Masked =

Do nothing, WFs continue waiting

No

=?

Yes

Wake all waiting WFs

Figure 7: The Monitor Log and how the monitor(), mwait(), and emulated xGMI write events interact with this structure to trigger thread wakeups.

monitored flag that will trigger a wake-up event, allowing the same mechanism to represent a wide range of synchronization primitives such as mutexes, semaphores, and barriers. This design extends beyond the x86 monitor instruction, which simply observes writes within a fixed 8-byte range from a register-specified address. By allowing for explicit specification of monitored size and wake conditions, the GPU monitor() pseudo op provides a richer and more controllable interface for exploring synchronization mechanisms. The mwait() pseudo op complements monitor() by suspending execution of a calling wavefront until one of its monitored memory addresses meets a wakeup condition. It takes a single argument, addr, the address being waited on, and records the calling wavefront’s ID in the corresponding Monitor Log entry. Once issued, the wavefront is marked to be descheduled by the GPU scheduler, freeing compute resources for other active wavefronts. During simulation, each memory write that completes at the cache directory is compared against the entries in the Monitor Log. If a matching address is found and its associated wakeValue condition is satisfied, all wavefronts waiting on that entry are marked as ready and rescheduled by the command processor in the subsequent dispatch cycle. This design allows the simulator to model spin-yield behavior efficiently and deterministically, capturing the intended functionality of the SyncMon hardware described in [16]. More importantly, 8

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

gem5 Simulation Timeline of GEMV+AllReduce Fused GPU Kernel Setup Kernel(s) To Emulate Peer GPU Writes

Fused GEMV+AllReduce Kernel Context

1

Arm Monitor Table with Flag Addresses

WFs Scheduled, Fused Kernel Execution Cont'd

Fused Kernel Execution Cont'd 2 Mwait Deschedules Spinning WFs

WTT Armed with Pending Writes

Default gem5 Initialization Initialize Write Tracking Table Initialize Monitor Table

gem5 Simulation Context

Fused Kernel Execution Start

3 WTT issues an xGMI emualted write

gem5 Sim Continues

Functional Execution

Fused Kernel Completes gem5 Teardown

Detailed Timing Execution

Reduced Memory Traffic from Spin-Yield

65948

65914

66109

65927

65972

66222

32

33

35

37

40

65919

30

31

66120

66004

29

66011

27

28

66000

65919

26

65923

24

Modelled Output Tile Computation Time ( sec)

25

65949

65928

23

65922

21

22

66126

65926

65911

18

20

65909

17

19

66088

65963

66026

14

16

65923

13

15

65920

65917

12

65920

10

11

65933

65919

9

66400

7

8

66119

65960

6

66037

4

5

66119

65927

3

65909

Non-Flag Flag 773 755 754 728 747 752 766 740 747 777 746 748 744 768 781 745 761 727 774 760 737 761 706 777 768 770 779 726 737 768 761 740 760 758 782 788

2

80 75 70 65 60 55 50 45 40

1

Kilo Reads (#)

Figure 8: A timeline of the simulated execution of the fused GEMV+AllReduce kernel, annotated to show the instrumentation points used to evaluate SyncMon-inspired synchronization features within the proposed simulation framework.

Figure 9: Read request reduction when adopting a spin-yield pattern synchronization pattern. Axes and color format is consistent with Figure 6. called within the while loop from lines 15–16 in Figure 3, while finer-grained tracking of individual flags can be implemented to emulate Hoare-style semantics. Again, this case study is not intended to reproduce SyncMon [16] at hardware fidelity, but rather to demonstrate how our enhanced simulation framework enables rapid, scalable exploration of advanced GPU synchronization mechanisms in multi-GPU systems.

5.1

the benefit of the proposed mechanism. These results validate the expected behavior described in the original SyncMon proposal and demonstrate that Eidola can accurately capture the performance implications of advanced synchronization mechanisms.

Key Takeaway: Eidola enables detailed evaluation of communication and synchronization mechanisms.

Enabling Architectural Research

We evaluate this extension in our multi-GPU simulation environment under the same test configuration described in 4. As shown in Figure 9, once SyncMon support is enabled, the number of memory reads issued by the simulated GPU no longer scales linearly with the modeled wait time. This contrasts sharply with the baseline spin-wait behavior shown in Figure 6, where increased wait times led to a proportional increase in polling reads due to continuous flag checking. With SyncMon-enabled execution, the number of reads associated with polling on flag variables remains effectively constant across all configurations, ranging between 728 and 788 reads. This is a significant departure from the baseline, where polling traffic grows with the duration of the wait. The bounded number of flagrelated reads reflects the transition from an active spin-wait to a spin-yield synchronization model, where waiting wavefronts are descheduled and only resume execution upon receiving a notification event. Non-flag-related memory reads remain largely unchanged, staying consistent with the values observed in Figure 6 at approximately 66K reads. This indicates that the introduction of SyncMon does not perturb the underlying computational behavior or memory access patterns of the kernel. Instead, it selectively reduces only the unnecessary memory traffic generated by synchronization, isolating

5.2

Simulation Scaling with Input Size

To evaluate how Eidola behaves under increasing computational workload, we study the relationship between application input size and overall simulation time. In particular, we vary the input matrix dimension M, which corresponds to the number of output rows computed by the fused GEMV+AllReduce kernel. Increasing M proportionally increases the amount of arithmetic work performed by each GPU, as more multiply-accumulate operations must be completed per kernel invocation. This experiment allows us to verify that Eidola preserves expected scaling trends as application complexity grows. Figure 10 compares total gem5 simulation time as the input dimension M is varied. Across all configurations, simulation time increases approximately linearly with M, as indicated by 𝑟 2 values ranging from 0.76 to 0.98 for the fitted trendlines. This behavior is consistent with expectations: as M grows, each GPU executes a larger number of compute operations, leading to a proportional increase in detailed timing simulation. Importantly, both the baseline configuration (solid lines) and the mwait-instrumented configuration (dashed lines) exhibit similar scaling trends. This indicates that the additional synchronization modeling introduced by Eidola, including emulated xGMI writes 9

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

24H

Simulation Time

20H 16H 12H 8H 4H 0

5000

10000

15000 20000 Input Dimension M (#)

25000

30000

3 eGPUs r2=0.76 3 eGPUs r2=0.96 (mwait) 7 eGPUs r2=0.91 7 eGPUs r2=0.96 (mwait) 15 eGPUs r2=0.88 15 eGPUs r2=0.98 (mwait) 31 eGPUs r2=0.87 31 eGPUs r2=0.95 (mwait) 63 eGPUs r2=0.84 63 eGPUs r2=0.92 (mwait) 127 eGPUs r2=0.77 127 eGPUs r2=0.85 (mwait) 255 eGPUs r2=0.79 255 eGPUs r2=0.81 (mwait)

where 𝑡 1𝐺𝑃𝑈𝑀 represents the estimated time to simulate a single GPU in detailed timing mode, and 𝑡𝑒𝐺𝑃𝑈𝑀 captures the per-eGPU overhead associated with maintaining additional communication state and replaying inter-GPU events. Here, 𝑡𝑀 is the measured simulation time for a given configuration, and 𝑒𝐺𝑃𝑈 𝑠 is the number of emulated GPUs. Simulation Time Scaling with eGPUs

Figure 10: Plots showing gem5 simulation time of the GEMV+AllReduce kernel with varying input matrix dimension M. Vertical axis shows wall clock time of gem5 simulations. Horizontal axis represents number of rows simulated in input matrix. Solid lines show results from default application with emulated xGMI writes. Dashed lines show results from mwait-instrumented application.

10

30

8

25 20

6

15

4

10

2

5 50

100 150 200 Number of eGPUs (#)

250

0

0

50

100 150 200 Number of eGPUs (#)

250

M=1024 M=1280 M=1536 M=1792 M=2048 M=3072 M=3584 M=4096 M=5120 M=6144 M=7168 M=8192 M=10240 M=12288 M=14336 M=16384 M=20480 M=24576 M=28672 M=32768

Figure 11: gem5 simulation time as number of eGPUs varies. Vertical axis represents relative simulation time normalized to estimation of simulating single GPU. Horizontal axis varies the number of emulated GPUs. Figure 11a) shows results from the application with xGMI emulated write instrumentation alone, while Figure 11b) shows results including mwait-instrumentation. Because simulation time is also a function of input size, all parameters are derived per input dimension M, as discussed in Section 5.2. After fitting the linear model in Equation 1, we normalize the measured simulation time by the estimated 𝑡 1𝐺𝑃𝑈𝑀 and plot the results in Figure 11. This normalization enables direct comparison against an idealized baseline in which only a single GPU is simulated in detail. Figures 11a) and 11b) show that simulation time grows sublinearly with the number of eGPUs across all configurations. For small numbers of eGPUs, normalized execution time approaches 1.0, indicating that the simulation cost is dominated by the detailed modeling of the primary GPU. As the number of eGPUs increases, total simulation time increases, but at a much slower rate than linear scaling. Even in the largest configuration with 255 eGPUs and mwait instrumentation enabled, normalized execution time ranges between 7.3× and 35.9×, far below the 256× cost that would be expected if all GPUs were simulated in full detail. This behavior confirms that the overhead introduced by Eidola is modest relative to the cost of detailed GPU simulation. The estimated 𝑡𝑒𝐺𝑃𝑈𝑀 values indicate that the additional work required to maintain the Write Tracking Table (WTT) and replay communication events is small and scales sub-linearly with the number of eGPUs. While an event-driven implementation using gem5’s native event queues could further reduce this overhead by eliminating per-cycle polling, the current design already provides a favorable trade-off between efficiency, transparency, and ease of debugging. Overall, these results demonstrate that Eidola enables scalable multi-GPU simulation by avoiding the need to fully model every GPU in the system, while still capturing the essential communication behavior required for architectural analysis.

Key Takeaway: Eidola preserves expected applicationlevel scaling.

Simulation Scaling with Emulated GPUs

The fused GEMV+AllReduce kernel inherently requires multiple GPUs to execute, as its computation depends on distributed data exchange and synchronization across devices. In practice, the kernel requires at least two GPUs to run and achieves its highest efficiency when executed on a power-of-two number of GPUs. As a result, directly measuring the simulation time of a single GPU execution (𝑡 1𝐺𝑃𝑈𝑀 ) is not possible for this workload. This presents a challenge when attempting to quantify the overhead introduced by Eidola relative to a baseline single-GPU detailed simulation. To address this, we vary the number of emulated GPUs (eGPUs) from 3 to 255 and measure the total wall-clock simulation time across configurations. We then use these measurements to extrapolate the effective cost of simulating a single GPU and the incremental overhead associated with each additional eGPU. Specifically, we group trials by input dimension 𝑀, fit a linear model across configurations, and extrapolate to 𝑒𝐺𝑃𝑈 = 1 to estimate both 𝑡 1𝐺𝑃𝑈𝑀 and 𝑡𝑒𝐺𝑃𝑈𝑀 . This approach allows us to isolate the contribution of Eidola’s communication modeling infrastructure to overall simulation time, even in the absence of a directly measurable single-GPU baseline. Formally, we model total simulation time 𝑡𝑀 as: 𝑡𝑀 = 𝑡 1𝐺𝑃𝑈𝑀 + 𝑒𝐺𝑃𝑈 𝑠 ∗ 𝑡𝑒𝐺𝑃𝑈𝑀

b) Mwait Enabled

35

0

and optional mwait-based behavior, does not fundamentally alter the relationship between workload size and simulation cost. Instead, the dominant factor remains the execution time of the simulated GPU pipeline. These results confirm that Eidola preserves the expected computational scaling characteristics of the underlying application while introducing only modest overhead for communication modeling.

5.3

a) Default Application

12 Simulation Time Normalized to Simulating 1 GPU

Simulation Time Scaling with Input Size

(1) 10

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

the behavior of non-target GPUs while preserving detailed simulation for selected devices, enabling scalable modeling of inter-GPU communication traffic and synchronization effects in large multiGPU systems. Eidola and MGPUSim address different points in the multi-GPU simulation design space. MGPUSim provides detailed, event-driven simulation of all GPUs, which is valuable for full-system studies but limits scalability as GPU count increases. Eidola instead abstracts non-target GPUs and applies detailed timing simulation only to a selected device, enabling scalable exploration of inter-GPU communication and synchronization effects on GPU microarchitecture. This is particularly effective for computationally balanced workloads such as those in [30], where examining a single GPU can provide system-level insights. While both simulators rely on event-based mechanisms, Eidola’s key contribution is decoupling execution detail from communication fidelity. By replaying cycle-level communication events from abstracted GPUs, Eidola supports simulations with hundreds of GPUs while retaining fine-grained visibility into a target GPU. Finally, MGPUSim models AMD GPUs based on GCN3, whereas Eidola targets modern CDNA-based GPUs (MI100, MI200, MI300), representative of current hyperscale systems.

Key Takeaway: Eidola introduces modest simulation overheads as system size increases.

5.4

Case Study Summary

The SyncMon study exemplifies how our proposed methodology can be directly applied to support GPU architectural research. Across the preceding experiments, we demonstrate that Eidola enables both accurate modeling of communication-driven behavior and scalable simulation of large multi-GPU systems. In Section 5.1, we validated Eidola’s ability to model architectural mechanisms through the SyncMon case study. By replacing spin-wait polling with a notification-based synchronization primitive, we observed that flag-related memory reads remain bounded (728–788 reads) rather than scaling with wait time, while non-flag memory traffic remains stable at approximately 66K reads. This demonstrates that Eidola can isolate and quantify the impact of synchronization mechanisms on memory behavior without perturbing the underlying computation. In Section 5.2, we showed that simulation time scales linearly with input size, confirming that Eidola preserves the expected relationship between application workload and execution cost. This ensures that increasing computational intensity does not introduce non-physical artifacts in simulation, allowing researchers to confidently study workload-dependent trends. Finally, in Section 5.3, we further demonstrated that simulation time grows sub-linearly with the number of emulated GPUs. By extrapolating a single-GPU baseline and isolating the per-eGPU overhead, we showed that Eidola avoids the prohibitive cost of fully simulating all devices while still capturing inter-GPU communication behavior. Even at large scales, the additional overhead introduced by communication replay and Write Tracking Table (WTT) management remains modest relative to detailed GPU pipeline simulation. This scalability is critical for studying modern distributed workloads, which routinely span tens to hundreds of GPUs. Taken together, these results highlight Eidola’s effectiveness as a research platform. By providing cycle-level, configurable modeling of multi-GPU systems, Eidola opens the door to exploring emerging hardware features in realistic, large-scale environments. As multi-GPU systems continue to underpin large-scale model training, such capabilities will be increasingly important for guiding both academic research and industrial design.

6

7

Future Work

An important direction for future work is extending validation to larger, scale-out multi-GPU systems. While the timing profiles used in this study were collected from a single node containing four GPUs, our simulator modifications are designed to support configurations with tens to hundreds of GPUs. Evaluating the methodology across more complex intra-node and inter-node topologies would provide greater insight into how communication patterns scale across distributed environments. This extended validation would further establish the simulator as a reliable platform for studying emerging architectural solutions in high-performance, distributed GPU systems. The fused GEMV+AllReduce kernel examined in this study is a balanced workload. We selected it because it stresses fine-grained synchronization, which is central to modern transformer training. The same work [30] also presents other fused kernels, such as embedding pooling + All-to-All and GEMM + All-to-All, which can be evaluated using Eidola without modification. Importantly, Eidola imposes no constraints on workload balance. User-specified profiling enables support of arbitrary and asymmetric communication patterns, ideal for workloads exhibiting producer–consumer data patterns. It is our hope that Eidola will enable research spanning a broad diversity of GPU workloads that exploit both heterogeneity and asymmetry.

Related Work

Several simulation frameworks have explored GPU and heterogeneous system modeling, but they target different aspects than our work. Multi2Sim [45] focuses on CPU–GPU interaction and host-driven communication, and does not model peer-to-peer GPU communication that dominates modern distributed training workloads. GPGPU-Sim [9] is a detailed simulator designed to study single-GPU microarchitectural behavior, but it lacks native support for multi-GPU systems and inter-GPU synchronization. MGPUSim [42] provides a detailed multi-GPU simulation environment and is the closest to our work; however, it models all GPUs in a cluster at fine granularity. In contrast, Eidola abstracts

8

Conclusion

The growing demand for large-scale machine learning has made multi-GPU systems central to modern training infrastructure, yet understanding the performance impact of inter-GPU communication remains challenging. This work extends gem5 to provide configurable, timing-accurate modeling of multi-GPU systems, Eidola, using real application traces annotated with precise timestamps to emulate peer-to-peer GPU writes. 11

Ranganath Selagamsetty, Matthew Poremba, Bradford Beckmann, Joshua San Miguel, and Mikko Lipasti

Volume 1 (Lake Tahoe, Nevada) (NIPS’12). Curran Associates Inc., Red Hook, NY, USA, 1223–1231. [15] Advanced Micro Devices. [n. d.]. ROCm OpenSHMEM (rocSHMEM). Accessed November 10, 2025. [16] Alexandru Duţu, Matthew D. Sinclair, Bradford M. Beckmann, David A. Wood, and Marcus Chow. 2020. Independent Forward Progress of Work-groups. In 2020 ACM/IEEE 47th Annual International Symposium on Computer Architecture (ISCA). 1022–1035. https://doi.org/10.1109/ISCA45697.2020.00087 [17] L. Eeckhout, K. de Bosschere, and H. Neefs. 2000. Performance analysis through synthetic trace generation. In 2000 IEEE International Symposium on Performance Analysis of Systems and Software. ISPASS (Cat. No.00EX422). 1–6. https://doi.org/ 10.1109/ISPASS.2000.842273 [18] Joaquín Ferrer, Juan M. Cebrian, Ricardo Fernández-Pascual, and Manuel E. Acacio. 2025. Precise characterization of coherence activity in multicores using gem5. In The Journal of Supercomputing, Vol. 81. https://doi.org/10.1007/s11227025-07434-0 [19] Arnaud Fiorini and Michel R. Dagenais. 2022. Visualization of profiling and tracing in CPU-GPU programs. Concurrency and Computation: Practice and Experience 34, 23 (2022), e7188. https://doi.org/10.1002/cpe.7188 arXiv:https://onlinelibrary.wiley.com/doi/pdf/10.1002/cpe.7188 [20] Hanchen Jin, Zichao Yue, Zhongyuan Zhao, Yixiao Du, Chenhui Deng, Nitish Srivastava, and Zhiru Zhang. 2025. Vesper: A Versatile Sparse Linear Algebra Accelerator With Configurable Compute Patterns. IEEE Transactions on Computer-Aided Design of Integrated Circuits and Systems 44, 5 (2025), 1731–1744. https://doi.org/10.1109/TCAD.2024.3496882 [21] Matthew Leinhauser, Jeffrey Young, Sergei Bastrakov, Rene Widera, Ronnie Chatterjee, and Sunita Chandrasekaran. 2021. Performance Analysis of PIConGPU: Particle-in-Cell on GPUs using NVIDIA’s NSight Systems and NSight Compute. Technical Report. Oak Ridge National Laboratory (ORNL), Oak Ridge, TN (United States). https://doi.org/10.2172/1761619 [22] Shen Li, Yanli Zhao, Rohan Varma, Omkar Salpekar, Pieter Noordhuis, Teng Li, Adam Paszke, Jeff Smith, Brian Vaughan, Pritam Damania, and Soumith Chintala. 2020. PyTorch Distributed: Experiences on Accelerating Data Parallel Training. arXiv:2006.15704 [cs.DC] https://arxiv.org/abs/2006.15704 [23] Wenhai Lin, Yiquan Chen, Jiexiong Xu, Zhen Jin, Peiyu Liu, Shishun Cai, Yuzhong Zhang, Jingchang Qin, Yiquan Lin, and Wenzhi Chen. 2024. CINDA: Don’t Ignore Instructions When Cloning Memory Access Behavior. In 2024 IEEE 24th International Symposium on Cluster, Cloud and Internet Computing (CCGrid). 507–513. https://doi.org/10.1109/CCGrid59990.2024.00063 [24] Qunyou Liu, Marina Zapater, and David Atienza. 2025. Gem5-AcceSys: Enabling System-Level Exploration of Standard Interconnects for Novel Accelerators. In Proceedings of the 62nd Annual ACM/IEEE Design Automation Conference (San Francisco, California, United States) (DAC ’25). IEEE Press, Article 444, 7 pages. https://doi.org/10.1109/DAC63849.2025.11133394 [25] Ethan Luis McDonough, Joel E. Denny, and Johannes Doerfert. 2026. Profile Generation for GPU Targets. In OpenMP: Balancing Productivity and Performance Portability, Yonghong Yan, Michael Klemm, Bronis R. de Supinski, Erik Saule, Jannis Klinkenberg, and Swaroop Pophale (Eds.). Springer Nature Switzerland, Cham, 99–113. [26] Aditya Nandakumar, Shobha Vissapragada, Ashish Panday, Wen Xie, Matt Ouellette, Zhenyu Gu, and Ram Sivaramakrishnan. 2025. WHITEPAPER: AT-SCALE AI TRAINING ON AMD INSTINCT™ MI350/MI300X SERIES GPUS. Technical Report. Advanced Micro Devices, Inc. [27] Reena Panda, Xinnian Zheng, Jiajun Wang, Andreas Gerstlauer, and Lizy K. John. 2017. Statistical Pattern Based Modeling of GPU Memory Access Streams. In Proceedings of the 54th Annual Design Automation Conference 2017 (Austin, TX, USA) (DAC ’17). Association for Computing Machinery, New York, NY, USA, Article 81, 6 pages. https://doi.org/10.1145/3061639.3062320 [28] Josna Philomina, Rekha K. James, Shirshendu Das, Palash Das, and Daleesha M Viswanathan. 2026. NeSTAR: Hardware Trojans and its mitigation strategy in NoC routers. Integration 107 (2026), 102603. https://doi.org/10.1016/j.vlsi.2025. 102603 [29] Jason Power, Joel Hestness, Marc S. Orr, Mark D. Hill, and David A. Wood. 2015. gem5-gpu: A Heterogeneous CPU-GPU Simulator. IEEE Computer Architecture Letters 14, 1 (2015), 34–36. https://doi.org/10.1109/LCA.2014.2299539 [30] Kishore Punniyamurthy, Khaled Hamidouche, and Bradford M. Beckmann. 2024. Optimizing Distributed ML Communication with Fused Computation-Collective Operations. In SC24: International Conference for High Performance Computing, Networking, Storage and Analysis (Atlanta, GA, USA). IEEE Press, 1–17. https: //doi.org/10.1109/SC41406.2024.00094 [31] Vishnu Ramadas, Daniel Kouchekinia, Ndubuisi Osuji, and Matthew D. Sinclair. 2023. Closing the Gap: Improving the Accuracy of gem5’s GPU Models. (2023). https://par.nsf.gov/biblio/10468163 [32] Vishnu Ramadas, Daniel Kouchekinia, and Matthew D Sinclair. 2024. Further Closing the GAP: Improving the Accuracy of gem5’s GPU Models. (2024). https: //par.nsf.gov/biblio/10542852 [33] Jeff Rasley, Samyam Rajbhandari, Olatunji Ruwase, and Yuxiong He. 2020. DeepSpeed: System Optimizations Enable Training Deep Learning Models with Over

Eidola reproduces communication-induced memory traffic with fine-grained control, enabling detailed analysis of synchronization delays. We validate this capability by demonstrating linear scaling of spin-wait traffic with write delay and by implementing key features of the SyncMon proposal, where spin-yield synchronization reduces memory reads as expected. Eidola represents an initial step toward scalable simulation of multi-GPU workloads with finegrained synchronization. The case study in Section 5 demonstrates both its scalability and usefulness. These results show that our approach supports realistic modeling of inter-GPU interactions and facilitates architectural exploration of emerging synchronization mechanisms. As distributed training continues to scale, such tools will be essential for bridging devicelevel design and system-level performance modeling.

References [1] 2021. Introducing PyTorch Profiler – the new and improved performance tool. Retrieved November 6, 2025 from https://pytorch.org/blog/introducing-pytorchprofiler-the-new-and-improved-performance-tool/ [2] 2022. ROC-profiler and debugger: An Overview of AMD ROCmTM Tools. Retrieved November 6, 2025 from https://extremecomputingtraining.anl.gov/wp-content/ uploads/sites/96/2022/11/ATPESC-2022-Track-6-Talk-4-Tandon-AMD.pdf [3] 2023. NVIDIA DGX SuperPOD: Next Generation Scalable Infrastructure for AI Leadership, Reference Architecture. Retrieved November 6, 2025 from https://docs.nvidia.com/https:/docs.nvidia.com/dgx-superpod-referencearchitecture-dgx-h100.pdf [4] 2025. AMD Instinct™ MI300 Series Cluster Reference Architecture Guide. Retrieved November 6, 2025 from https://www.amd.com/content/dam/amd/en/documents/ instinct-tech-docs/other/instinct-mi300-series-cluster-reference-guide.pdf [5] 2025. NVIDIA Nsight Systems. Retrieved November 6, 2025 from https://developer. nvidia.com/nsight-systems [6] Jennifer M. Anderson, Lance M. Berc, Jeffrey Dean, Sanjay Ghemawat, Monika R. Henzinger, Shun-Tak A. Leung, Richard L. Sites, Mark T. Vandevoorde, Carl A. Waldspurger, and William E. Weihl. 1997. Continuous profiling: where have all the cycles gone? ACM Trans. Comput. Syst. 15, 4 (Nov. 1997), 357–390. https: //doi.org/10.1145/265924.265925 [7] Mario Badr, Carlo Delconte, Isak Edo, Radhika Jagtap, Matteo Andreozzi, and Natalie Enright Jerger. 2020. Mocktails: Capturing the Memory Behaviour of Proprietary Mobile Architectures. In 2020 ACM/IEEE 47th Annual International Symposium on Computer Architecture (ISCA). 460–472. https://doi.org/10.1109/ ISCA45697.2020.00046 [8] Mario Badr and Natalie Enright Jerger. 2014. SynFull: Synthetic traffic models capturing cache coherent behaviour. In 2014 ACM/IEEE 41st International Symposium on Computer Architecture (ISCA). 109–120. https://doi.org/10.1109/ISCA. 2014.6853236 [9] Ali Bakhoda, George L. Yuan, Wilson W. L. Fung, Henry Wong, and Tor M. Aamodt. 2009. Analyzing CUDA workloads using a detailed GPU simulator. In 2009 IEEE International Symposium on Performance Analysis of Systems and Software. 163–174. https://doi.org/10.1109/ISPASS.2009.4919648 [10] Nathan Binkert, Bradford Beckmann, Gabriel Black, Steven K. Reinhardt, Ali Saidi, Arkaprava Basu, Joel Hestness, Derek R. Hower, Tushar Krishna, Somayeh Sardashti, Rathijit Sen, Korey Sewell, Muhammad Shoaib, Nilay Vaish, Mark D. Hill, and David A. Wood. 2011. The gem5 simulator. SIGARCH Comput. Archit. News 39, 2 (Aug. 2011), 1–7. https://doi.org/10.1145/2024716.2024718 [11] Long Chen, Oreste Villa, Sriram Krishnamoorthy, and Guang R. Gao. 2010. Dynamic load balancing on single- and multi-GPU systems. In 2010 IEEE International Symposium on Parallel & Distributed Processing (IPDPS). 1–12. https://doi.org/10.1109/IPDPS.2010.5470413 [12] Tianshi Chen, Zidong Du, Ninghui Sun, Jia Wang, Chengyong Wu, Yunji Chen, and Olivier Temam. 2014. DianNao: a small-footprint high-throughput accelerator for ubiquitous machine-learning. In Proceedings of the 19th International Conference on Architectural Support for Programming Languages and Operating Systems (Salt Lake City, Utah, USA) (ASPLOS ’14). Association for Computing Machinery, New York, NY, USA, 269–284. https://doi.org/10.1145/2541940.2541967 [13] Derek Christ, Lukas Steiner, Matthias Jung, and Norbert Wehn. 2024. PIMSys: A Virtual Prototype for Processing in Memory. In Proceedings of the International Symposium on Memory Systems (MEMSYS ’24). Association for Computing Machinery, New York, NY, USA, 26–33. https://doi.org/10.1145/3695794.3695797 [14] Jeffrey Dean, Greg S. Corrado, Rajat Monga, Kai Chen, Matthieu Devin, Quoc V. Le, Mark Z. Mao, Marc’Aurelio Ranzato, Andrew Senior, Paul Tucker, Ke Yang, and Andrew Y. Ng. 2012. Large scale distributed deep networks. In Proceedings of the 26th International Conference on Neural Information Processing Systems 12

Eidola: Modeling Multi-GPU Network Communication Traffic in Distributed AI Workloads

100 Billion Parameters. In Proceedings of the 26th ACM SIGKDD International Conference on Knowledge Discovery & Data Mining (Virtual Event, CA, USA) (KDD ’20). Association for Computing Machinery, New York, NY, USA, 3505–3506. https://doi.org/10.1145/3394486.3406703 [34] Achref Rebai, Mubarak Adetunji Ojewale, Anees Ullah, Marco Canini, and Suhaib A. Fahmy. 2024. SqueezeNIC: Low-Latency In-NIC Compression for Distributed Deep Learning. In Proceedings of the 2024 SIGCOMM Workshop on Networks for AI Computing (Sydney, NSW, Australia) (NAIC ’24). Association for Computing Machinery, New York, NY, USA, 61–68. https://doi.org/10.1145/ 3672198.3673801 [35] Gang Ren, Eric Tune, Tipp Moseley, Yixin Shi, Silvius Rus, and Robert Hundt. 2010. Google-Wide Profiling: A Continuous Profiling Infrastructure for Data Centers. IEEE Micro 30, 4 (2010), 65–79. https://doi.org/10.1109/MM.2010.68 [36] Kyle Roarty and Matthew D. Sinclair. 2020. Modeling Modern GPU Applications in gem5. gem5 Users Workshop (2020). https://par.nsf.gov/biblio/10192411 [37] Gabin Schieffer, Ruimin Shi, Stefano Markidis, Andreas Herten, Jennifer Faj, and Ivy Peng. 2024. Understanding Data Movement in AMD Multi-GPU Systems with Infinity Fabric. In SC24-W: Workshops of the International Conference for High Performance Computing, Networking, Storage and Analysis. 567–576. https: //doi.org/10.1109/SCW63240.2024.00079 [38] A. N. Semakin. 2021. Simulation of a multi-core computer system in the gem5 simulator. AIP Conference Proceedings 2318, 1 (02 2021), 090006. https://doi.org/ 10.1063/5.0035841 [39] Alexander Sergeev and Mike Del Balso. 2018. Horovod: fast and easy distributed deep learning in TensorFlow. arXiv:1802.05799 [cs.LG] https://arxiv.org/abs/ 1802.05799 [40] Mohammad Shoeybi, Mostofa Patwary, Raul Puri, Patrick LeGresley, Jared Casper, and Bryan Catanzaro. 2020. Megatron-LM: Training Multi-Billion Parameter Language Models Using Model Parallelism. arXiv:1909.08053 [cs.CL] https: //arxiv.org/abs/1909.08053 [41] Wei Sun, Ang Li, Sander Stuijk, and Henk Corporaal. 2024. How Much Can We Gain From Tensor Kernel Fusion on GPUs? IEEE Access 12 (2024), 126135–126144. https://doi.org/10.1109/ACCESS.2024.3411473 [42] Yifan Sun, Trinayan Baruah, Saiful A. Mojumder, Shi Dong, Xiang Gong, Shane Treadway, Yuhui Bao, Spencer Hance, Carter McCardwell, Vincent Zhao, Harrison Barclay, Amir Kavyan Ziabari, Zhongliang Chen, Rafael Ubal, José L. Abellán, John Kim, Ajay Joshi, and David Kaeli. 2019. MGPUSim: enabling multi-GPU performance modeling and optimization. In Proceedings of the 46th International Symposium on Computer Architecture (Phoenix, Arizona) (ISCA ’19). Association for Computing Machinery, New York, NY, USA, 197–209. https://doi.org/10.1145/3307650.3322230 [43] Zhuo Tang, Lifan Du, Xuedong Zhang, Li Yang, and Kenli Li. 2022. AEML: An Acceleration Engine for Multi-GPU Load-Balancing in Distributed Heterogeneous Environment. IEEE Trans. Comput. 71, 6 (2022), 1344–1357. https: //doi.org/10.1109/TC.2021.3084407 [44] Hugo Touvron, Louis Martin, Kevin Stone, Peter Albert, Amjad Almahairi, Yasmine Babaei, Nikolay Bashlykov, Soumya Batra, Prajjwal Bhargava, Shruti Bhosale, Dan Bikel, Lukas Blecher, Cristian Canton Ferrer, Moya Chen, Guillem Cucurull, David Esiobu, Jude Fernandes, Jeremy Fu, Wenyin Fu, Brian Fuller, Cynthia Gao, Vedanuj Goswami, Naman Goyal, Anthony Hartshorn, Saghar Hosseini, Rui Hou, Hakan Inan, Marcin Kardas, Viktor Kerkez, Madian Khabsa, Isabel Kloumann, Artem Korenev, Punit Singh Koura, Marie-Anne Lachaux, Thibaut Lavril, Jenya Lee, Diana Liskovich, Yinghai Lu, Yuning Mao, Xavier Martinet, Todor Mihaylov, Pushkar Mishra, Igor Molybog, Yixin Nie, Andrew Poulton, Jeremy Reizenstein, Rashi Rungta, Kalyan Saladi, Alan Schelten, Ruan Silva, Eric Michael Smith, Ranjan Subramanian, Xiaoqing Ellen Tan, Binh Tang, Ross Taylor, Adina Williams, Jian Xiang Kuan, Puxin Xu, Zheng Yan, Iliyan Zarov, Yuchen Zhang, Angela Fan, Melanie Kambadur, Sharan Narang, Aurelien Rodriguez, Robert Stojnic, Sergey Edunov, and Thomas Scialom. 2023. Llama 2: Open Foundation and Fine-Tuned Chat Models. arXiv:2307.09288 [cs.CL] https://arxiv.org/abs/2307.09288 [45] Rafael Ubal, Byunghyun Jang, Perhaad Mistry, Dana Schaa, and David Kaeli. 2012. Multi2Sim: a simulation framework for CPU-GPU computing. In Proceedings of the 21st International Conference on Parallel Architectures and Compilation Techniques (Minneapolis, Minnesota, USA) (PACT ’12). Association for Computing Machinery, New York, NY, USA, 335–344. https://doi.org/10.1145/2370816.2370865 [46] Guibin Wang, YiSong Lin, and Wei Yi. 2010. Kernel Fusion: An Effective Method for Better Power Efficiency on Multithreaded GPU. In 2010 IEEE/ACM Int’l Conference on Green Computing and Communications & Int’l Conference on Cyber, Physical and Social Computing. 344–350. https://doi.org/10.1109/GreenComCPSCom.2010.102 [47] Yipeng Wang, Ganesh Balakrishnan, and Yan Solihin. 2015. MeToo: Stochastic Modeling of Memory Traffic Timing Behavior. In 2015 International Conference on Parallel Architecture and Compilation (PACT). 457–467. https://doi.org/10. 1109/PACT.2015.36 [48] Haicheng Wu, Gregory Diamos, Srihari Cadambi, and Sudhakar Yalamanchili. 2012. Kernel Weaver: Automatically Fusing Database Primitives for Efficient GPU Computation. In 2012 45th Annual IEEE/ACM International Symposium on

Microarchitecture. 107–118. https://doi.org/10.1109/MICRO.2012.19 [49] Kan Wu, Zejia Lin, Mengyue Xi, Zhongchun Zheng, Wenxuan Pan, Xianwei Zhang, and Yutong Lu. 2025. GoPTX: Fine-grained GPU Kernel Fusion by PTXlevel Instruction Flow Weaving. In 2025 62nd ACM/IEEE Design Automation Conference (DAC). 1–7. https://doi.org/10.1109/DAC63849.2025.11132627 [50] Ehsan Yousefzadeh-Asl-Miandoab, Ties Robroek, and Pinar Tozun. 2023. Profiling and Monitoring Deep Learning Training Tasks. In Proceedings of the 3rd Workshop on Machine Learning and Systems (Rome, Italy) (EuroMLSys ’23). Association for Computing Machinery, New York, NY, USA, 18–25. https://doi.org/10.1145/ 3578356.3592589 [51] Yichao Yuan, Advait Iyer, Lin Ma, and Nishil Talati. 2024. Vortex: Overcoming Memory Capacity Limitations in GPU-Accelerated Large-Scale Data Analytics. Proc. VLDB Endow. 18, 4 (12 2024), 1250–1263. https://doi.org/10.14778/3717755. 3717780

13