Deep-learning compilers like TVM let engineers optimize GPU kernels by declaratively reshaping loops — split this loop, reorder that one, bind the third to threads — which is easy to write but too blunt to express the tricks that make hand-written libraries fast (like double-buffering, which needs to touch a loop's body at fine grain), and it forces hours-long tuning because the search space is tied to input dimensions; Hidet (Toronto + AWS) flips the abstraction to make the assignment of computation-to-threads a first-class object embedded directly in the program, so any fine-grained optimization becomes expressible, fusion happens automatically after scheduling, and the search space is tied to hardware instead of input size — beating ONNX Runtime and tuned TVM by up to 1.48× while cutting tuning time 11–20×.
The argument in one glance
-
Problem
Writing fast GPU kernels by hand does not scale
Every deep-learning model is a graph of operators — convolutions, matrix multiplications, attention — and each must be compiled into a GPU kernel that runs near the hardware's peak. Hand-writing these kernels in CUDA gets the last drop of performance (that is what cuDNN and cuBLAS do), but it is laborious, error-prone, demands deep GPU expertise, and must be redone for every new operator, input shape, and fusion pattern. With operators multiplying and accelerators growing more complex, hand-tuning cannot keep up.
-
What people did
Compile kernels automatically by separating math from schedule
The dominant answer (Halide, then TVM) splits a kernel into two pieces: the COMPUTATION DEFINITION (the math — element (i,j) of the output is this sum) and the SCHEDULE (how it runs — this loop order, this thread binding). Developers write declarative loop-oriented scheduling primitives — split a loop, reorder loops, fuse them, bind one to threads — and an auto-tuner (AutoTVM, Ansor) searches loop tile sizes to fit each input shape. Anyone can now optimize without writing raw CUDA.
-
Why it fell short
Reshaping loops is too blunt, and the search is tied to input size
Loop-oriented scheduling manipulates the loop STRUCTURE but cannot reach inside a loop's BODY at fine grain — so the optimizations that make hand-written libraries fast (double-buffering, thread-block swizzle, direct Tensor-Core instructions, multi-stage async prefetch) simply cannot be expressed. Fusion is a second pain: a fused sub-graph is scheduled by its "anchor" operator's template, so that template must anticipate every possible fusion pattern, ballooning its complexity. And the search space is INPUT-CENTRIC: tile factors must divide the input dimensions, so the space explodes with those dimensions (up to 100 million schedules for one convolution), it misses good non-dividing tile sizes, and it even fails outright on prime dimensions — taking 15 hours to tune one CNN.
-
This paper's idea
Make the computation-to-thread assignment a first-class program object
Hidet (University of Toronto + AWS) replaces loop-reshaping with the TASK-MAPPING paradigm. It observes that everything loop-scheduling does reduces to two jobs — assign computations (tasks) to processing units (workers, e.g. threads), and order the tasks each worker runs — so it makes that assignment-and-ordering a first-class object, a "task mapping", embedded directly in the tensor program. Now scheduling happens IN the program, at any granularity, so every hand-library optimization becomes expressible. Two payoffs follow. POST-SCHEDULING FUSION: schedule the anchor operator alone, then automatically graft neighbor operators on as prologues/epilogues — no template needs to anticipate fusion. HARDWARE-CENTRIC SPACE: because task mappings use predicated (bounds-checked) loads, tile sizes need not divide the input, so the search space depends on the HARDWARE (under 200 schedules) not the input — 100,000× smaller. Result: up to 1.48× faster than ONNX Runtime and tuned TVM, with 11–20× less tuning time.
Why this problem is worth solving
A deployed deep-learning model is executed billions of times,
so shaving latency off its kernels saves energy, money, and —
in autonomous driving or interactive serving — keeps the
system usable. The field split into two camps to produce
those kernels. Hand-written kernel libraries (cuDNN, cuBLAS,
CUTLASS) hit near-peak performance because their authors can
reach for any low-level trick, but they cover only a fixed
menu of operators and fusion patterns, and extending them is
expensive expert work.
systems (TVM, XLA) that
take an operator’s mathematical definition and automatically
generate a GPU kernel for it, rather than calling a
hand-written library routine. Their promise is generality:
they can compile ANY operator and, crucially, FUSE adjacent
operators into one kernel to cut memory traffic — a
flexibility libraries lack. Their historical weakness is that
the auto-generated code often trails the hand-tuned
libraries, precisely because the scheduling language they use
cannot express the libraries’ best tricks
writing a kernel as
two decoupled parts. The COMPUTATION defines only the math:
what each output element equals. The SCHEDULE defines only
the execution: loop order, tiling, which loop maps to GPU
threads, what lives in fast memory. The same computation can
be paired with many schedules to explore performance, and an
auto-tuner searches for the best schedule per input shape.
This separation is why a non-expert can optimize kernels at
all — but it is only as powerful as the language the schedule
is written in
The paper’s thesis is that the schedule language everyone
uses —
the declarative style where
you optimize by transforming the LOOP STRUCTURE: “split this
length-1024 loop into 16×64”, “reorder these two loops”,
“fuse these loops into one”, “bind this loop to the thread
index”. Each primitive rewrites the nest of loops, and the
tuner picks the numeric factors. It is intuitive and
composable, and it made auto-scheduling practical — but it
operates only on the loops themselves, never on the fine
detail of what happens inside a loop iteration
What people did before, and where it broke
- Hand-written kernel libraries (cuDNN, cuBLAS). Peak performance via any low-level trick, but a fixed operator menu, limited fusion, and expensive to extend — the very generality problem compilers exist to solve.
- C++ template libraries (CUTLASS). Generate kernels for many input shapes on the fly and tune hyper-parameters, but do not reduce the difficulty of writing a new operator and fuse only a small predefined set.
- Loop-oriented compilers + tuners (TVM with AutoTVM / Ansor). The state of the art, and the paper’s main foil. Broadly general, but loop-reshaping cannot express fine-grained optimizations (double-buffering and friends), so it trails libraries on the compute-heavy operators; its anchor-template fusion is hard to extend; and its input-centric search takes hours-to-days and can miss or fail to find good schedules.
- Auto-schedulers (Ansor). Generate schedule templates from rules to avoid hand-writing them — but a new rule must be compatible with every existing rule and general across all operators, so extending the rule set is itself hard, and the loop-only expressiveness ceiling remains.
- The gap. No compiler let a developer schedule at arbitrary granularity (matching hand libraries), fuse without anchor-template gymnastics, AND search a compact input-independent space. Those three wants share one root cause — the loop-oriented abstraction — and Hidet replaces it.
How Hidet works
Schedule with task mappings, fuse after, search hardware-centric
-
Import model, optimize the graph
A model from PyTorch or ONNX becomes a computation graph. Standard graph-level passes (constant folding, partitioning into fusible sub-graphs) run first, and each sub-graph gets an anchor operator (its most compute-heavy member — a convolution or matrix multiply) to schedule.
-
Task mapping: assign tasks to workers new
The core abstraction. A "task mapping" is a first-class object saying which computations (tasks — e.g. the 512 element-loads of a tile) go to which processing units (workers — e.g. 128 threads) and in what order. Two primitive mappings compose: spatial(16,8) assigns tasks to distinct workers in parallel; repeat(4,1) makes one worker do several tasks in sequence. Composed — repeat(4,1) × spatial(16,8) — they say "128 threads, each loading 4 elements." The developer iterates over the tasks a worker owns by calling the mapping with the thread index, then implements each task by its index.
-
Schedule directly in the program, at any grain new
Because the mapping lives IN the tensor program (not applied as an external loop rewrite), the developer can drop to plain loops on some dimensions, interleave loads and computes inside a loop body, and express what loop-scheduling cannot: double-buffering, thread-block swizzle, direct Tensor-Core MMA instructions, multi-stage prefetch. The optimizations of hand libraries become first-class code.
-
Post-scheduling fusion new
Fusion is decoupled from scheduling. Schedule the anchor operator ALONE — its template never has to anticipate fusion. Then automatically graft neighbors on: operators with no reduction (elementwise, transpose, reshape) attach BEFORE as prologues (injective) or AFTER as epilogues (bijective). One anchor schedule now serves every fusion pattern, instead of a template that must cover them all.
-
Hardware-centric schedule space new
Task mappings use predicated loads (each access is bounds-checked), so a tile size need not divide the input dimension — "partial tiling" is legal. That decouples the search space from the input: tile sizes are chosen from hardware-friendly values (64×64, 128×64, …), giving under 200 schedules for matrix multiply versus up to 100 million for AutoTVM. The space is small enough to enumerate exhaustively in about a minute — and it does not fail on prime dimensions.
-
Lower and generate CUDA
The scheduled, fused tensor program in Hidet's IR is lowered, optimized, and emitted as CUDA C, compiled by the CUDA toolchain. Hidet needs only two hand-written schedule templates (matrix multiply and reduction) to cover all evaluated models — convolutions are converted to matrix multiply, most other operators handled by rules.
The idea
Change what the schedule language describes, and the expressiveness ceiling lifts. The paper’s core insight is diagnostic before it is constructive: loop-oriented scheduling is limited not because it lacks primitives but because it describes the wrong thing — the shape of loops rather than the assignment of work. Every loop primitive (split, reorder, fuse, bind) ultimately does one of two jobs: decide which processing unit runs which computation, or decide the order a unit runs its computations. Hidet names those two jobs directly as a task mapping and makes it the program’s first-class scheduling object. Because the mapping is embedded in the program rather than applied as an external loop rewrite, the developer can interleave operations inside a loop body — which is exactly where double-buffering, prefetch, and direct Tensor-Core instructions live. The optimizations were never impossible on the hardware; they were unsayable in the language. Give the language the right subject and they become ordinary code.
Decouple fusion from scheduling so neither has to know about the other. Loop-oriented fusion couples the two: the anchor operator’s schedule template must be written to handle every sub-graph it might end up fused into, which is what makes those templates monstrous and hard to extend. Post-scheduling fusion inverts the order — schedule the anchor as if it were alone, then graft neighbors on afterward, classified cleanly by their data-flow shape (injective operators can be prologues, bijective ones can be epilogues). This is the separation-of-concerns move the site keeps meeting: a template author reasons about ONE operator’s performance; the fusion machinery reasons about connectivity; neither is entangled with the other, so both stay simple and extensible.
Bind the search space to the hardware, not the input — a smaller space that is also better. The most counterintuitive result is that Hidet’s search space is five orders of magnitude SMALLER than AutoTVM’s yet contains BETTER schedules. The reason is that input-centric spaces are large for a bad reason: they enumerate every tile factor that divides the input, which balloons combinatorially, misses non-dividing tile sizes that are often optimal, and collapses to empty on prime dimensions. By using predicated loads to permit partial tiling, Hidet frees tile sizes to be hardware-friendly constants independent of the input, so the space is tiny, exhaustively searchable in a minute, robust to any input size, and — because those hardware-aligned tiles are what actually run fast — of higher average quality. Small, correct, and better all at once: the input dependence was pure liability.
Cooperatively loading a tile
128 threads load a 64×8 tile — loop-scheduling versus task mapping
From the paper's running example: in a matrix-multiply kernel, all threads of a block cooperatively load a 64×8 fragment of matrix A into shared memory. 512 element-loads, 128 threads, 4 loads each.
Step 1 of 5
Loop-scheduling: rewrite the nest
Start from a default double loop (for i in 64, for k in 8). To assign the 512 loads to 128 threads, apply declarative primitives: split the i-loop into 16×4, fuse the inner part with k, bind that fused loop to threadIdx.x. The tuner picks the numeric factors. It works for THIS assignment — but the whole vocabulary is loop rewrites, so anything needing to touch the loop body (double-buffer the load, prefetch the next tile) is inexpressible.
split/fuse/bind manipulate loop STRUCTURE only — never the fine detail inside an iteration this is why hand-library tricks like double-buffering cannot be written in this style at all
Why it works
- Loop-oriented scheduling can only reshape loops, never their bodies → task mappings embed the schedule in the program so operations can be interleaved at any grain → double-buffering, prefetch, and direct Tensor-Core instructions become expressible, closing the gap to hand libraries.
- All loop primitives reduce to assigning tasks to workers and ordering them → naming that assignment as a composable first-class object (spatial × repeat) captures everything loop-scheduling did, plus what it could not → one abstraction replaces the whole primitive set with more power.
- A fused template that must anticipate every fusion pattern is complex and brittle → scheduling the anchor alone and grafting injective prologues / bijective epilogues afterward → fusion and scheduling stay decoupled and each stays simple and extensible.
- Predicated loads permit tiles that do not divide the input → tile sizes become hardware-friendly constants independent of input dimensions → the search space shrinks ~100,000×, is exhaustively searchable, robust to prime sizes, and richer in fast schedules.
- Minutes-not-hours tuning → the compiler can be invoked inside graph-level optimization and neural-architecture search loops that need kernel latencies fast → compilation moves from an offline bottleneck to a usable inner loop.
What it costs
- A steeper authoring model. Task mappings are more expressive but lower-level than declarative loop primitives; writing them demands understanding worker assignment and the hardware, closer to CUDA thinking than to “split and reorder.” The paper mitigates this by needing only two hand-written templates, but the paradigm asks more of a template author.
- Template coverage still required. Hidet leans on two expert-written schedule templates (matrix multiply, reduction) and converts other operators to reuse them; operators that fit neither, or need a genuinely new template, still need expert effort.
- GPU-focused evaluation. The ideas are claimed general (CPUs, TPUs), but the implementation and all measurements are NVIDIA-GPU CUDA; other backends are asserted, not shown.
- Hardware-centric space is a curated bet. Its small size comes from restricting tiles to hardware-friendly values; that captures the fast region for the evaluated operators, but a workload whose optimum lies outside the curated tile set would be missed — the space is small because it is opinionated.
- Gains are workload-dependent. Up to 1.48× but 1.22× average, and on some models (a lightweight MobileNet case) it trails a baseline — the headline is a best case, and the broad win is the tuning-time reduction plus competitive latency, not a uniform speedup.
The evaluation
Five representative models (ResNet-50, Inception-V3, MobileNet-V2, BERT, GPT-2) on an NVIDIA RTX 3090, against PyTorch, ONNX Runtime, and TVM equipped with AutoTVM and Ansor, measuring both latency and tuning cost.
Two scheduling paradigms for a DL compiler
| Property | Loop-oriented (TVM + AutoTVM / Ansor) | Hand libraries (cuDNN / CUTLASS) | Hidet (task-mapping) |
|---|---|---|---|
| Fine-grained optimizations (double-buffer, direct MMA, prefetch) | |||
| Operator + fusion generality | Fixed operators, limited predefined fusion | ||
| Search space | N/A (hand-tuned) | ||
| Tuning time | None at compile time (pre-tuned by experts) | ||
| End-to-end latency (5 models) | Competitive but trails on compute-heavy ops | Near-peak on their supported operators |
Check yourself
Check your understanding
1. Loop-oriented scheduling has a rich set of primitives (split, reorder, fuse, bind) and years of refinement, yet cannot express double-buffering — an optimization hand-written libraries use routinely. Predict the fundamental reason, and how task mappings escape it.
2. Hidet's schedule space for matrix multiplication is ~100,000× SMALLER than AutoTVM's, yet contains BETTER schedules. Predict why a smaller space can be higher quality, and what makes AutoTVM's large.
3. In loop-oriented compilers, fusing Conv2d-BN-ReLU requires the convolution's schedule template to handle the whole sub-graph. Hidet instead schedules the convolution alone and fuses afterward. Why is this decoupling the better structure?
4. Task mappings are built from two primitives, spatial and repeat, that compose (e.g. repeat(4,1) × spatial(16,8)). Predict what distinguishes the two and why just these two suffice to replace the whole loop-primitive vocabulary.