Papers
Topics
Authors
Recent
Search
2000 character limit reached

GPU-Tile-Sim: A Tile-Centric GPU Simulation Framework for LLM Hardware-Software Co-Design

Published 13 Jul 2026 in cs.DC | (2607.11262v1)

Abstract: Modern LLM workloads increasingly rely on optimized GPU kernels through hardware-software co-design. These kernels achieve high-performance through fine-grained dependency scheduling and computation-memory overlap. As such, they incur new challenges on existing GPU performance models. Instruction-driven simulators are costly to adapt to evolving architectures, while analytical models are too coarse to capture kernels' characteristics. We propose GPU-Tile-Sim, a tile-centric GPU simulation framework for LLM hardware-software co-design. The key insight is that modern LLM kernel performance is governed less by individual instruction latency than by the dependency structure that controls execution order and overlap. Accordingly, GTSim represents kernel execution as a warp-level tile graph whose nodes capture tile-level operations and whose edges encode data and ordering constraints. Using this representation, we design an automatic tile-graph frontend and a graph-driven simulation backend. We evaluate GTSim on representative GEMM, attention, and end-to-end LLM inference workloads. On A100 and H100 across both conventional and highly optimized kernels, GTSim achieves high performance-modeling accuracy (MAPE, Mean Absolute Percentage Error, 1.22%--8.71%). We further extend GTSim to Blackwell with preliminary validation, and demonstrate its effectiveness in analyzing software and architectural design choices.

Summary

  • The paper introduces GTSim, a tile-centric simulator that replaces instruction-driven simulation with dependency-aware tile graphs for modeling LLM kernels.
  • It employs a graph-driven backend with throughput and latency models to simulate dense, fused, and asynchronous kernel execution accurately.
  • Empirical evaluations show GTSim achieving MAPE as low as 1.22% for GEMM and consistently low errors on attention workloads across modern GPU architectures.

GPU-Tile-Sim: A Tile-Centric Simulator for LLM-Oriented GPU Hardware-Software Co-Design

Motivation: LLM Kernel Patterns, Architectural Evolution, and Modeling Challenges

The paper introduces GTSim, a GPU simulation framework that departs from instruction-driven simulation, instead leveraging a tile-centric, dependency-aware abstraction to model performance-critical LLM kernels. Modern LLM workloads exhibit execution patterns characterized by either single-stage iterative kernels, such as those underlying dense/fused GEMM and feed-forward operations, or multi-stage coupled kernels embodying complex attention mechanisms, fused computational stages, and intricate data dependencies (Figure 1). Figure 1

Figure 1: Representative kernel patterns in LLM workloads, contrasting standard transformer block structure with iterative and coupled kernel execution.

These kernels are increasingly scheduled to exploit architectural advances in the Ampere, Hopper, and Blackwell GPU generations—embracing techniques such as asynchronous data movement (e.g., cp.async, TMA), kernel fusion, multi-buffered software pipelines, and warp specialization. This evolution is reflected in the shifting design of GPU SMs and memory hierarchies (Figure 2). Figure 2

Figure 2: GPU architectural evolution, highlighting the transition from tightly coupled warp-scoped execution to distributed, decoupled, and explicitly managed multi-stage execution across generations.

The shift toward explicit intra-kernel coordination and dependency-driven overlap introduces modeling challenges that undermine both instruction-driven simulators (limited architectural extensibility) and traditional analytical or mapping-based models (insufficient expressiveness for fine-grained dependencies and kernel fusions). Kernels in modern LLM workloads often deploy data-movement/compute scheduling patterns and synchronization behaviors that are not representable via prior coarse abstractions or single-warp representative models.

GTSim Framework Overview

GTSim combines a tile-graph-based frontend and a graph-driven, throughput-oriented simulation backend for LLM kernel evaluation and co-design (Figure 3). It eschews per-instruction simulation, instead representing kernels as directed acyclic graphs where nodes correspond to warp-level tile operations and edges encode explicit data and order dependencies. Figure 3

Figure 3: GTSim high-level workflow, illustrating the pipeline from tile-centric input to graph-driven simulation execution.

This design enables GTSim to model explicit synchronization, fine-grained overlap, and heterogeneous execution roles within and across warp groups, supporting both manually crafted and automatically derived input from intermediate representations such as TileLang IR. The simulator is further parameterized by throughput and latency models for compute, memory, NoC, and SM-level schedulers.

Tile Graph Abstraction and Kernel Representation

The warp-centric tile graph encodes operations, data regions (tiles), execution group assignments, and their dependencies (Figure 4). Nodes are annotated with execution semantics (operation kind, hardware mapping, tile layout), while edges are either data dependencies (producer-consumer) or order constraints (arising from synchronization or pipeline reuse). Figure 4

Figure 4: Tile graph abstraction: (a) Edge and node semantics; (b) Fused GEMM graph instantiating producer/consumer warps; (c) Various software pipeline organizations, including warp-specialized and multi-buffered variants.

This abstraction allows direct, fine-grained expression of kernel fusion, warp specialization (producer/consumer/epilogue roles), and software pipelining by appropriate assignment of nodes and dependency edges. Figure 5 presents example pipeline organizations for a fused GEMM + SiLU kernel, demonstrating flexibility in representing distinct execution regimes. Figure 5

