Skip to content

Faster microfloat conversions: generated bit-twiddles and branch-free narrowing - #35

Merged
AntonOresten merged 5 commits into
mainfrom
generated-twiddles
Oct 21, 2026
Merged

AntonOresten merged 5 commits into
mainfrom
generated-twiddles

Conversation

@AntonOresten

@AntonOresten AntonOresten commented Sep 25, 2026 •

Copy link
Copy Markdown
Member

Faster conversions in three parts. All match cvt_generic bit for bit:

  1. Microfloat → microfloat: bit-twiddles generated from the cvt_generic results replace the hand-written Float4_E2M1FN => Float8_E4M3(FN) specializations. They cover every pair, including user types.
  2. Widening into Float16, BFloat16 and Float32: generated per type the same way.
  3. Narrowing from Float32: a new branch-free path, cvt_twiddle, for all six rounding modes and both overflow policies.

How it works

@cvt_table already builds the full result table of a pair from cvt_generic. It now also fits that table with linear pieces (i << k) + c over the source bits, chosen with branch-free compares and selects:

  • Exact widenings: one piece covers the source normals, and each binade of source subnormals adds one more.
  • Saturated and NaN results: constant pieces.
  • Sign: when the sign carries over unchanged, it moves separately as one extra shift.

A table that fits in few pieces becomes a twiddle; otherwise it stays a lookup. Because the code is built from the reference results, it is correct by construction. All 12 exact built-in pairs twiddle at cost 1–5 (pieces, plus one for the sign move). E2M1→E2M3 is a single shift. For E2M1→E4M3FN the output reproduces the old hand-written code. specializations.jl and TWIDDLED_PAIRS are removed. Packed conversions of any length compile to straight-line code through the lanewise path. For x2 this is identical to the old hand-written packed method.

Widening: @microfloat registers a generated cvt(F, x) per type, fitted over the bits of Float16, BFloat16 or Float32. Zero and all subnormals share one "head" piece evaluated on the FPU: put the subnormal significand under the exponent of 2^(1−bias), then subtract 2^(1−bias). The result is exact and normalized, with no subnormal operands. Float64 extends the Float32 result.

Per-target cutoff

New public hook max_twiddle_cost(): 8 on the host, overridden to 2 in CUDACoreExt. Each generated method contains both forms, and the choice folds when the method is compiled.

Microfloat → microfloat Widening
CPU (Cortex-X925, vectorized loop) 4.5–29× faster than the table lookup 2–4× faster for most type/destination pairs; E4M3/E3M4 → Float16 about 10% slower
GPU (GB10) table (one ld.global.nc) beats twiddles above cost 2 lookup or native PTX beats twiddles

The table only wins on the CPU above about 25 pieces. On the GPU, a lone shift still beats the table.

Narrowing from Float32

cvt_twiddle is the new default cvt for Float32 sources, and so for Float16, BFloat16 and every other source that converts through Float32 first.

  • Target normals: an integer add rounds the Float32 bits at the target's last significand bit, carrying into the exponent, and a subtraction adjusts the exponent bias.
  • Target subnormals: |x| is scaled to units of the smallest subnormal and rounded on the FPU.
  • Everything else (overflow, NaN, sign) is a select, and the constants fold from the type's layout. For E4M3FN it compiles to about 20 instructions with no branches.
  • Errors: the error checks are the only branches, and they fold away for formats that cannot fail.

Results and errors match cvt_generic over all 2³² Float32 inputs for every built-in format × 6 modes × 2 policies. This was checked offline; skipped inputs are only those that throw in both paths. The test suite uses boundary, tie and random inputs, on the GPU as well.

vs cvt_generic
CPU, signed formats with NaN (E5M2, E4M3, E3M4, E4M3FN) ~10× faster (0.22 vs 2.4 ns per value; the loop vectorizes)
CPU, FP4/FP6, E8M0, E5M3FNU ~2× faster (1.1 vs 2.8 ns); the NaN or negative-input throw that the semantics require blocks loop vectorization
GPU (GB10) ~3× faster; CUDACoreExt now falls back to it where no native instruction applies

Note: for formats without significand bits (E8M0), cvt_generic's RoundNearest rounds ties away from zero rather than to even: its odd-last-bit test reads the implicit leading 1. cvt_twiddle reproduces this exactly and documents it in a comment. Changing it would change E8M0 results, so it's left for a separate change.

Tests

  • The fitted twiddles reproduce every table: all type pairs including the unsigned test types, 6 rounding modes, SAT/OVF, and every widening destination.
  • The costs of the 12 exact pairs are pinned.
  • Packed x2/x4/x8 conversions are compared against the generic path.
  • Widening is compared bit-exactly against cvt_generic, now including Float16.
  • On the GPU, every built-in pair and bit pattern is compared with the host, which covers the device lookup path.

