From 6284b6374d21b63cc1da9a6dec8454781ea6015f Mon Sep 17 00:00:00 2001 From: ZaneHam Date: Tue, 28 Jul 2026 22:11:32 +1200 Subject: [PATCH] feat: lower fp32 to soft-float calls on RV32 --- CHANGELOG.md | 6 +++ Makefile | 2 +- runtime/soft_fp_internal.h | 26 +++++++---- src/ir/bir_softfp.c | 95 ++++++++++++++++++++++++++++++++++++++ src/ir/bir_softfp.h | 10 ++++ src/main.c | 34 ++++++++++++++ 6 files changed, 164 insertions(+), 9 deletions(-) create mode 100644 src/ir/bir_softfp.c create mode 100644 src/ir/bir_softfp.h diff --git a/CHANGELOG.md b/CHANGELOG.md index a187cdf..279326f 100644 --- a/CHANGELOG.md +++ b/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) diff --git a/Makefile b/Makefile index 83b0528..2ea75ec 100644 --- a/Makefile +++ b/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 \ diff --git a/runtime/soft_fp_internal.h b/runtime/soft_fp_internal.h index 58a21f7..a9321dc 100644 --- a/runtime/soft_fp_internal.h +++ b/runtime/soft_fp_internal.h @@ -4,10 +4,12 @@ #include #include -/* 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 */ diff --git a/src/ir/bir_softfp.c b/src/ir/bir_softfp.c new file mode 100644 index 0000000..31732cd --- /dev/null +++ b/src/ir/bir_softfp.c @@ -0,0 +1,95 @@ +#include "bir_softfp.h" +#include +#include + +/* + * 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; +} diff --git a/src/ir/bir_softfp.h b/src/ir/bir_softfp.h new file mode 100644 index 0000000..83d6ea9 --- /dev/null +++ b/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 */ diff --git a/src/main.c b/src/main.c index 91c3639..cd5d13f 100644 --- a/src/main.c +++ b/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