Figure 5: Four pipeline organizations for fused GEMM + SiLU, differing in producer/consumer assignment and pipeline depth.

Tile-graph construction is automated for TileLang IR—preserving structural and synchronization information post software pipeline and warp-role lowering—facilitating coverage of both regular and complex dynamic LLM kernels (Figure 6). Figure 6

Figure 6: Automatic construction of tile graphs from compiler IR, translating loop, pipeline, and role structure into graph nodes, edges, and execution groups.

Graph-Driven Simulation: Backend Execution and Hardware Modeling

GTSim's backend advances kernel execution by issuing graph nodes (tile operations) to appropriately modeled hardware resources (compute, memory, NoC) based on explicit data and order dependencies (Figure 7). Scheduling and execution exploit concurrency by decomposing nodes into throughput-limited sub-operations aligned with pipeline width and resource models. Figure 7

Figure 7: Simulation workflow, showing threadblock and warp scheduling, node issue, resource binding, and dependency-driven activation.

Compute nodes are mapped to pipelines modeling Tensor Cores, CUDA cores, and other units, parameterized by latency, issue width, and sustainable throughput. Memory operations are derived from tile descriptors, supporting both conventional (address-decomposed) and bulk transfer (TMA) paths. The NoC/DSMEM model enables explicit block-to-block communication for DSMEM-enabled fusion workloads.

This throughput-oriented, resource-aware abstraction yields fidelity sufficient for performance modeling of both regular and highly asynchronous LLM kernels, without the overheads of full instruction trace execution.

Empirical Evaluation: Fidelity and Scalability

GTSim is validated across GEMM and attention kernels (including FP8, fused, and multi-stage), as well as end-to-end LLM inference on A100, H100, and preliminary Blackwell configurations.

For GEMM workloads, GTSim achieves MAPE as low as 1.22% (A100/H100) and remains below 5.4% across all evaluated variants, outperforming TileFlow and LLMCompass by wide margins, particularly on fused and low-precision kernels (Figure 8). Figure 8

Figure 8: Parity plots for GEMM workloads; GTSim tracks real hardware closely across architectures, substantially outperforming analytical baselines.

On representative attention workloads (FlashAttention-3, Flash-Decoding, FlashMLA), GTSim consistently yields MAPE in the 2.37–6.50% range, while baseline models often exceed 100% error due to their inability to capture fused, fine-grained execution (Figure 9). Figure 9

Figure 9: GTSim achieves consistently low MAPE for attention kernels, capturing fused and overlapped execution structure missed by coarser models.

End-to-end Llama-3-8B inference is modeled on H100 hardware: GTSim attains kernel-level average MAPE of 8.71% for prefill and decode, and remains stable across batch size and sequence length variations, indicating robust modeling capacity (Figure 10). Figure 10

Figure 10: GTSim matches measured Llama-3-8B inference cycle breakdowns and can scale to long-KV decode scenarios.

Efficiency tests demonstrate 3.5–4.6× higher simulation speed than Accel-Sim while maintaining comparable accuracy, and analysis reveals that explicit modeling of order and synchronization edges in the tile graph is necessary to attain GTSim’s accuracy.

Case Studies: Co-Design Analysis Enabled by GTSim

Software Pipelining

A study on fused GEMM + SiLU demonstrates how different pipeline organizations (naive, WS cooperative, WS ping-pong, WS 3-stage) impact throughput depending on whether the kernel is memory or compute-bound (Figure 11). The 3-stage pipeline achieves optimal utilization by decoupling epilogue and mainloop assignments, enabled by GTSim’s explicit dependency modeling. Figure 11

Figure 11: Software pipelining impacts depend strongly on regime (memory- or compute-bound) and the explicitness of synchronization/overlap in the pipeline design.

NoC-Enabled Fusion

Evaluation of cluster-based fusion using NoC vs. DRAM for intermediate data reveals up to 1.94× speedup, and topology-aware threadblock mapping yields up to 1.91× lower NoC residence time for small clusters (Figure 12). Figure 12

Figure 12: NoC communication and topology-aware mapping demonstrably improve fused-kernel execution, particularly in cluster-based regimes.

Blackwell Architecture Adaptation

GTSim’s modular design enables rapid support for Blackwell’s TMEM and new Tensor Core semantics, validated by a comparison of FlashAttention-3 and FlashAttention-4 kernels on B200 (Figure 13). Performance gains in FA4 are attributed to larger threadblock granularity and more decoupled intra-threadblock workflows, elucidating how architectural shifts favor kernel redesign. Figure 13

Figure 13: FA4 outperforms FA3 on B200 due to increased threadblock granularity and pipeline decoupling, both of which are accurately captured by GTSim.

Implications and Future Developments

GTSim’s tile-centric, dependency-aware abstraction enables high-fidelity performance modeling and rapid architectural adaptation, facilitating principled hardware-software co-design for evolving LLM workloads. It directly models fused, asynchronous, and warp-specialized kernels found in SOTA foundation models and efficiently accommodates future architectural paradigms such as those in Blackwell.

The approach supports input-dependent dynamic kernels, cluster-fused execution, and hardware-mapped communication primitives, highlighting the suitability of tile-graph simulation as a basis for next-generation AI accelerator evaluation frameworks. The implications extend to hardware-software DSLs, dynamic kernel generation, and future LLM system stack co-design.

