Lower fp32 to soft-float calls on RV32

CHANGELOG.md

@@ -41,6 +41,12 @@ Booth — Changelog
 
 ### Tensix
 
+- lower fp32 arithmetic and conversions to soft-float calls on RV32, which has
+  no F extension. The runtime never compiled as device code before this:
+  `#ifndef __device__` erased the qualifier on every build, because it is a
+  keyword and never a macro. Float kernels still need #148
+  (Zane Hambly, 2026-07-28)
+
 - `--tt-chip` selects wormhole or blackhole, so L1 and text limits follow the
   target part instead of being fixed
   (Zane Hambly, 2026-07-23)

Makefile

@@ -58,7 +58,7 @@ endif
 
 SOURCES = src/main.c \
           src/fe/bc_err.c src/fe/bc_render.c src/fe/preproc.c src/fe/lexer.c src/fe/parser.c src/fe/sema.c \
-          src/ir/bir.c src/ir/bir_print.c src/ir/bir_lower.c src/ir/bir_mem2reg.c src/ir/bir_cfold.c src/ir/bir_dce.c src/ir/bir_struct.c src/ir/bir_insert.c src/ir/bir_sroa.c src/ir/bir_inline.c \
+          src/ir/bir.c src/ir/bir_print.c src/ir/bir_lower.c src/ir/bir_mem2reg.c src/ir/bir_cfold.c src/ir/bir_dce.c src/ir/bir_struct.c src/ir/bir_insert.c src/ir/bir_sroa.c src/ir/bir_inline.c src/ir/bir_softfp.c \
           src/tdf/tdf.c src/tdf/tdf_lower.c src/tdf/tdf_fission.c src/tdf/tdf_place.c src/tdf/tdf_noc.c \
           src/amdgpu/amd_rplan.c src/amdgpu/isel.c src/amdgpu/emit.c src/amdgpu/ra_ssa.c src/amdgpu/encode.c src/amdgpu/enc_tab.c src/amdgpu/sched.c src/amdgpu/verify.c \
           src/tensix/isel.c src/tensix/emit.c src/tensix/coarsen.c src/tensix/datamov.c src/tensix/noc.c \

runtime/soft_fp_internal.h

@@ -4,10 +4,12 @@
 #include <stdint.h>
 #include <string.h>
 
-/* Erase __device__ on the host so tests see plain C; under Booth it marks each
- * function callable from kernel code (the float lowering BIR_CALLs into them). */
-#ifndef __device__
+/* Erase __device__ on the host so tests see plain C. Keyed off the predefined
+ * macros, not __device__ itself, which is a keyword and never a macro. */
+#if !defined(__CUDACC__) && !defined(__BARRACUDA__) && !defined(__HIPCC__)
 #define __device__
+#else
+#define SFP_ON_DEVICE 1
 #endif
 
 /* Internal helpers; public API is in soft_fp.h, not for user code.
@@ -55,29 +57,37 @@ typedef struct {
 /* ---- Bit reinterpret helpers ----
  * memcpy between float and uint32_t: strict-aliasing-safe and folds to a register
  * move, unlike unions or pointer casts. */
-static inline uint32_t sfp_to_bits(float f)
+__device__ static inline uint32_t sfp_to_bits(float f)
 {
+#ifdef SFP_ON_DEVICE
+    return (uint32_t)__float_as_int(f);
+#else
     uint32_t b;
     memcpy(&b, &f, sizeof(b));
     return b;
+#endif
 }
 
-static inline float sfp_from_bits(uint32_t b)
+__device__ static inline float sfp_from_bits(uint32_t b)
 {
+#ifdef SFP_ON_DEVICE
+    return __int_as_float((int)b);
+#else
     float f;
     memcpy(&f, &b, sizeof(f));
     return f;
+#endif
 }
 
 /* ---- Unpack / pack ---- */
