First-class fp16/bf16/fp8/fp4 element types for both backends

[johnnynunez]

Motivation

Reduced-precision formats could not be used as FKL element types: PerThreadRead<ND::_2D, __half> does not compile (no VectorTraits/thread-fusion registration, ~15 std::is_fundamental gates), and CPU-only builds have no such types at all. The only way in was the attention/ workaround: ad-hoc Read IOps that dequantize to float inside exec(). This PR adds the types to the library so any pipeline can read, compute (fp16/bf16), cast, saturate and write them on both backends.

What this adds

Five element types in namespace fk — fp16, bf16, fp8_e4m3, fp8_e5m2, fp4_e2m1 — plus their 1–4 channel vector companions (fp16_2, bf16_4, fp8_e4m3_4, …), usable in executeOperations / Cast / SaturateCast / thread fusion, e.g.:

Ptr2D<fp16> input(w, h), output(w, h);
executeOperations<TransformDPP<>>(stream,
    PerThreadRead<ND::_2D, fp16>::build(input.ptr()),
    Cast<fp16, float>::build(),
    Mul<float>::build(2.0f),
    SaturateCast<float, fp16>::build(),
    PerThreadWrite<ND::_2D, fp16>::build(output.ptr()));

The same source builds as _cpp (pure g++, zero CUDA headers) and _cu.

Design

• fk-owned types, identical on every backend. A single parametric constexpr softfloat core (MiniFloatFormat + one RN-even encoder/decoder pair) implements all five formats with NVIDIA-exact special-value policies (SATFINITE for fp8/fp4; e4m3 NaN=0x7F, no Inf; e2m1 NaN→+6, no NaN/Inf). No per-backend divergence of aggregate-ness/traits/operators, fully constexpr, and adding fp6/e8m0 later is a table entry. Layouts are bit-identical to __half/__nv_bfloat16/__nv_fp8_*/__nv_fp4_e2m1 (static_asserted against the real headers in the _cu utest).
• Native fast paths under nvcc. fp16/bf16 conversions branch on __builtin_is_constant_evaluated() to the native ctors at runtime (hardware cvt on device); toNative()/fromNative() are free bit-casts. fp8/fp4 are pure software on all paths, so core headers never include cuda_fp8.h/cuda_fp4.h and the API does not vary with the installed toolkit (kills the FK_HAS_FP8-style __has_include trap; sm_89+ cvt intrinsics are a follow-up).
• One-way implicit conversions. Implicit operator float() + explicit ctors: h * 2.0f and float f = h work, and the float-vs-__half ambiguity minefield can't happen. fp16/bf16 have full constexpr arithmetic/comparisons (promote-compute-demote in binary32 is correctly rounded: p ≥ 2·p_target+2). fp8/fp4 are conversion-only, exactly like CUDA.
• fp4 is one element per byte (low nibble, canonicalized on construction), matching CUDA's scalar __nv_fp4_e2m1. The packed 2-per-byte format cannot fit the RawPtr one-addressable-element model and remains a fused-read use case.
• Registration without touching the standard lists. New RFV* type lists append to VAll; the thread-fusion tables get an appended block (fp16→fp16_2 like short→short2; 1-byte formats →x4 like char→char4), pinned by exhaustive static_asserts including legacy row boundaries (the table is positional; a shifted entry would corrupt vectorized loads silently).
• Gate widenings: std::is_fundamental → fk::validScalar in AreSS/AreSV/AreVS, vector_at, cxp::Exec, cmp_*, max/min/abs, Discard, ArrayVector; cmp_*/saturate_cast promote reduced operands explicitly to float; isnan/isinf/abs classify by bits (fp8/fp4 have no operators, and e4m3/e2m1 NaN is undetectable via s != s).
• Hardening that outlives this PR: std::numeric_limits specializations (raw-bit constexpr) make fk::maxValue/minValue and the cxp limits work unmodified — and vlimits.h now static_asserts is_specialized instead of silently producing zero limits for unknown types (numeric_limits<__half> is unspecialized in CUDA 13.3; a SaturateCast to it would previously have clamped everything to 0 with no diagnostic). CanCompound now also checks the element-level operator exists (before, fp8x2 += fp8x2 would have been trait-visible and hard-errored inside the operator body). RGB2Gray no longer falls off the end of a non-void function for non-standard outputs.

Verification (RTX PRO 6000, CUDA 13.3 + g++ 11.5)