Conclusion

GTSim establishes a unified, extensible framework for accurate, efficient simulation of modern GPU workloads, particularly those critical to LLM inference and training. By explicitly modeling dependency structures through warp-centric tile graphs and leveraging a modular, graph-driven simulation backend, it overcomes the extensibility and modeling gaps of previous approaches. The framework sets a new bar for LLM hardware-software co-design tools, supporting systematic exploration of emerging kernel paradigms and architectural innovations (2607.11262).

Paper to Video (Beta)

No one has generated a video about this paper yet.

Whiteboard

No one has generated a whiteboard explanation for this paper yet.

Explain it Like I'm 14

Explaining “GPU-Tile-Sim: A Tile-Centric GPU Simulation Framework for LLM Hardware-Software Co-Design”

Overview: What this paper is about

This paper introduces a new way to “pretend run” (simulate) how big AI programs, like LLMs, run on powerful graphics chips called GPUs. The tool is called GPU-Tile-Sim (GTSim). It helps engineers test ideas for faster software and better hardware without having to build real chips or write low-level code for every tiny instruction.

Think of a busy kitchen: many cooks (GPU “warps”) work together, passing ingredients (data) and dishes (results) between stations. Modern AI code runs fast only if this teamwork is carefully planned so cooks don’t wait around. GTSim is like a smart kitchen simulator. It doesn’t track every hand motion of a cook; instead, it looks at how teams coordinate, what steps depend on each other, and how work overlaps to keep everyone busy.

Goals: What questions the paper tries to answer

The authors aim to:

  • Predict how fast LLM kernels (the heavy-duty parts of the program, like matrix multiplications and attention) will run on different modern GPUs.
  • Capture the real reasons for speed or slowness: not tiny instruction timings, but how steps depend on each other and how data movement overlaps with computation.
  • Make a simulator that’s accurate, but also easy to update when new GPU features come out (like new data-move engines or bigger teamwork groups).

Approach: How the method works (in simple terms)

Modern LLM code gets speed by splitting big jobs into “tiles,” overlapping data loading with computing, and assigning different roles to different warps (teams). GTSim models this directly.

  • Tile graph idea:
    • Picture a to-do map made of boxes and arrows.
    • Each box (a “node”) is a tile-sized job a warp or group of warps does, like “load this block of data to fast memory,” “multiply two small matrices,” or “write results back.”
    • Arrows show dependencies:
    • A solid arrow means “this step needs data produced by that earlier step.”
    • A dashed arrow means “do this after that for ordering/safety,” even if no data is shared (like waiting until a shared buffer is free).
  • Warp-centric view:
    • GPUs schedule work in units called “warps” (small teams of threads).
    • GTSim assigns each node to the warp or warp group that does it, so it naturally represents producer–consumer teams, specialized roles, and group operations.
  • Frontend and backend:
    • Frontend: It can automatically build the tile graph from a higher-level kernel description (using a DSL called TileLang). This keeps the big-picture structure (pipelines, roles, buffers, sync) clear.
    • Backend: Instead of simulating every instruction, it “schedules” ready nodes based on their dependencies and available hardware resources. It uses throughput-focused models for:
    • Compute units (like Tensor Cores) to estimate how quickly tiles are processed.
    • Memory systems (shared memory, L2 cache, DRAM) to estimate how fast tiles can be moved.
    • On-chip networks (NoC) and distributed shared memory (DSMEM) for communication between nearby GPU cores (SMs).
    • For memory moves, it expands tile descriptions into realistic memory transactions. For special engines like TMA (a bulk copy engine), it models their steady streaming behavior.
    • For compute, it can split a node into sub-operations so different warps’ work can overlap, respecting pipeline widths and throughput.

In short: GTSim runs the graph of tile tasks, releasing nodes when their arrows are satisfied, and accounts for how fast the hardware can push tiles through compute and memory “conveyor belts.”

Key results: What they found and why it matters

The authors tested GTSim on common LLM building blocks:

  • Matrix multiplications (GEMM)
  • Attention kernels (like FlashAttention variants)
  • End-to-end LLM inference (e.g., running Llama-3-8B)

They compared its predictions to real measurements on NVIDIA A100 and H100 GPUs and found:

  • Very low error on kernels: about 1.22% to 6.50% difference from real hardware.
  • For full LLM inference, error stayed low: about 7.06% to 8.71%.
  • It beat other higher-level analytical tools (like TileFlow and LLMCompass) on accuracy, especially for modern, highly optimized kernels.

Why this is important:

  • It shows that focusing on the dependency structure and overlap (the teamwork plan), not tiny instruction timings, is the right level for modern LLM kernels.
  • It gives developers a practical, trustworthy tool to explore designs before coding or buying hardware.

Implications: What this could change

  • Faster co-design: Hardware and software teams can try new ideas (like deeper pipelines, different warp roles, or new data-move strategies) and quickly see performance impacts.
  • Easier upgrades: When new GPU features arrive (like Hopper’s warp-group matrix ops or Blackwell’s newer tensor ops and memory), GTSim can add new node types and resource models without rewriting a giant instruction simulator.
  • Better kernels: By simulating attention, GEMM, and fused kernels with realistic overlap and sync behavior, developers can discover bottlenecks (memory vs. compute), tune buffering, and choose the best mapping of work to warps and SMs.
  • Future-ready: The authors even did a preliminary check on the next generation (Blackwell/B200) and used GTSim to guide early kernel ideas, suggesting it can keep up with fast hardware evolution.

In everyday terms: GTSim is like a smart kitchen planner that helps you design the most efficient cooking line for a huge banquet. It focuses on who hands what to whom and when, and how many dishes the ovens and fridges can handle at once. Because it plans at the right level, it’s accurate and easy to update when you buy new ovens or change the menu.

Knowledge Gaps

Unresolved gaps, limitations, and open questions

Below is a consolidated list of concrete gaps that remain uncertain or unexplored, framed to guide follow-up research and engineering.

  • Frontend generality: The automatic graph-extraction path is tied to TileLang IR; there is no implemented and validated frontend for Triton, CUDA/C++ (without an external lifting tool), PTX/SASS, or other DSLs. How to robustly and automatically lift diverse kernel sources to the same tile graph without manual effort?
  • Static DAG assumption: The simulation assumes a static, acyclic tile graph. How to model dynamic, data-dependent control flow (variable sequence lengths, MoE routing, block-sparse attention, early-exit), dynamic loop bounds, or runtime-dependent dependencies without unrolling everything a priori?
  • Multi-GPU scope: The simulator explicitly targets single-GPU kernels and omits inter-GPU communication (NVLink/NVSwitch, PCIe) and collectives (NCCL). How to extend the model to pipeline/tensor/data-parallel LLM inference/training and capture overlap with communication and collective synchronization?
  • On-chip NoC/DSMEM fidelity: The DSMEM and SM-cluster NoC are modeled with simplified bandwidth/latency/queueing abstractions. What topology-specific effects (e.g., routing, crossbar arbitration, link contention, multicast, remote-SMEM bank/port limits) materially affect kernel timing, and how should they be parameterized and validated?
  • Memory hierarchy simplifications: L2 is modeled as fully associative with LRU and there is no explicit modeling of L1/texture caches, sectoring, MSHRs, cache partitioning, or L2 compression. What accuracy is gained by introducing realistic associativity, banking, sector/line sizes, miss status handling, and compression-aware bandwidth?
  • Shared memory and DSMEM banking: The model does not account for SMEM/DSMEM bank conflicts, multicast, or port/bandwidth limits during high-throughput loads/stores and TMA fills. How often do bank conflicts throttle fused/pipelined kernels, and how should this be captured?
  • TMA modeling limits: The TMA engine is parameterized with fixed 128B transactions and a sustained issue rate (~100 B/cycle). What is the impact of (a) the number of concurrent TMA engines per SM/cluster, (b) descriptor setup and fence overheads, (c) burstiness, (d) interactions with L2 compression/cache operators, and (e) interleaving TMA with LD/ST traffic?
  • Tensor-core pipeline realism: Compute is modeled throughput-first, without operand-collector constraints, scoreboard hazards, register-file bandwidth/port limits, or instruction-issue quirks. Which of these microarchitectural effects materially alter overlapped pipelines (e.g., wgmma waves) and how to include them without losing extensibility?
  • Blackwell-specific features: Modeling of TMEM residency, tcgen05 single-thread issue semantics, and CTA-pair/TPC execution is only preliminarily validated. What are the precise ordering/availability constraints and bandwidth/latency envelopes for TMEM and CTA-pair collectives, and how should tile graphs represent them?
  • Warp scheduler fidelity: The backend uses policy approximations (e.g., greedy-then-oldest). How to calibrate scheduling policies per generation (A100/H100/B200), including barrier handling, fairness, and queue arbitration, and quantify sensitivity of predictions to scheduler choices?
  • Occupancy and resource constraints: The paper does not detail modeling of register-file allocation, barrier registers, SMEM capacity/partitioning, TMA-descriptor footprints, or per-SM warp/CTA slot constraints. How to incorporate these to predict occupancy-bound regimes and CTA residency accurately?
  • Coalescing and alignment: Tile-derived memory transactions may miss sub-warp coalescing inefficiencies, misalignment penalties, gather/scatter patterns (e.g., ragged attention), and sector utilization. What abstractions best capture alignment- and stride-induced bandwidth loss?
  • Address translation and memory faults: TLBs, page sizes, page-walks, and HMM/page migration are not modeled. Do TLB misses or page placement affect LLM kernels at realistic batch/sequence sizes, and how should they be introduced?
  • Control-plane and stream-level overlap: Kernel launch latency, CUDA Graphs, concurrent kernel execution, and multi-stream copy/compute overlap are not captured. How to extend the framework to predict end-to-end inference/training latency where host/device orchestration and stream concurrency matter?
  • Error attribution and sensitivity analysis: The paper reports MAPE but does not analyze error sources. Which components (compute throughput, cache model, NoC, scheduler, TMA) dominate residual error across kernels, and how robust are predictions to parameter miscalibration?
  • Simulation performance and scaling: There is no quantitative report of simulation throughput, memory footprint, or scaling with model size/sequence length/num-CTAs. What are the runtime/accuracy trade-offs versus instruction-level simulators and analytical models for large-scale LLM inference and training?
  • Training workloads: Evaluation focuses on GEMM/attention and end-to-end inference. How well does the model handle training backpropagation (activation recomputation, optimizer/gradient ops), optimizer fusion, and communication-compute overlap (e.g., ZeRO, DDP)?
  • Power/energy modeling: The framework is performance-only. How to add power/energy estimators (per-unit activity, DVFS states) to enable co-design under performance-per-watt constraints?
  • Reproducibility and parameterization: Details of microbenchmarks and calibrated parameters (e.g., L2/DRAM bandwidth, TMA rates, NoC latencies) are not fully disclosed. Can the authors release parameter sets per GPU generation and a standardized calibration suite?
  • Frontend–backend semantic drift: The graph is derived from a mid-level IR after pipeline/warp-role lowering. How to detect and correct drift between IR-level intent and final SASS (e.g., register allocation side-effects, scheduler/bundling changes) that alter overlap/synchronization?
  • Automatic extraction from binaries: There is no path to lift PTX/SASS traces or CUPTI/Nsight profiles into tile graphs. Can a trace-to-graph reconstruction be developed to broaden applicability to closed-source or hand-tuned kernels?
  • Persistent kernels and long-running loops: For persistent/threadblock-parked kernels and ring-buffer pipelines, how to model steady-state without full unrolling while preserving buffer-reuse/order constraints and startup/tail effects?
  • KV-cache and memory residency: The model does not explicitly capture KV-cache placement, tiling across HBM/L2/SMEM, or prefetching policies in decoding. How do cache location and prefetch schedule affect latency/bandwidth balance in attention and decoding kernels?
  • Quantization breadth: Only a subset of low-precision formats is discussed; int4/NF4 and custom dequantization pipelines are not evaluated. How to model dequantization/fusion kernels (including tensor core int/fp mixes) and their memory-compute overlap?
  • Portability beyond NVIDIA: The abstraction centers on NVIDIA features (cp.async, TMA, wgmma, DSMEM). What adaptations are needed for AMD/Intel GPUs (different wavefront/warp semantics, MFMA, LDS, SDMA, XNACK), and does the tile graph remain expressive?
  • Topology-aware fusion automation: The paper sketches NoC-enabled fusion and topology-aware mapping, but lacks an automated mapping/scheduling algorithm. Can the simulator be coupled with a mapper that searches tile placement and fusion layouts under NoC/DSMEM constraints?
  • Validation breadth for LLM kernels: Coverage of fused attention variants (FlashAttention 3/4, FlashDecoding, MLA) is claimed, but details on irregular layouts, causal masks, and paged KV caches are limited. Which attention/dataflow variants remain poorly predicted and why?

