Lower fp32 to soft-float calls on RV32

[Zaneham]

Draft. The pass works, the runtime lowers, and float kernels still do not
compile end to end because of #148.

The Tensix baby cores are RV32IM with no F extension, so float has to become
calls into runtime/soft_fp.c. That file has been sitting in the tree
host-tested and never actually wired to anything, and rv_isel said as much:

Any subsequent BIR_FADD or other float arithmetic will refuse, because we
have no soft-float runtime yet

So this is the wiring.

bir_softfp rewrites BIR_FADD and friends into BIR_CALL to the
libgcc-named entry points. Every op it handles is a one-for-one rewrite in
place, so nothing renumbers. It runs after cfold, so constant float folds away
rather than turning into a call nobody needed. FCMP is deliberately absent: it
needs a call plus a compare against zero, which is an insertion rather than a
rewrite, and it refuses in isel until that lands.

The driver appends runtime/soft_fp.c to the translation unit when the kernel
mentions float, so BIR_CALL can resolve a callee by index into funcs[].
Gated on float use, or every integer kernel pays the text budget for arithmetic
it never does.

Three things were wrong in the runtime and none of them had ever been noticed,
because nothing had compiled it as device code:

• #ifndef __device__ erased the qualifier on every build. It is a keyword
  and never a macro, so the condition is always true. The whole runtime was
  being compiled as host code and skipped by lowering.
• The bit-reinterpret helpers and the header declarations needed the qualifier
  too.
• memcpy is not a device builtin, so the float/int punning now uses
  __float_as_int and __int_as_float, which Booth already had.

What is left is #148. sfp_unpacked_t has a _pad[] member, and indexing an
array member of a struct produces a GEP with a non-pointer result type. That is
a general frontend bug with a seven-line repro, nothing to do with soft-float,
and it blocks this the moment the runtime actually runs.

Suite is green at 358. I have deliberately not included two speculative fixes
to bir_lower that looked right and changed the output by nothing, since an
unverified edit to lowering is worse than no edit.