DIVISION ELIMINATION

Abstract

Integer division on current NVIDIA GPUs costs roughly 20–40 cycles, and tensor indexing performs several divisions per access. On indexing-bound kernels this overhead dominates. We present a compile-time-first approach: divisors known at compile-time are encoded as template parameters, and a Granlund–Montgomery multiply-and-shift specialization is selected at instantiation time. Measured on an RTX 5070 Ti against runtime idiv, the operation-level speedup is 12× to 30× across 16 divisors at both 32-bit and 64-bit widths. We also find that the dominant win comes from Granlund–Montgomery itself rather than its delivery mechanism: variants that pass the magic struct as a kernel argument, read it from __constant__ memory, or template-specialize it are all statistically tied, which clarifies what the compile-time encoding does and does not buy.

1. Background

Granlund and Montgomery's 1994 result is well known: division by a runtime-invariant integer can be replaced by a multiply-high and a shift, with constants derived from the divisor. Compilers apply this transformation whenever the divisor is a literal visible to the optimizer — x / 8 becomes a shift, x / 10 becomes a multiply-and-shift at -O2. This is not what this paper is about.

The relevant case is when the divisor is a compile-time constant in the programmer's mental model but not in the compiler's. A transformer's hidden dimension, head count, or KV stride is fixed before any request begins, but by the time it reaches a CUDA kernel it has been threaded through several function boundaries, a kernel launch, and the host/device ABI. The optimizer no longer sees a constant — it sees a register or a kernel parameter — and the transformation is no longer available to it. The kernel pays for the full IDIV instruction even though the divisor never actually varies.

Cathedral Architecture (see Philosophy) treats this as a type-system problem: any value that is fixed once per request should be a template parameter, not a runtime argument. When the divisor is a template parameter, the constants the Granlund–Montgomery transformation needs can be computed at compile time, and the kernel emits a multiply-and-shift directly. This paper measures what that costs and what it buys.

2. Method

2.1 Hardware and Toolchain

• GPU: NVIDIA GeForce RTX 5070 Ti (Compute Capability 12.0, Blackwell)
• Host OS: Windows 10.0.26200
• Host compiler: MSVC 19.44.35227.0
• Device compiler: NVCC, Release configuration

2.2 Variants Measured

Each kernel performs the same workload — three back-to-back loops of 65,536 divisions per thread, accumulating into output[idx] — and differs only in how the divisor is delivered to the kernel:

• native-division — stock operator/ with the divisor passed as a size_t kernel argument. Compiler emits hardware IDIV.
• magic-division-rt — Granlund–Montgomery magic struct passed by value as a kernel argument. Compiler sees a multiply-high and shift, but the magic constants live in registers.
• magic-division-rt-const — Magic struct stored in __constant__ memory, accessed by global lookup. Compiler emits multiply-and-shift; constants come from the constant cache.
• magic-division-ct — Divisor is a template parameter; the magic struct is materialized as a static constexpr inside the kernel via collect_values(). Compiler sees the constants as immediates.

The four-way comparison isolates two questions. Native vs. any-magic measures the cost of hardware IDIV. RT vs. RT-const vs. CT measures whether the compile-time encoding does anything Granlund–Montgomery alone does not.

2.3 Methodology

Measurements were taken with bnch_swt in GPU-event mode. Each variant ran 2,000 launches with 20 measured iterations after stabilization. Each launch operated on a fresh page of input data (input + iter * N) to prevent cache reuse from contaminating timings. Inputs are __restrict__-qualified to allow the compiler to vectorize loads where possible. The kernel writes output[idx] inside the loop body, which prevents the optimizer from hoisting the division out of the loop — confirmed by the differentiated timings across variants. Sixteen divisors were measured at both 32-bit and 64-bit widths: power-of-two cases (32, 64, 256, 512, 1024, 2048, 4096), common transformer dimensions (768, 11008, 14336), and arbitrary values (100, 997, 1337, 102039132, and the 64-bit maximum-range case 10¹⁹).

3. Results

3.1 Native vs. Magic — Operation Cost of IDIV

Across every divisor and both widths, native division is the slowest variant by a wide margin. The slow path is consistently 1300–1720 MB/s regardless of divisor; the magic variants are 15,000–52,000 MB/s. The native throughput floor is dictated by the GPU's IDIV latency, not by any property of the divisor.

Two patterns are visible. First, the 64-bit case shows a bimodal distribution: power-of-two divisors that fit a shift-only path (32, 64, 256, 2048) reach ~50 GB/s, while general magic-number divisors plateau around 38 GB/s. Second, the 32-bit case shows the same split but at lower absolute throughput, because the kernel emits more arithmetic per division at narrower widths. Across the entire matrix, native division never closes the gap.

3.2 The Three Magic Variants Are Statistically Tied