Practical Applications

Immediate Applications

The following applications can be deployed now with modest integration effort, leveraging the paper’s tile-centric graph abstraction, TileLang-based frontend, and throughput-oriented backend.

  • Kernel design and debugging assistant for fused LLM kernels (software, AI infrastructure)
    • Use GTSim to rapidly compare pipeline organizations (e.g., WS 2-buffer vs x-buffer), tilings, TMA vs LD/ST, and DSMEM usage for GEMM/attention/epilogues.
    • Potential tool/workflow: “TileGraph Visualizer” integrated into CUTLASS/TileLang/Triton developer workflows to surface dependency bottlenecks and overlap opportunities.
    • Assumptions/dependencies: TileLang IR or an equivalent frontend (e.g., TLX/CuBridge) to generate tile graphs; calibrated throughput parameters for target GPUs.
  • Compiler autotuning cost model in the loop (software, ML compilers)
    • Drive schedule/fusion/tiling search with GTSim as a fast, accurate oracle to prune suboptimal candidates before on-hardware trials.
    • Potential tool/workflow: “sim-in-the-loop autotuner” for TileLang (and adapters for Triton/MLIR).
    • Assumptions/dependencies: Single-GPU focus; accurate microbenchmark-based calibration for A100/H100 (preliminary for Blackwell/B200).
  • Cloud LLM inference capacity planning and SLO validation (cloud, finance, healthcare)
    • Forecast latency/throughput/GPU-hours for specific models (e.g., Llama-3-8B) across A100/H100 SKUs; evaluate kernel variants and fusion strategies before deployment.
    • Potential product: “LLM serving planner” dashboard that runs what-if simulations on kernel mixes and batch sizes.
    • Assumptions/dependencies: Workloads dominated by GEMM/attention-like kernels; no inter-GPU communication modeling; relies on vendor-specific parameters.
  • Performance regression gating in CI (software engineering)
    • Catch kernel performance regressions across CUDA/toolchain updates by comparing simulated cycles against baselines per architecture.
    • Potential workflow: CI step that runs GTSim on representative kernels and flags >X% deviation.
    • Assumptions/dependencies: Versioned calibration datasets; stable tile-graph extraction per compiler version.
  • NoC/DSMEM-aware fusion and cluster mapping (software, performance engineering)
    • Use the on-chip NoC model to choose threadblock cluster sizes, DSMEM placement, and synchronization to avoid hotspots and stalls on Hopper-like systems.
    • Potential tool: “NoC-aware fusion planner” for attention/decoder kernels using DSMEM.
    • Assumptions/dependencies: Simplified intra-GPU NoC and DSMEM model; no inter-GPU links.
  • Blackwell readiness assessment for existing kernels (semiconductor, software)
    • Predict the impact of tcgen05/TMEM/CTA-pair execution on current Hopper kernels; identify pipeline-depth and operand-residency adjustments.
    • Potential tool: “Blackwell readiness checker” for kernel/library maintainers.
    • Assumptions/dependencies: Preliminary Blackwell validation; parameters subject to change.
  • Teaching labs for modern GPU execution (academia, education)
    • Hands-on labs demonstrating warp specialization, software pipelining, and dependency-driven overlap via editable tile graphs.
    • Potential product: “GTSim Lab” notebooks bundled with curated kernels and parameter sets.
    • Assumptions/dependencies: Lightweight packaging and classroom-friendly examples.
  • Early-stage architectural parameter sweeps (semiconductor)
    • Explore sensitivity to Tensor Core pipeline widths, TMA issue rates, L2/DRAM bandwidth, and warp-scheduler policies without instruction-level simulators.
    • Potential workflow: Rapid DSE (design space exploration) over tile-level resource knobs.
    • Assumptions/dependencies: Throughput-oriented abstractions; abstracts away micro-ops and scoreboard details.
  • Robotics/edge GPU kernel tuning (robotics)
    • For Jetson-class (Ampere) devices, simulate alternative tilings/pipelines to meet real-time deadlines for on-device LLM or vision-language kernels.
    • Potential workflow: Pre-deployment tuning of time-critical kernels with hardware-constrained profiles.
    • Assumptions/dependencies: Requires calibration for specific Jetson SKUs; IR frontend availability.
  • Energy and cost estimation via performance proxy (energy, policy, procurement)
    • Combine simulated runtimes with power curves to approximate energy per token and TCO for procurement/compliance reporting.
    • Potential tool: “Perf→Energy estimator” integrating PUE and carbon-intensity factors.
    • Assumptions/dependencies: External power/thermal models; mapping performance-to-power for each GPU SKU.