-sfp_unpacked_t sfp_unpack(uint32_t bits);
-uint32_t       sfp_pack  (sfp_unpacked_t u);
+__device__ sfp_unpacked_t sfp_unpack(uint32_t bits);
+__device__ uint32_t sfp_pack(sfp_unpacked_t u);
 
 /* ---- Rounding helpers ----
  * sfp_round_normal rounds a wide mantissa (24 target bits + shift_amount guard
  * bits below) to nearest-even and packs. shift_amount locates the round/sticky
  * lane, e.g. 24 for a 48-bit multiply product. */
-uint32_t sfp_round_normal(uint8_t sign, int32_t exp,
+__device__ uint32_t sfp_round_normal(uint8_t sign, int32_t exp,
                           uint64_t wide_mant, int shift_amount);
 
 #endif /* BARRACUDA_SOFT_FP_INTERNAL_H */

src/ir/bir_softfp.c

@@ -0,0 +1,95 @@
+#include "bir_softfp.h"
+#include <stdio.h>
+#include <string.h>
+
+/*
+ * bir_softfp: fp32 to soft-float calls. RV32IM has no F extension.
+ *
+ * Every op here rewrites in place, so nothing renumbers. FCMP is absent
+ * because it needs a call plus a compare, which is an insertion.
+ *
+ * I was listening to Let It Happen by Tame Impala when writing this file.
+ */
+
+#define SF_NO_FUNC 0xFFFFFFFFu
+
+/* Hello, I see we're both floating on by. */
+static const struct { uint16_t op; const char *fn; uint8_t nargs; } SFMAP[] = {
+    { BIR_FADD,   "__addsf3",      2 },
+    { BIR_FSUB,   "__subsf3",      2 },
+    { BIR_FMUL,   "__mulsf3",      2 },
+    { BIR_FDIV,   "__divsf3",      2 },
+    { BIR_SITOFP, "__floatsisf",   1 },
+    { BIR_UITOFP, "__floatunsisf", 1 },
+    { BIR_FPTOSI, "__fixsfsi",     1 },
+    { BIR_FPTOUI, "__fixunssfsi",  1 },
+};
+#define SFMAP_N (sizeof(SFMAP) / sizeof(SFMAP[0]))
+
+static uint32_t sf_find(const bir_module_t *M, const char *name)
+{
+    for (uint32_t i = 0; i < M->num_funcs; i++) {
+        if (strcmp(&M->strings[M->funcs[i].name], name) == 0) return i;
+    }
+    return SF_NO_FUNC;
+}
+
+/* A call planted inside __addsf3 recurses until the stack runs out, and a baby
+   core has very little to run out of. */
+static int sf_is_runtime(const bir_module_t *M, uint32_t fi)
+{
+    const char *n = &M->strings[M->funcs[fi].name];
+    for (uint32_t i = 0; i < SFMAP_N; i++) {
+        if (strcmp(n, SFMAP[i].fn) == 0) return 1;
+    }
+    return strncmp(n, "sfp_", 4) == 0;
+}
+
+int bir_softfp(bir_module_t *M)
+{
+    uint32_t callee[SFMAP_N];
+    for (uint32_t i = 0; i < SFMAP_N; i++) callee[i] = sf_find(M, SFMAP[i].fn);
+
+    for (uint32_t fi = 0; fi < M->num_funcs; fi++) {
+        if (sf_is_runtime(M, fi)) continue;
+        const bir_func_t *F = &M->funcs[fi];
+
+        for (uint32_t b = 0; b < F->num_blocks; b++) {
+            const bir_block_t *B = &M->blocks[F->first_block + b];
+
+            for (uint32_t k = 0; k < B->num_insts; k++) {
+                bir_inst_t *I = &M->insts[B->first_inst + k];
+
+                for (uint32_t m = 0; m < SFMAP_N; m++) {
+                    if (I->op != SFMAP[m].op) continue;
+
+                    if (callee[m] == SF_NO_FUNC) {
+                        fprintf(stderr,
+                                "bir_softfp: %s needs %s, which is not in the "
+                                "module (soft-float runtime not linked in?)\n",
+                                bir_op_name(I->op), SFMAP[m].fn);
+                        return BC_ERR_VERIFY;
+                    }
+                    if (I->num_operands != SFMAP[m].nargs) {
+                        fprintf(stderr,
+                                "bir_softfp: %s has %u operands, expected %u\n",
+                                bir_op_name(I->op), I->num_operands,
+                                SFMAP[m].nargs);
+                        return BC_ERR_VERIFY;
+                    }
+
+                    /* BIR_CALL wants the callee in operand 0. */
+                    for (uint32_t a = SFMAP[m].nargs; a > 0u; a--) {
+                        I->operands[a] = I->operands[a - 1u];
+                    }
+                    I->operands[0]   = callee[m];
+                    I->num_operands  = (uint8_t)(SFMAP[m].nargs + 1u);
+                    I->subop         = 0;
+                    I->op            = BIR_CALL;
+                    break;
+                }
+            }
+        }
+    }
+    return BC_OK;
+}

src/ir/bir_softfp.h

@@ -0,0 +1,10 @@
+#ifndef BARRACUDA_BIR_SOFTFP_H
+#define BARRACUDA_BIR_SOFTFP_H
+
+#include "bir.h"
+
+/* fp32 to libgcc-named soft-float calls, for targets with no FPU. Refuses if
+   the runtime is not already in the module. */
+int bir_softfp(bir_module_t *M);
+
+#endif /* BARRACUDA_BIR_SOFTFP_H */

src/main.c

@@ -9,6 +9,7 @@
 #include "bir_mem2reg.h"
 #include "bir_cfold.h"
 #include "bir_dce.h"
+#include "bir_softfp.h"
 #include "amdgpu.h"
 #include "sched.h"
 #include "verify.h"
@@ -113,6 +114,12 @@ static int run_bir_backends(bir_module_t *bir, const backend_cfg_t *cfg)
     if (!cfg->no_cfold)   bir_cfold(bir);
     if (!cfg->no_dce)     bir_dce(bir);
 
+    /* After cfold, so constant float folds instead of becoming a call. */
+    if (cfg->mode_rv_elf) {
+        int src = bir_softfp(bir);
+        if (src != BC_OK) return src;
+    }
+
     /* Wrap the BIR in a TDF module and lower it. For AMD and NVIDIA
      * this is a degenerate passthrough, the lowering hands the same
      * BIR pointer straight back, and the cost is one memset plus
@@ -710,6 +717,33 @@ int main(int argc, char *argv[])
     if (read_file(file, source_buf, BC_MAX_SOURCE, &src_len) != BC_OK)
         return 1;
 
+    /* One translation unit: BIR_CALL resolves a callee by index into funcs[].
+       Only when the kernel wants float, or every integer kernel pays the text. */
+    if (mode_rv_elf && strstr(source_buf, "float") != NULL) {
+        /* A quoted include resolves against the user's kernel now, not us. */
+        if (num_include_paths + 2 <= PP_MAX_INCLUDE_PATHS) {
+            include_paths[num_include_paths++] = "runtime";
+            include_paths[num_include_paths++] = "../runtime";
+        }
+        static const char *const rt[] = {
+            "runtime/soft_fp.c",
+            "../runtime/soft_fp.c",
+        };
+        uint32_t rt_len = 0;
+        uint32_t i = 0;
+        for (; i < sizeof(rt) / sizeof(rt[0]); i++) {
+            if (read_file(rt[i], source_buf + src_len,
+                          BC_MAX_SOURCE - src_len, &rt_len) == BC_OK) break;
+        }
+        if (i == sizeof(rt) / sizeof(rt[0])) {
+            fprintf(stderr,
+                    "error: --rv-elf needs runtime/soft_fp.c and could not "
+                    "find it; run from the repo root\n");
+            return 1;
+        }
+        src_len += rt_len;
+    }
+
     /* ---- TRITON NOTES -------------------------------------------------
      * The Triton frontend is a parallel input path that does not share
      * the C99 preprocessor or lexer. When --triton is on, we route the