• utest_reduced_float_native_parity (_cu): bit-exact parity with NVIDIA's converters for all 65536 fp16 and all 65536 bf16 codes (software decoder vs __half2float/__bfloat162float), 2M sampled float encodes + a directed corpus vs the native ctors, all 256 fp8 codes and encode parity vs __nv_cvt_float_to_fp8(__NV_SATFINITE), all 16 fp4 codes + a 40k encode sweep vs __nv_fp4_e2m1 (including NaN→+6 and nibble packing), plus constexpr-vs-runtime agreement inside the same TU.
• utest_reduced_float_types (_cpp+_cu): layout/trait/constexpr asserts, exhaustive software roundtrips (which caught a real bug during development: bf16 subnormals decode below the binary32 normal range), a double-precision arithmetic oracle for the operator surface, RN-even tie cases per format, and a CPU-purity tripwire (#error if cuda_fp16.h leaks into a _cpp TU — CUDA-present CI cannot catch that otherwise).
• utest_reduced_float_pipelines (_cpp+_cu): executeOperations for all five scalar types (odd width → non-thread-divisible path), vector pipelines, loose vs .then()-fused bit-equality, a TF::ENABLED thread-fusion run, and TestCaseBuilder cases (the builder gained a third specialization for reduced scalar I/O; equalValues compares reduced types bit-exactly since a float tolerance quantizes to 0 on fp8/fp4 grids).
• Full suite: 74/74 ctest targets pass on both backends; an adversarial review pass over the diff was run and every confirmed finding fixed (including a compile regression where float2 a; a *= 2.0L; would have stopped working because VBase<long double> is unregistered — covered now by the split CanCompound specializations).

Out of scope / follow-ups

1. Migrating attention/'s Bf16AttentionRead/Fp8TokenDequantRead onto these types — needs an "identity read of element type E" trait first, so the mma/decode cp.async fast paths keep their compile-time dispatch (they currently match on the exact workaround op types).
2. sm_89+/sm_120 hardware cvt intrinsics for the fp8/fp4 runtime paths (behind the already-established __CUDA_ARCH__ branch point in the ctors).
3. A CUDA-less CI job for the _cpp targets, to machine-check the no-CUDA-headers invariant (with the toolkit installed, cuda_fp16.h compiles silently under plain g++, so current CI can't prove it).
4. Mixed fk::fp16↔__half expressions are ambiguous by design (both convert implicitly to arithmetic types) — documented at the toNative()/fromNative() interop functions.

🤖 Generated with Claude Code

[Copilot]

Pull request overview

This PR introduces first-class reduced-precision floating-point element types (fp16, bf16, fp8_*, fp4_*) to FKL so they can be used end-to-end in executeOperations pipelines on both CPU and CUDA backends, and updates core traits/utilities so these types participate in vectorization, thread-fusion, constexpr math, and saturating casts.

Changes:

• Adds new reduced-float scalar + vector element types (including traits, conversions, numeric limits, and CUDA interop for fp16/bf16).
• Extends core type-gating/vector utilities and thread-fusion tables to treat reduced floats as valid scalars/vectors.
• Adds comprehensive unit/integration tests for type properties and pipeline correctness on _cpp and _cu, plus native parity tests on _cu.

Reviewed changes

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

Show a summary per file

File	Description
utests/core/execution_model/utest_reduced_float_pipelines.h	New pipeline-level tests covering reduced-float reads/casts/muls/saturate/writes and thread fusion.
utests/core/data/utest_reduced_float_types.h	New type/trait/limits/constexpr/arithmetic tests for reduced-float scalars and vectors.
utests/core/data/utest_reduced_float_native_parity.h	New CUDA-only parity tests comparing software conversions against NVIDIA native converters.
tests/operation_test_utils.h	Updates equality and TestCaseBuilder coverage to support reduced-float scalar I/O.
include/fused_kernel/core/utils/vlimits.h	Hardens maxValue/minValue to error out when std::numeric_limits is unspecialized.
include/fused_kernel/core/utils/vector_utils.h	Registers reduced-float vector types/traits and widens scalar/vector gating to validScalar.
include/fused_kernel/core/execution_model/thread_fusion.h	Extends thread-fusion mapping tables with reduced-float entries and adds boundary pinning asserts.
include/fused_kernel/core/data/vector_types.h	Adds fk-owned reduced-float vector aggregates and includes reduced-float scalar definitions.
include/fused_kernel/core/data/reduced_float_types.h	New reduced-float scalar types, software encode/decode core, traits, numeric_limits, and CUDA interop helpers.
include/fused_kernel/core/data/ptr_nd.h	Removes a problematic post-throw return value (now returns T{}) in host at() path.
include/fused_kernel/core/data/array.h	Widens ArrayVector scalar constraints from fundamental-only to validScalar.
include/fused_kernel/core/constexpr_libs/constexpr_vector_exec.h	Allows constexpr op execution on reduced-float scalars via validScalar.
include/fused_kernel/core/constexpr_libs/constexpr_saturate.h	Promotes reduced-float sources to float for saturate_cast correctness and operability.
include/fused_kernel/core/constexpr_libs/constexpr_cmath.h	Adds reduced-float handling to isnan/isinf/cmp/max/min/abs via bit classification and promotion.
include/fused_kernel/algorithms/image_processing/color_conversion.h	Fixes RGB2Gray to support reduced-float outputs and avoids fallthrough.
include/fused_kernel/algorithms/basic_ops/vector_ops.h	Treats reduced-float scalars as valid scalar outputs for vector ops.
include/fused_kernel/algorithms/basic_ops/math.h	Extends Abs to reduced floats via bit-based absolute value.