Add fast_divmod magic-number division helper

[kashif]

Adds a magic-number division helper so a runtime divisor becomes a multiply and two shifts instead of a hardware integer divide, which GPUs do not have.

fastdivmod_magic(divisor) precomputes the (magic, shift) pair on the host. fast_divmod(dividend, divisor, magic, shift) does the device-side divmod, and FastDivmod wraps both for a compile-time constant divisor. It is valid for non-negative dividends below 2^31, the same contract as cutlass::FastDivmod.

The intended use is tile schedulers that do one divmod per work tile from a grid dimension known only at launch, where the compiler cannot fold the divisor. For a single divmod the generated code on gfx1150 drops from 57 to 25 instructions.

Tests check the host magic against Python // and % over a range of divisors, plus device tests comparing fast_divmod and the FastDivmod class to native divmod on GPU.

[Copilot]

Pull request overview

Adds a host/device “magic-number” division utility to avoid expensive integer division sequences on GPUs, exposing both a low-level divmod function and a convenience wrapper for constant divisors. This fits into FlyDSL’s expr numeric utilities and is intended to help kernel scheduling/indexing code where divisors are only known at launch.

Changes:

• Added fastdivmod_magic(divisor) host helper to compute (magic, shift) for 0 <= n < 2**31.
• Added fast_divmod(dividend, divisor, magic, shift) device-side divmod implementation plus FastDivmod wrapper for compile-time constant divisors.
• Added unit + device tests validating the host math and GPU behavior against native // and %.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated no comments.

File	Description
python/flydsl/expr/numeric.py	Implements fastdivmod_magic, fast_divmod, and FastDivmod in the numeric DSL layer.
python/flydsl/expr/typing.py	Re-exports the new helpers/types through flydsl.expr’s public surface (__all__).
tests/unit/test_fast_divmod.py	Adds L0 host verification and L2 GPU kernels comparing results to native divmod.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[sjfeng1999]

This is not enough. FastDivmod need to support dynamic divisor value and could be a valid leaf in the int_tuple.

[kashif]

thanks for the pointer! checking