The full suite passes locally, GPU tests included (50,588).

@cvt_table keeps its name, though it now emits twiddles as well as tables.

🤖 Generated with Claude Code

AntonOresten and others added 2 commits September 25, 2026 19:23
@cvt_table now fits the cvt_generic results with linear pieces
(i << k) + c over the source bits and emits a branch-free twiddle when
few pieces suffice, a lookup otherwise. This covers all 12 exact
widenings among the built-in types (and user types), replacing the
hand-written Float4_E2M1FN => Float8_E4M3(FN) specializations and
TWIDDLED_PAIRS. Packed conversions of any length inherit the twiddle
through the lanewise path.

The cutoff is max_twiddle_cost(): 8 on the host, where twiddles
vectorize and beat the gather by 4-29x, and 2 on CUDA devices, where a
cached constant-table load beats longer twiddles.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The piecewise fit now works on any unsigned table width, so @microfloat
registers a generated cvt(F, x) per type that computes Float16, BFloat16
and Float32 results with a bit-twiddle over the destination bits. Zero
and all subnormals share one head piece evaluated on the FPU (the
significand under the exponent of 2^(1-bias), minus 2^(1-bias)), so
every built-in type costs 3-5 instead of up to 9 linear pieces.
Float64 extends the Float32 result, which beats 64-bit-lane twiddles.

On the CPU this is 2-4x faster than the BFloat16 lookup for most
type/destination pairs. Devices keep the lookup (or native PTX) through
max_twiddle_cost, which measured faster there.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@codecov

codecov Bot commented Sep 25, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 98.08917% with 3 lines in your changes missing coverage. Please review.
✅ Project coverage is 84.46%. Comparing base (0b4f053) to head (e26e93c).

Files with missing lines Patch % Lines
ext/CUDACoreExt/conversion.jl 0.00% 3 Missing ⚠️
Additional details and impacted files
@@            Coverage Diff             @@
##             main      #35      +/-   ##
==========================================
+ Coverage   80.49%   84.46%   +3.96%     
==========================================
  Files          15       14       -1     
  Lines         482      605     +123     
==========================================
+ Hits          388      511     +123     
  Misses         94       94              

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

cvt_twiddle is the new default cvt for Float32 sources, and so for
Float16, BFloat16 and every other source that converts to Float32 first.
Target normals round the Float32 bits with an integer add at the
target's last significand bit and rebias with a subtraction; target
subnormals round |x| scaled to units of the smallest subnormal on the
FPU. Overflow, NaN and sign are selects, and the constants fold from the
type's layout. The error checks are the only branches and fold away for
formats that cannot fail.

It covers all six rounding modes and both overflow policies, and matches
cvt_generic bit for bit, errors included: checked exhaustively over all
2^32 Float32 inputs for every built-in format, mode and policy.

On the CPU it is ~10x faster for signed formats with NaN (the loop
vectorizes) and ~2x for unsigned and FiniteOnly formats, whose error
branch prevents vectorization. On the GPU it is ~3x faster than the
generic path, and CUDACoreExt now falls back to it where no native
instruction applies.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@AntonOresten AntonOresten changed the title Generated bit-twiddles for microfloat conversions and widening Faster microfloat conversions: generated bit-twiddles and branch-free narrowing Sep 26, 2026
AntonOresten and others added 2 commits September 26, 2026 20:23
`any(isnan, xs)` over the lanes of a pair is not unrolled in device code:
it spills the lanes to local memory and loops around the single cvt.
That made native FP6/FP4 pairs ~2.7x slower than the portable twiddle
on the GB10; with a `|` reduction they are ~3x faster than it. Same fix
for the sign guard of ue8m0 pairs.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
On x86 without native BFloat16 (Julia 1.11+), inlining cvt_twiddle
into T(::BFloat16) let LLVM take the sign from a `bitcast bfloat to
i16` of a value it had also extended with fpext. The x86 backend then
reads the low half of the extended register, which is zero, so the sign
OR folded away: FiniteOnly targets lost the sign of every negative
BFloat16 input (the CI failures on x64 1.11, 1.12 and pre).

BFloat16 sources now widen by shifting their bits, so no bfloat
arithmetic reaches the conversion. Verified under x86-64 emulation on
Julia 1.11 and 1.12: every BFloat16 and Float16 bit pattern, every
built-in format, mode and policy matches the generic path. Adds a
regression test over all BFloat16 inputs.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@AntonOresten
AntonOresten merged commit 46108af into main Oct 21, 2026
11 checks passed
@AntonOresten
AntonOresten deleted the generated-twiddles branch October 21, 2026 16:15
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant