prototype: port teeny's portable half/bfloat16 to C++11 (core/half.h)

[balbasty]

Answers the question "should we steal teeny's portable half/bf16, and would it work on C++11?"

Short version: yes — the header ports cleanly and is done here, verified on clang++, g++ and nvcc. It is a header and nothing more. Nothing includes it, no dtype dispatch gained a 16-bit arm, and the two things that actually block half/bf16 support — the CUDA atomics of #100 and the instantiation budget — are untouched and still open.

Brought up to current main (247714e, through #155) by merging rather than rebasing, so the push stays a fast-forward and no one's checkout of this branch is rewritten. One conflict, in the top-level Makefile: main added a test-atomics target to the same .PHONY line this branch adds test-half to. Both are kept, and nothing else in the branch collided.

Scope — what this is not

This is worth being blunt about, because the analysis below is long enough to read like a plan:

• It does not wire half or bf16 into the dispatch pyramid. No case kDLFloat/16 arm was added anywhere. The dispatch layer still throws on a 16-bit tensor, and the ten tests/lib-cpu/test_*.cpp suites that assert float16 must throw still pass unchanged.
• Half/bf16 CUDA atomics will need CAS fallbacks — and a 4-byte alignment precondition on user DLPack pointers #100 is not addressed. Half/bf16 CUDA atomics need CAS fallbacks below sm_70 (half) and sm_80 (bf16), and — the part that is easy to miss — a 4-byte alignment precondition on user-supplied DLPack pointers, because a 2-byte CAS reads and writes the enclosing aligned word. anyAtomicAdd(half*, half) is still a hard compile error in this branch, deliberately.
• The instantiation-budget problem is not addressed. Adding a dtype is a clean multiplier over the whole dispatch layer, and the CUDA build has no headroom to spend at two dtypes, let alone four. perf(cuda): split reg_flow into per-(family, ndim) translation units #147 measures and cuts the worst single translation unit; that is a prerequisite, not a solution.

What it is: a self-contained header, an exhaustive conversion test, two nvcc compile probes, and one new make target. make test is untouched.

1. C++11: yes, definitively

Verified under clang++ and g++, -std=c++11 and -std=c++14, -Wall -Wextra -Wpedantic clean, and under nvcc via make test-impl-cuda. Every blocker in the original reading was real and shallow:

• <cuda/std/type_traits> / <cuda/std/cstdint> → <type_traits> plus the <cstdint> that core/cuda_switch.h already provides. libcu++ (CCCL) enforces a C++17 floor; this repo is C++11 (C++14 for nvcc) and has no CCCL dependency.
• cs::enable_if_t<...> → typename std::enable_if<...>::type. Trivial.
• _TNY_API → FF_CUHOSTDEV.
• The union pun is UB in ISO C++ (defined in C; a documented extension on gcc/clang/msvc/nvcc). Replaced with memcpy, which those compilers fold to the same register move. Free correctness, taken.

One correction to the original reading. uint16_t bits{} is C++11-legal, but the NSDMI is not what costs aggregate-ness — the user-provided templated converting constructor already does that. What the NSDMI actually costs is a trivial default constructor, which means __shared__ half tile[N] will not compile under nvcc and half x; silently zero-stores. Dropped it; trivial default construction and trivial copyability are both restored and both static_asserted. compile_probe_half_portable.cu compiles a __shared__ ff::half tile[32] under nvcc precisely so this cannot regress.

2. What actually bites (found by compiling, not by reading)

The header ports mechanically. Pointing it at the real kernels does not. These were found while prototyping a half arm into the dispatch — that arm is not part of this PR, but what it turned up is the reason this header alone is not "half support":

• reduce_t OP scalar_t is ambiguous. The implicit converting constructor makes operator*(half, half) viable and the implicit operator float() makes the built-in viable; neither wins. Fixed entirely inside the header with heterogeneous overloads (FF_HALF_MIXED), as PyTorch's c10::Half does — no kernel edited. Making the constructor explicit instead trades this class of error for a different one (no viable conversion from double to half at every scalar_t x = <literal>), so it is not the answer.
• A latent bug in bounds.h.cget casts to val_t and then returns scalar_t, discarding the widening the call site asked for. Harmless for the current dtypes (float→double→float is exact), but with scalar_t = half the value round-trips through 16 bits. Not fixed here — it touches shared bound helpers — but it is a real bug on its own merits.
• ff::is_floating_point is false on the host (it is only specialised under __CUDACC__), which routes half into _mod's integer specialisation and its x % d — a hard compile error.
• std::numeric_limits is unspecialised, so max(), lowest() and infinity() all silently return 0.0 with no diagnostic.

The last three are silent-wrong-answer risks the moment half appears anywhere, and each is small and independently fixable.

3. Verification

• make test-half — exhaustive. All 2^16 half and all 2^16 bf16 patterns decoded, and all 2^32 float patterns encoded in each direction, bit-compared against the compiler's own _Float16 / __bf16. Every subnormal, every overflow boundary and every round-to-nearest-even tie is covered by construction rather than by a hand-picked list. NaN is exempted (software canonicalises the payload; hardware preserves it). make test-half STRIDE=N subsamples for a smoke run. Measured on the merged tree: 8.6 billion comparisons, clean under clang++ -std=c++11 (415 s) and g++ -std=c++14 (380 s).
  The oracle earned its keep: one named expectation in the test was wrong and it caught it — FLT_MAX rounds up to bf16 +inf, and bf16's largest finite is 0x7f7f. The header was right.
• nvcc:compile_probe_half.cu (native path — the __half/__nv_bfloat16 aliases and templated device kernels) and compile_probe_half_portable.cu (the software structs compiled as device code under FF_PORTABLE_HALF, including the __shared__ array). Both are picked up automatically by the existing make test-impl-cuda glob.
• The gate does not move.tools/test-baseline.sh --legs default,lib on this branch, diffed row for row against tools/test-baseline.expected: 15 rows, 0 differ — 13 suites / 59,886 checks / 0 failures on the default leg, plus the 2 hub suites (14 checks). All five test-cpu legs and test-hub are green in CI besides. Nothing in this PR is reachable from any of them: core/half.h is included by exactly two files, both new, neither built by make test.
  (--check itself refuses a two-leg run — it compares whole reports and the recording covers six legs — so the fifteen rows were diffed against the recording directly.)
• No GPU in CI, so nothing here is a claim about runtime behaviour on a device — only that it compiles as device code and that the host conversions are bit-exact against the native types.

4. Conventions

Brought up to what the tree now requires, which is the only substantive change since the last review:

• #pragma once on line 1, with the whole-file #ifndef FF_HALF guard removed (refactor: #pragma once in every header #145). The file had acquired both, which is the one combination the convention rules out.
• FF_NAMESPACE_BEGIN(FF_NS) / FF_NAMESPACE_END(FF_NS) — FF was renamed FF_NS in refactor: prefix every remaining public macro with FF_ #91.
• FF_CUHOSTDEV, not the old unprefixed CUHOSTDEV (refactor: prefix every remaining public macro with FF_ #91).
• <fastfields/…> for the public interface, in the header and in all three test/probe files (refactor: spell the public interface <fastfields/...> #146).

normalise-header-guards.py, normalise-include-delimiters.py, rename-macros.py and check-cuda-launches.py all pass --check.

5. Licence and provenance

balbasty/teeny is MIT, © 2023-2025 Yaël Balbastre; this repository is MIT, © 2025 Yaël Balbastre. Same licence, same holder — no new obligation. The header carries a provenance block naming teeny, the licence, and every deviation made for fastfields. No LICENSE change is needed.

6. Recommendation

Take the header; do not read it as a step towards half support that is now underway. It is cheap to carry, it retires the "would it even work on C++11" question with evidence, and it is inert. The work that would make half real — #100's atomics and alignment precondition, the instantiation budget, a GPU in CI, and a deliberate rewrite of the ten float16-must-throw cases — is all still ahead and none of it is started.

lint (clang-format, changed lines) is green on this branch. The lint (source conventions) job #156 has since added did not exist when this run started, so it has not run here; the three checkers it invokes were run against this tree by hand and pass (section 4).

[balbasty]

ok