GPU-Tile-Sim: A Tile-Centric GPU Simulation Framework for LLM Hardware-Software Co-Design
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.
Paper Prompts
Sign up for free to create and run prompts on this paper using GPT-5.
Top Community Prompts
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"
Collections
Sign up for free to add this paper to one or more collections.