Long-Term Applications

These applications require additional research, broader integration, or new capabilities (e.g., interconnect modeling, power models, standardization).

  • Closed-loop HW/SW co-design platform (semiconductor, software)
    • Automated search that jointly optimizes kernel tile graphs (fusion, roles, buffers) and architectural knobs (pipelines, NoC bandwidth) using GTSim-as-oracle.
    • Potential product: RL-driven co-designer that proposes hardware changes and codegen changes together.
    • Assumptions/dependencies: Extended APIs for architectural variability; large-scale calibration and search infrastructure.
  • Multi-GPU and distributed LLM simulation (cloud, HPC)
    • Extend to NVLink/NVSwitch/PCIe interconnects and collective ops to plan tensor/pipeline parallelism, sharding, and topology-aware fusion.
    • Potential tool: Cluster-level “inference topology planner” for latency/throughput targets.
    • Assumptions/dependencies: Accurate inter-GPU network and synchronization models; much larger tile graphs.
  • Cross-vendor, vendor-neutral tile-graph IR and simulators (software, standards)
    • Standardize a tile-graph representation consumable by TileLang, Triton, MLIR, and map to NVIDIA/AMD/Intel/TPU backends for portable co-design.
    • Potential product: “TileGraph-IR” spec and reference simulators.
    • Assumptions/dependencies: Community and vendor buy-in; robust frontends (e.g., TLX-like, CuBridge-like) across ecosystems.
  • Online runtime advisors for adaptive kernels (systems, software)
    • Distill GTSim into fast surrogates to adjust pipeline depth, warp roles, and prefetch distance in production using live counters.
    • Potential product: “Kernel autopilot” that autotunes at runtime.
    • Assumptions/dependencies: Low-overhead inference of surrogate models; safe reconfiguration mechanisms.
  • Power/thermal-aware co-optimization (energy, data centers)
    • Integrate detailed power and thermal models to co-design for energy-per-token and carbon-aware scheduling under datacenter constraints.
    • Potential tool: Energy/SLA co-optimizer for LLM serving.
    • Assumptions/dependencies: Validated power/thermal models and telemetry; coupling with facility-level models (PUE, cooling).
  • Static checking and verification of fused kernels (software, security)
    • Use explicit data/order edges to prove absence of races/deadlocks and to validate buffer reuse and synchronization correctness.
    • Potential product: “tilegraph-lint” integrated into compilers and CI for correctness and performance safety.
    • Assumptions/dependencies: Formalization of semantics and sound analysis; compiler hooks.
  • Auto-fusion and kernel synthesis under NoC constraints (software)
    • Co-synthesize fusion patterns, tilings, and DSMEM usage that maximize utilization without violating NoC bandwidth/latency budgets.
    • Potential tool: “NoC-constrained fusion synthesizer” emitting TileLang/Triton kernels.
    • Assumptions/dependencies: Stronger NoC/DSMEM models; robust codegen and validation.
  • Sector-specific LLM appliances and kernels (healthcare, finance, education)
    • Co-design kernels (and potentially microarchitectural profiles) tuned for domain constraints (e.g., low-latency triage in healthcare, compliance logging in finance, on-prem education appliances).
    • Potential product: Verticalized accelerators and kernel libraries with certified performance envelopes.
    • Assumptions/dependencies: Custom hardware partnerships; regulatory and validation pathways.
  • Rich performance explainability and IDE integration (education, industry)
    • Interactive UIs to navigate dependency graphs, resource heatmaps, and pipeline timelines for developers and students.
    • Potential product: IDE plugins for VS Code/JetBrains with live simulation overlays.
    • Assumptions/dependencies: Productization effort; stable APIs from simulators to visualization.