The more interesting result: across the great majority of measured divisors, magic-division-rt, magic-division-rt-const, and magic-division-ct are statistically indistinguishable (95% confidence, by bnch_swt's reporting). The compile-time encoding does not produce faster code than the runtime-delivered Granlund–Montgomery struct, once that struct exists in any form.

The few exceptions occur where one variant lands in a slightly different microarchitectural path — for the 64-bit case at divisors 512 and 1024, magic-division-rt-const performs anomalously worse, and at divisor 64 the same variant pulls ahead. These differences are inside the run-to-run noise band for the rest of the matrix and we do not read them as evidence of a systematic effect.

The honest reading: the dominant win is Granlund–Montgomery itself. Once the multiply-and-shift is reaching the SASS, where the magic constants come from — registers, constant cache, or compile-time immediates — does not measurably change throughput on this workload.

3.3 What the Compile-Time Encoding Does Buy

The throughput numbers are the same; the engineering surface is not.

• No runtime initialization. The magic-division-rt and magic-division-rt-const variants require a host-side collect_values() call to compute the magic constants and either a kernel-argument marshalling or a cudaMemcpyToSymbol to deliver them. The CT variant computes them at compile time and emits them directly into the kernel. For an inference engine that launches kernels at request granularity, this removes a class of one-time setup steps from the hot path.
• Power-of-two cases collapse to a shift automatically. The CT variant's if constexpr (is_power_of_2(static_divisor)) branch emits a single SHR for power-of-two divisors with no runtime check. The RT variants execute the same magic-number multiply-and-shift regardless. This is visible in the 64-bit data: at power-of-two divisors, the CT variant reaches the ~50 GB/s tier alongside the others, but its instruction stream is shorter — the compiler emits a shift instead of a UMULHI + SHR.
• Type-level enforcement. A divisor that is a template parameter cannot be accidentally overwritten or passed through the wrong kernel argument slot. A divisor that is a runtime value can be both.

3.4 What It Doesn't Buy

• Throughput. On this benchmark, the CT variant is not faster than the RT variants. Anyone who has Granlund–Montgomery and a way to deliver the magic struct to the kernel will see the same numbers.
• End-to-end model speedup. These figures are operation-level on an indexing-bound microkernel. A compute-bound matmul kernel, where division is amortized over hundreds of cycles of MMA, will see proportionally smaller end-to-end gains. The relevant question is the fraction of kernel runtime spent on indexing arithmetic, not whether division is fast in isolation.

4. Discussion

4.1 The Compiler Already Does This — When It Can

x / 768 with a literal 768 is already a magic-number multiply-and-shift at -O2. The whole contribution of this work is the case where the compiler cannot do this: when the divisor is fixed in the programmer's mental model but threaded through enough indirection that the optimizer has lost it. Encoding the divisor as a template parameter restores the constant to the compiler's view of the world, and the existing transformation it already knows how to perform becomes applicable again.

A reasonable critique: "Why not just inline more aggressively, or use __forceinline__ and a literal, instead of templates?" In practice, the divisors live in a configuration struct that flows through layer construction, kernel launch, and several helper templates before reaching the arithmetic. Inlining alone does not propagate the constant across a __global__ boundary — kernel parameters are not constant-folded. Template parameters are.

4.2 Why the Three Magic Variants Tie

The RTX 5070 Ti (Blackwell, sm_120) has enough integer throughput that the multiply-and-shift form is bandwidth-bound rather than instruction-bound for this workload. The magic struct occupies 16 bytes; passing it by value adds two registers to the kernel signature, reading it from __constant__ memory hits the constant cache on the first warp and stays there. Neither is a measurable cost relative to the multiply-high instruction itself. We would expect the gap between variants to widen on older architectures with tighter register pressure or worse constant-cache behavior, and to widen further when the magic struct is one of many compile-time values rather than the only one.

4.3 Generalizing the Pattern

Division is the easiest case to measure because the speedup is large and the workload is easy to construct. The pattern — identify values fixed once per request, encode them in the type system, let the compiler specialize — applies anywhere the compiler has lost track of a constant it could otherwise act on. Memory strides, masking constants, modular arithmetic, branch conditions that depend on configuration: each is a candidate for the same treatment. The size of the win in each case depends on what the corresponding hardware instruction costs and how much of the kernel runtime is spent on it.

5. Reproducibility

The full benchmark harness is available at github.com/nihilai-collective. The div_mod_logic implementation lives in Nihilus; the benchmark driver uses bnch_swt's CUDA event-timing mode. All measurements in this paper were taken with the harness, GPU, and toolchain described in Section 2.1 and can be reproduced by building the harness against a CUDA 12+ toolchain on Blackwell or later hardware. We have not measured on Ada, Ampere, or earlier; the qualitative result (native IDIV is the bottleneck) is expected to hold, but absolute throughput will differ.

6. Conclusion

On indexing-bound CUDA kernels, replacing hardware IDIV with Granlund–Montgomery multiply-and-shift produces a 12–30× operation-level speedup on the RTX 5070 Ti, across 16 divisors at both 32-bit and 64-bit widths. The compile-time encoding of the divisor is not the source of the speedup — the multiply-and-shift form is — but it is the cheapest delivery mechanism: no host-side setup, automatic shift-collapse for power-of-two cases, and type-level enforcement of which divisors are actually constant. The win is real, the source of the win is Granlund–Montgomery, and the contribution of Cathedral Architecture is making the transformation reachable in cases where the compiler had otherwise lost the constant.

"Whatever's already known when you write the code, let the compiler act on it." — Cathedral Architecture

References

1. Granlund, T. and Montgomery, P. L. Division by Invariant Integers using Multiplication. PLDI 1994.
2. Warren, H. S. Hacker's Delight, 2nd ed., Chapter 10. Addison-Wesley, 2012.
3. NVIDIA Corporation. CUDA C++ Programming Guide. Performance Guidelines: Arithmetic Instructions.
4. Nihilai Collective. Cathedral Architecture: An Overview. 2026. /philosophy/