Notes on overarching assumptions and dependencies:

  • Frontend availability: Immediate path is via TileLang IR; analogous paths (e.g., TLX for Triton, CuBridge for CUDA) are emerging but may require engineering.
  • Calibration: Accuracy depends on per-architecture microbenchmarks (A100/H100 validated; Blackwell preliminary). New features necessitate new node types and parameter tuning.
  • Scope: The current simulator targets single-GPU kernels with intra-GPU DSMEM/NoC modeling; inter-GPU networking is out of scope and needed for distributed applications.
  • Workload fit: Best aligned with LLM-style kernels (GEMM/attention/fused epilogues); highly irregular control flow or extreme cache-pathological codes may need enhanced cache/bank models.

Glossary

  • Accel-Sim: A GPU simulator that executes instruction-level traces to model microarchitectural behavior. Example: "instruction-driven simulators (e.g., GPGPU-Sim~\cite{bakhoda2009gpgpusim}, Accel-Sim~\cite{khairy2020accelsim}, gem5 and its GPU extensions~\cite{binkert2011gem5,power2015gem5gpu})"
  • AMALI: An interval-based analytical performance model that estimates execution from stall-delimited intervals. Example: "Interval-based models (e.g., MDM~\cite{wang2020mdm}, GCoM~\cite{cha2022gcom}, AMALI~\cite{cao2025amali})"
  • cp.async: An NVIDIA Ampere feature for asynchronous global-to-shared memory transfers that enables copy–compute overlap. Example: "This shift is enabled by cp.async, which supports asynchronous global-to-shared transfers"
  • CTA-pair: A Blackwell execution feature where two cooperative thread arrays coordinate within a tensor processing cluster for larger-scale tensor-core execution. Example: "CTA-pair execution within a tensor processing cluster (TPC) enables larger-scale cooperative tensor-core execution."
  • CUDA cores: Scalar/vector ALUs on NVIDIA GPUs that execute general-purpose instructions (as distinct from Tensor Cores). Example: "such as CUDA cores, Tensor Cores, SFUs, LD/ST units, or TMA."
  • Distributed Shared Memory (DSMEM): An on-chip mechanism allowing inter-SM shared-memory access within a cluster for cooperative execution. Example: "DSMEM-based cluster cooperation together enable coordinated execution across multiple warps and even across SMs."
  • Epilogue (operations): Post-GEMM per-element computations (e.g., activation, scaling) and stores fused into the kernel. Example: "integrating TMA loads, warp-group tensor-core computation, epilogue operations, and a final store"
  • FlashAttention: A family of fused attention kernels that optimize memory traffic and on-chip reuse via multi-stage pipelines. Example: "FlashAttention and its subsequent variants fuse multiple stages into a single kernel"
  • GEMM: General Matrix Multiply; a core high-performance kernel widely used in LLMs. Example: "projection and feed-forward layers, which are commonly implemented using optimized GEMM kernels"
  • GPGPU-Sim: A detailed instruction-driven GPU simulator used for architectural research. Example: "instruction-driven simulators (e.g., GPGPU-Sim~\cite{bakhoda2009gpgpusim}, Accel-Sim~\cite{khairy2020accelsim}, gem5 and its GPU extensions~\cite{binkert2011gem5,power2015gem5gpu})"
  • Greedy-then-oldest (scheduling policy): A warp/node selection policy that prioritizes ready work greedily, then breaks ties by age. Example: "scheduled using standard policies such as greedy-then-oldest~\cite{rogers2012ccws}"
  • Instruction-driven simulator: A performance model that drives execution by simulating individual instructions through detailed microarchitectural components. Example: "Instruction-driven simulators...take instruction streams such as PTX or SASS traces as input and drive execution by simulating the behavior of each instruction"
  • Interval-based model: An analytical model that partitions execution into intervals (e.g., between stalls) and estimates cycles per interval for a representative warp. Example: "Interval-based analytical models...approximate execution by partitioning a representative warp's execution into stall-delimited intervals"
  • Kernel fusion: Combining multiple operators into a single kernel to reduce global memory traffic and improve locality. Example: "Kernel fusion has become a central technique to reduce memory traffic and improve locality in LLM workloads."
  • LD/ST units: Load/Store units responsible for memory access instructions on the GPU. Example: "such as CUDA cores, Tensor Cores, SFUs, LD/ST units, or TMA."
  • Mapping-based framework: A performance modeling approach that represents programs via loop/mapping abstractions instead of instructions. Example: "Mapping-based frameworks (e.g., Timeloop~\cite{parashar2019timeloop}, MAESTRO~\cite{kwon2019maestro}, TileFlow~\cite{zheng2023tileflowfusion}, LLMCompass~\cite{zhang2024llmcompass})"
  • Mean Absolute Percentage Error (MAPE): An accuracy metric expressing average absolute error as a percentage of ground truth. Example: "GTSim achieves MAPE (Mean Absolute Percentage Error~\cite{hyndman2006forecast_accuracy}) from 1.22\% to 6.50\%"
  • Network-on-Chip (NoC): The on-chip interconnect and associated models for routing and bandwidth-limited data movement between GPU components. Example: "throughput-oriented compute, memory, and NoC abstractions"
  • Producer–consumer coordination: A programming pattern where warps specialize into data producers and consumers and coordinate via dependencies and synchronization. Example: "encourage programming styles based on producer--consumer coordination, warp specialization, and explicit data placement"
  • PTX: NVIDIA’s intermediate GPU assembly language used as an input trace for simulators and analytical models. Example: "take instruction streams such as PTX or SASS traces as input"
  • SASS: NVIDIA’s machine-level GPU assembly (ISA) used in low-level tracing and simulation. Example: "take instruction streams such as PTX or SASS traces as input"
  • SFU (Special Function Unit): Dedicated units for special mathematical functions (e.g., transcendental ops) on the GPU. Example: "such as CUDA cores, Tensor Cores, SFUs, LD/ST units, or TMA."
  • SiLU: Sigmoid Linear Unit; an activation function often fused with GEMM in LLM kernels. Example: "Representative software pipeline organizations for a fused GEMM + SiLU kernel."
  • Software pipelining: Organizing computation and data movement into overlapped stages (e.g., multi-buffered) to hide latency. Example: "Software pipelining is widely used to overlap data movement and computation within a kernel."
  • Static Single Assignment (SSA): A representation where each variable/tile version is assigned exactly once, easing dependency tracking. Example: "assigning versioned tile descriptors in a static single assignment (SSA)-like manner"
  • Streaming Multiprocessor (SM): The fundamental GPU execution unit that schedules and runs warps and thread blocks. Example: "kernels execute on streaming multiprocessors (SMs)"
  • Subpartition (SM subpartition): A per-SM subdivision with its own scheduler that issues instructions for assigned warps. Example: "Each SM has multiple subpartitions, and each subpartition has an independent scheduler"
  • Tensor Core: Specialized matrix-multiply units providing high-throughput low-precision compute on NVIDIA GPUs. Example: "recent GPU generations continue to raise peak compute throughput...such as Tensor Memory Accelerator, distributed share memory, and Tensor Memory"
  • Tensor Memory (TMEM): A Blackwell-era operand storage space decoupled from registers for tensor-core execution. Example: "Tensor Memory (TMEM) + SMEM"
  • Tensor Memory Accelerator (TMA): A hardware engine for bulk asynchronous tile transfers between global and shared memory. Example: "We model TMA operations as bulk tile transfers handled by a dedicated TMA front-end"
  • Tensor Processing Cluster (TPC): A grouping of SMs enabling cooperative tensor-core execution (e.g., CTA-pair) on Blackwell. Example: "CTA-pair execution within a tensor processing cluster (TPC) enables larger-scale cooperative tensor-core execution."
  • Throughput-oriented model: A performance model focusing on sustained throughput, issue width, and latency of high-level operations rather than per-instruction pipelines. Example: "throughput-oriented compute, memory, and NoC abstractions"
  • Tile descriptor: Metadata describing a tile’s storage, type, shape, layout, and coordinates used to derive memory transactions. Example: "The tile descriptor specifies the tile produced or accessed by a node"
  • Tile graph: A dependency graph where nodes are warp-level tile operations and edges encode data and ordering constraints. Example: "models GPU kernel execution as a dependency-driven warp-centric tile graph"
  • TileLang IR: A tile-level intermediate representation emitted by the TileLang DSL compiler for constructing the tile graph. Example: "currently instantiated with TileLang IR~\cite{wang2025tilelang}"
  • Triton TTGIR: An extended Triton IR that preserves role specialization and dependencies before lowering. Example: "extended TTGIR before downstream lowering"
  • tcgen05: A Blackwell tensor-core generation/primitive introducing single-thread issue semantics for tensor-core execution. Example: "Blackwell's tcgen05 introduces single-thread issue semantics for tensor-core execution"
  • Warp: The basic GPU scheduling unit comprising a group of threads that execute in lockstep. Example: "computation is organized around warps"
  • Warp specialization: Assigning different warps/warp-groups distinct roles (e.g., producer/consumer) to improve overlap and throughput. Example: "Warp specialization further enhances this overlap by assigning different warps or warp groups to distinct roles"
  • Warpgroup: A cooperating set of warps that collectively execute certain primitives (e.g., WGMMA) with coordinated semantics. Example: "warpgroup-level tensor-core execution (wgmma)"
  • WGMMA: Hopper warpgroup-level matrix multiply-accumulate primitive that coordinates multiple warps for tensor-core execution. Example: "warpgroup-level tensor-core execution (wgmma)"
  • Tile-based domain-specific language (DSL): A programming interface centered on tiles to express GPU kernels at a high level for analysis and transformation. Example: "a tile-based domain-specific language (DSL), currently instantiated with TileLang IR"

Open Problems

We haven't generated a list of open problems mentioned in this paper yet.

Collections

Sign up for free to add this paper to one or more collections.

Tweets

Sign up for free to view the 1 tweet with 74 likes about this paper.