ISA Agent

An LLM agent that learns an unseen vector DSP instruction set from its manual, writes fixed-point DSP kernels, and is checked by a real toolchain: clang/LLVM 18 for RISC-V vector, QEMU, bit-exact tests against NumPy. Plus an ONNX → MLIR → LLVM comparison path. All numbers are measured; see the repository for the raw data. GitHub

Pass rate by condition

All numbers below are generated by python -m isaagent.report from results/episodes.jsonl, results/baselines.json and results/mlir_flow.json. Nothing is typed in by hand.

Setup: 14 DSP kernels x 3 seeds per condition, local models served by Ollama on one NVIDIA L4 (temperature 0.4). Programs are compiled with clang 18 for RV64GCV and run on QEMU 8.2 (VLEN = 128). "Pass" = bit-exact (integer) or within tolerance (float) on every test case, including edge sizes and a forced-saturation case, with no writes past the output buffer. Brackets: Wilson 95% interval. Costs are dynamic instruction counts inside the kernel on QEMU, NOT cycles.

qwen2.5-coder:7b

ConditionPass (any attempt)Pass first tryFirst program compilesInvented identifiers in first programVectorised passesLLM calls / episodeOutput tokens / episodeSeconds / episodeTasks solved (any seed)
Public ISA (RVV intrinsics), no docs0/42 = 0% [0-8]0/42 = 0% [0-8]3/42 = 7% [2-19]33/42 = 79% [64-88]0/01.0368120/14
Unseen ISA, no docs11/42 = 26% [15-41]11/42 = 26% [15-41]15/42 = 36% [23-51]25/42 = 60% [44-73]0/111.025994/14
Unseen ISA, full manual in prompt16/42 = 38% [25-53]16/42 = 38% [25-53]30/42 = 71% [56-83]2/42 = 5% [1-16]16/161.017686/14
Unseen ISA, type-aware retrieval (core ops + examples)12/42 = 29% [17-44]12/42 = 29% [17-44]20/42 = 48% [33-62]2/42 = 5% [1-16]11/121.020995/14
Agent with type-aware retrieval8/42 = 19% [10-33]7/42 = 17% [8-31]16/42 = 38% [25-53]12/42 = 29% [17-44]7/85.01088435/14

qwen3-coder:30b

ConditionPass (any attempt)Pass first tryFirst program compilesInvented identifiers in first programVectorised passesLLM calls / episodeOutput tokens / episodeSeconds / episodeTasks solved (any seed)
Public ISA (RVV intrinsics), no docs2/42 = 5% [1-16]2/42 = 5% [1-16]2/42 = 5% [1-16]40/42 = 95% [84-99]0/21.0385141/14
Unseen ISA, no docs0/42 = 0% [0-8]0/42 = 0% [0-8]1/42 = 2% [0-12]41/42 = 98% [88-100]0/01.0367140/14
Unseen ISA, full manual in prompt13/42 = 31% [19-46]13/42 = 31% [19-46]31/42 = 74% [59-85]2/42 = 5% [1-16]13/131.0305155/14
Unseen ISA, retrieved manual sections0/42 = 0% [0-8]0/42 = 0% [0-8]0/42 = 0% [0-8]35/42 = 83% [69-92]0/01.0266120/14
Agent: retrieval + search tool + toolchain repair loop19/42 = 45% [31-60]1/42 = 2% [0-12]4/42 = 10% [4-22]28/42 = 67% [52-79]18/194.41902948/14
Agent + optimisation loop (instruction counts)21/42 = 50% [36-64]0/42 = 0% [0-8]0/42 = 0% [0-8]32/42 = 76% [61-87]21/215.2237712610/14
Unseen ISA, type-aware retrieval (core ops + examples)17/42 = 40% [27-56]17/42 = 40% [27-56]29/42 = 69% [54-81]11/42 = 26% [15-41]17/171.0295157/14
Agent with type-aware retrieval26/42 = 62% [47-75]15/42 = 36% [23-51]29/42 = 69% [54-81]4/42 = 10% [4-22]22/263.217678311/14
Agent with type-aware retrieval + optimisation loop24/42 = 57% [42-71]16/42 = 38% [25-53]28/42 = 67% [52-79]10/42 = 24% [13-39]24/244.5226810610/14

Per task, qwen2.5-coder:7b

Pass counts out of 3 seeds, and the best passing agent kernel's cost relative to compiler auto-vectorisation at the same register width (LMUL = 1; < 1.00 = fewer instructions than the compiler).

Taskno docsfull manualretrievalagentagent + opttyped retrievaltyped agenttyped agent + optbest agent insnsvs clang autovec (LMUL=1)vs hand-written XDSPvs scalar -O2
vadd_sat_q150/33/3---2/31/3-14101.201.000.11
dot_q150/33/3---3/33/3-12880.621.000.16
scale_q150/33/3---3/32/3-11550.691.000.11
window_q150/31/3---0/30/3-no vectorised pass---
fir_q150/30/3---0/30/3-no vectorised pass---
cmag2_q153/30/3---0/30/3-no vectorised pass---
cmul_q150/30/3---0/30/3-no vectorised pass---
argmax_abs_q150/30/3---0/30/3-no vectorised pass---
count_above_i320/33/3---0/30/3-no vectorised pass---
decim2_q152/30/3---0/30/3-no vectorised pass---
cfar_ca_i320/30/3---1/31/3-no vectorised pass---
axpy_f320/33/3---3/31/3-28181.361.000.34
rms_f323/30/3---0/30/3-no vectorised pass---
dc_remove_f323/30/3---0/30/3-no vectorised pass---

Per task, qwen3-coder:30b

Pass counts out of 3 seeds, and the best passing agent kernel's cost relative to compiler auto-vectorisation at the same register width (LMUL = 1; < 1.00 = fewer instructions than the compiler).

Taskno docsfull manualretrievalagentagent + opttyped retrievaltyped agenttyped agent + optbest agent insnsvs clang autovec (LMUL=1)vs hand-written XDSPvs scalar -O2
vadd_sat_q150/33/30/32/32/33/33/33/311570.990.820.09
dot_q150/33/30/32/33/33/33/33/311630.560.900.14
scale_q150/33/30/32/31/33/33/33/311550.691.000.11
window_q150/33/30/33/33/33/33/33/314110.681.000.11
fir_q150/30/30/30/30/31/30/31/3215130.762.210.29
cmag2_q150/30/30/32/32/30/30/30/316660.921.180.16
cmul_q150/31/30/32/33/30/32/32/325630.860.950.10
argmax_abs_q150/30/30/30/31/30/31/30/3160021.027.281.02
count_above_i320/30/30/33/32/30/33/32/317990.440.780.29
decim2_q150/30/30/30/30/30/30/30/3no vectorised pass---
cfar_ca_i320/30/30/30/30/31/31/31/3186740.511.510.35
axpy_f320/30/30/33/33/33/33/33/325651.240.910.31
rms_f320/30/30/30/30/30/31/30/320621.331.000.40
dc_remove_f320/30/30/30/31/30/33/33/316950.550.390.14

Optimisation loop

Episodes in the optimisation condition that passed: 43. The optimisation rounds found a faster correct kernel in 20 of them.

ModelTaskseedfirst correct (insns)after optimisationchange
qwen3-coder:30baxpy_f32125682566-0%
qwen3-coder:30baxpy_f32128182565-9%
qwen3-coder:30baxpy_f32233313075-8%
qwen3-coder:30baxpy_f32228182565-9%
qwen3-coder:30bcfar_ca_i3211891718674-1%
qwen3-coder:30bcmul_q15235872819-21%
qwen3-coder:30bcount_above_i32164106409-0%
qwen3-coder:30bcount_above_i32233341799-46%
qwen3-coder:30bdc_remove_f32148734365-10%
qwen3-coder:30bdot_q15012881163-10%
qwen3-coder:30bdot_q15112881163-10%
qwen3-coder:30bdot_q15112881163-10%
qwen3-coder:30bdot_q15212881163-10%
qwen3-coder:30bfir_q1512151421513-0%
qwen3-coder:30bvadd_sat_q15014101157-18%
qwen3-coder:30bvadd_sat_q15116671539-8%
qwen3-coder:30bvadd_sat_q15114101157-18%
qwen3-coder:30bvadd_sat_q15216661410-15%
qwen3-coder:30bvadd_sat_q15214101157-18%
qwen3-coder:30bwindow_q15017951540-14%

Most frequent invented identifiers (first program, XDSP conditions)

IdentifierCount
xd_wzero43
vld1q_s1619
xd_fzero16
xd_wmsb15
vint16x4_t13
saturate_int1613
vint32x4_t9
xd_hload_i168
xd_hstore_i168
xd_hzero7
v4sf6
ptrdiff_t6
vint16_t5
vx5
xd_wmadd5

Baselines (instructions for one call, bench input)

Tasknscalar C -O2clang -O3 autovecclang -O3 autovec, LMUL=1hand-written XDSP
vadd_sat_q1510241266159811741410
dot_q1510248196104220671288
scale_q1510241024685416861155
window_q15102412296105020741411
fir_q155127373519984281779735
cmag2_q1510241024390618021410
cmul_q15102425610150729792691
argmax_abs_q1510241566215662156622197
count_above_i3210246148206441132308
decim2_q1510241024371711651283
cfar_ca_i3251253313327563668912356
axpy_f3210248194104520682818
rms_f321024513091615552063
dc_remove_f32102412295169531014368

ONNX -> MLIR -> LLVM -> RISC-V (float FIR, 16 taps, 1024 outputs; axpy, 1024)

PathFIR insnsaxpy insnsnumerics vs onnxruntime
c_autovec32779-pass
c_autovec_m148138-pass
c_const_taps_autovec6863-pass
c_const_taps_autovec_m113646-pass
c_scalar128005-pass
expert_xdsp38662-pass
mlir_affine_vec4107244879max rel err 7.8e-08
mlir_affine_vec4@lmul1198114879max rel err 7.8e-08
mlir_loops+llvm_O390601310max rel err 7.8e-08
mlir_loops+llvm_O3@lmul1180192589max rel err 7.8e-08

Pick a run. Every attempt shows the program the model wrote and the feedback the tools returned.

Overview

XDSP is a vector DSP extension. A program processes arrays in chunks ("strip-mining"): each loop iteration asks the core how many lanes it may use (vl) and every vector operation takes that vl as its last argument. Lanes >= vl are not written by loads/stores. All XDSP operations are C functions declared in `xdsp.h`; include it with `#include "xdsp.h"`. Only the operations in this manual exist. The core also runs ordinary scalar C code, which is always allowed (for example loop remainders or sequential recurrences).

Data types

- `xd_h`: Vector of int16 lanes ("half" lanes). Lane count set by xd_hlen().
- `xd_w`: Vector of int32 lanes, one per xd_h lane ("wide" accumulator). Same vl as xd_h.
- `xd_s`: Vector of int32 lanes ("single" lanes). Lane count set by xd_slen().
- `xd_f`: Vector of float32 lanes. Uses the same vl as xd_s (xd_slen()).
- `xd_hmask`: Per-lane predicate produced by xd_h / xd_w comparisons.
- `xd_smask`: Per-lane predicate produced by xd_s / xd_f comparisons.

There are two lane families. The int16 family (xd_h, and xd_w which holds the int32 widening of each xd_h lane) uses vl from xd_hlen(). The int32/float family (xd_s, xd_f) uses vl from xd_slen(). Never mix a vl from one family with operations of the other.

Fixed-point conventions

Q15: an int16 x represents x / 32768. The product of two Q15 numbers is Q30 and fits in int32. Rounding operations (xd_hqmul, xd_hqmulx, xd_hsrar, xd_wnarrow, xd_havg) use round-to-nearest-up: add half an LSB, then shift arithmetically. Saturating operations clamp to [-32768, 32767] instead of wrapping. A typical Q15 dot product or FIR accumulates exact products with xd_wmac in an xd_w and converts back once with xd_wnarrow(acc, 15, vl).

Programming example: saturating vector add

```c
#include "xdsp.h"
void add_sat(const int16_t *a, const int16_t *b, int16_t *y, size_t n) {
    for (size_t vl; n > 0; n -= vl, a += vl, b += vl, y += vl) {
        vl = xd_hlen(n);
        xd_h va = xd_hload(a, vl);
        xd_h vb = xd_hload(b, vl);
        xd_hstore(y, xd_haddsat(va, vb, vl), vl);
    }
}
```

Programming example: reduction with a wide accumulator

```c
#include "xdsp.h"
int32_t sum_sq(const int16_t *x, size_t n) {
    xd_w acc = xd_wdup(0, xd_hmaxlen());      /* zero ALL lanes once */
    for (size_t vl; n > 0; n -= vl, x += vl) {
        vl = xd_hlen(n);
        xd_h v = xd_hload(x, vl);
        acc = xd_wmac(acc, v, v, vl);          /* lanes >= vl keep their partial sums */
    }
    return xd_wredsum(acc, xd_hmaxlen());       /* reduce every lane */
}
```

Operation reference

### Lane-count control

#### xd_hlen

`size_t xd_hlen(size_t n)`

Returns vl = min(n, XD_HLANES), the number of int16 lanes to process in this loop iteration. Use it for xd_h and xd_w operations. Typical strip-mining loop: `for (size_t vl; n > 0; n -= vl, p += vl) { vl = xd_hlen(n); ... }`.

#### xd_hmaxlen

`size_t xd_hmaxlen(void)`

Returns XD_HLANES, the maximum number of int16 lanes per xd_h register (8 on the reference core, but never hard-code it).

#### xd_slen

`size_t xd_slen(size_t n)`

Returns vl = min(n, XD_SLANES) for int32 (xd_s) and float32 (xd_f) operations.

#### xd_smaxlen

`size_t xd_smaxlen(void)`

Returns XD_SLANES, the maximum number of lanes per xd_s / xd_f register (4 on the reference core).

### int16 memory

#### xd_hload

`xd_h xd_hload(const int16_t *p, size_t vl)`

Loads vl consecutive int16 values from p into lanes 0..vl-1.

#### xd_hstore

`void xd_hstore(int16_t *p, xd_h v, size_t vl)`

Stores lanes 0..vl-1 of v to p[0..vl-1]. Memory beyond p[vl-1] is not touched.

#### xd_hload_step

`xd_h xd_hload_step(const int16_t *p, ptrdiff_t step, size_t vl)`

Strided load: lane i = p[i * step]. `step` is in ELEMENTS (not bytes); it may be negative.

#### xd_hload_iq

`void xd_hload_iq(const int16_t *p, xd_h *i_out, xd_h *q_out, size_t vl)`

De-interleaving load of vl complex samples stored as I,Q,I,Q,... Lane k of *i_out = p[2k], lane k of *q_out = p[2k+1]. Reads 2*vl int16 values.

#### xd_hstore_iq

`void xd_hstore_iq(int16_t *p, xd_h i_in, xd_h q_in, size_t vl)`

Interleaving store: p[2k] = lane k of i_in, p[2k+1] = lane k of q_in, for k < vl.

### int16 arithmetic

#### xd_hdup

`xd_h xd_hdup(int16_t x, size_t vl)`

Broadcasts scalar x into lanes 0..vl-1.

#### xd_hadd

`xd_h xd_hadd(xd_h a, xd_h b, size_t vl)`

Lane-wise a + b, wrapping modulo 2^16.

#### xd_hsub

`xd_h xd_hsub(xd_h a, xd_h b, size_t vl)`

Lane-wise a - b, wrapping modulo 2^16.

#### xd_haddsat

`xd_h xd_haddsat(xd_h a, xd_h b, size_t vl)`

Lane-wise saturating a + b: result clamped to [-32768, 32767].

#### xd_hsubsat

`xd_h xd_hsubsat(xd_h a, xd_h b, size_t vl)`

Lane-wise saturating a - b: result clamped to [-32768, 32767].

#### xd_havg

`xd_h xd_havg(xd_h a, xd_h b, size_t vl)`

Lane-wise rounded average (a + b + 1) >> 1, computed without intermediate overflow (arithmetic shift, exact for all int16 inputs).

#### xd_hmul

`xd_h xd_hmul(xd_h a, xd_h b, size_t vl)`

Lane-wise low 16 bits of a * b (wrapping). For fixed-point use xd_hqmul.

#### xd_hqmul

`xd_h xd_hqmul(xd_h a, xd_h b, size_t vl)`

Q15 fractional multiply with rounding and saturation: result = sat16((a * b + 16384) >> 15). Only (-32768) * (-32768) saturates (to 32767).

#### xd_hqmulx

`xd_h xd_hqmulx(xd_h a, int16_t k, size_t vl)`

Q15 multiply of every lane by scalar k: sat16((a * k + 16384) >> 15).

#### xd_hmax

`xd_h xd_hmax(xd_h a, xd_h b, size_t vl)`

Lane-wise signed maximum.

#### xd_hmin

`xd_h xd_hmin(xd_h a, xd_h b, size_t vl)`

Lane-wise signed minimum.

#### xd_habssat

`xd_h xd_habssat(xd_h a, size_t vl)`

Lane-wise saturating absolute value: |a|, with |-32768| = 32767.

#### xd_hsra

`xd_h xd_hsra(xd_h a, unsigned sh, size_t vl)`

Arithmetic shift right by sh (0..15), truncating toward minus infinity.

#### xd_hsrar

`xd_h xd_hsrar(xd_h a, unsigned sh, size_t vl)`

Rounding arithmetic shift right: (a + (1 << (sh-1))) >> sh for sh >= 1, no overflow.

#### xd_hslide

`xd_h xd_hslide(xd_h a, size_t k, size_t vl)`

Lane i of result = lane i+k of a (lanes past the register end read as 0). Rarely needed: reloading from memory at p+k is usually cheaper.

### int16 compare, select, reduce

#### xd_hgt

`xd_hmask xd_hgt(xd_h a, xd_h b, size_t vl)`

Mask lane i = (a[i] > b[i]).

#### xd_hgtx

`xd_hmask xd_hgtx(xd_h a, int16_t x, size_t vl)`

Mask lane i = (a[i] > x).

#### xd_heqx

`xd_hmask xd_heqx(xd_h a, int16_t x, size_t vl)`

Mask lane i = (a[i] == x).

#### xd_hsel

`xd_h xd_hsel(xd_hmask m, xd_h if_true, xd_h if_false, size_t vl)`

Lane i = m[i] ? if_true[i] : if_false[i]. Note the argument order: mask first.

#### xd_hcount

`size_t xd_hcount(xd_hmask m, size_t vl)`

Number of set mask lanes among lanes 0..vl-1.

#### xd_hfirst

`long xd_hfirst(xd_hmask m, size_t vl)`

Index of the lowest set mask lane among 0..vl-1, or -1 if none is set.

#### xd_hredmax

`int16_t xd_hredmax(xd_h a, size_t vl)`

Maximum over lanes 0..vl-1 (returns -32768 if vl == 0).

#### xd_hredsum_w

`int32_t xd_hredsum_w(xd_h a, size_t vl)`

Sum of lanes 0..vl-1 computed in int32 (widening, exact unless the int32 sum overflows).

### Wide accumulators (xd_w)

#### xd_wdup

`xd_w xd_wdup(int32_t x, size_t vl)`

Broadcasts x into lanes 0..vl-1 of a wide accumulator. To zero an accumulator for a strip-mined loop use xd_wdup(0, xd_hmaxlen()).

#### xd_wwiden

`xd_w xd_wwiden(xd_h a, size_t vl)`

Sign-extends each int16 lane to int32.

#### xd_wmul

`xd_w xd_wmul(xd_h a, xd_h b, size_t vl)`

Lane-wise exact product a * b as int32.

#### xd_wmulx

`xd_w xd_wmulx(xd_h a, int16_t k, size_t vl)`

Lane-wise exact product a * k as int32.

#### xd_wmac

`xd_w xd_wmac(xd_w acc, xd_h a, xd_h b, size_t vl)`

acc[i] += a[i] * b[i] (int32, wrapping) for i < vl. Lanes i >= vl keep their old value (tail-undisturbed), so one accumulator can be reused across a strip-mined loop whose last iteration has a shorter vl.

#### xd_wmacx

`xd_w xd_wmacx(xd_w acc, int16_t k, xd_h a, size_t vl)`

acc[i] += k * a[i] (int32, wrapping) for i < vl; lanes i >= vl unchanged. Note argument order: scalar coefficient before the vector.

#### xd_wadd

`xd_w xd_wadd(xd_w a, xd_w b, size_t vl)`

Lane-wise int32 a + b (wrapping).

#### xd_wsub

`xd_w xd_wsub(xd_w a, xd_w b, size_t vl)`

Lane-wise int32 a - b (wrapping).

#### xd_wsra

`xd_w xd_wsra(xd_w a, unsigned sh, size_t vl)`

Arithmetic shift right of int32 lanes by sh (0..31).

#### xd_wnarrow

`xd_h xd_wnarrow(xd_w a, unsigned sh, size_t vl)`

Narrow to int16 with rounding and saturation: sat16((a + (1 << (sh-1))) >> sh) for sh >= 1 (sh = 0: plain saturation). This is the standard way to turn a Q30 product sum back into Q15 (sh = 15).

#### xd_wredsum

`int32_t xd_wredsum(xd_w a, size_t vl)`

Sum of lanes 0..vl-1 (int32, wrapping). After a strip-mined xd_wmac loop, reduce with vl = xd_hmaxlen() so that lanes filled in earlier, longer iterations are included.

#### xd_wload

`xd_w xd_wload(const int32_t *p, size_t vl)`

Loads vl int32 values into an xd_w (vl from xd_hlen).

#### xd_wstore

`void xd_wstore(int32_t *p, xd_w v, size_t vl)`

Stores lanes 0..vl-1 of an xd_w as int32 (vl from xd_hlen).

#### xd_wgtx

`xd_hmask xd_wgtx(xd_w a, int32_t x, size_t vl)`

Mask lane i = (a[i] > x) for wide lanes.

### int32 (xd_s)

#### xd_sload

`xd_s xd_sload(const int32_t *p, size_t vl)`

Loads vl int32 values (vl from xd_slen).

#### xd_sstore

`void xd_sstore(int32_t *p, xd_s v, size_t vl)`

Stores lanes 0..vl-1.

#### xd_sdup

`xd_s xd_sdup(int32_t x, size_t vl)`

Broadcast.

#### xd_sadd

`xd_s xd_sadd(xd_s a, xd_s b, size_t vl)`

Lane-wise a + b (wrapping).

#### xd_ssub

`xd_s xd_ssub(xd_s a, xd_s b, size_t vl)`

Lane-wise a - b (wrapping).

#### xd_smul

`xd_s xd_smul(xd_s a, xd_s b, size_t vl)`

Lane-wise low 32 bits of a * b.

#### xd_smulx

`xd_s xd_smulx(xd_s a, int32_t k, size_t vl)`

Lane-wise low 32 bits of a * k.

#### xd_ssra

`xd_s xd_ssra(xd_s a, unsigned sh, size_t vl)`

Arithmetic shift right.

#### xd_smax

`xd_s xd_smax(xd_s a, xd_s b, size_t vl)`

Lane-wise maximum.

#### xd_sgt

`xd_smask xd_sgt(xd_s a, xd_s b, size_t vl)`

Mask lane i = (a[i] > b[i]).

#### xd_sgtx

`xd_smask xd_sgtx(xd_s a, int32_t x, size_t vl)`

Mask lane i = (a[i] > x).

#### xd_ssel

`xd_s xd_ssel(xd_smask m, xd_s if_true, xd_s if_false, size_t vl)`

Lane i = m[i] ? if_true[i] : if_false[i]. Mask first, like xd_hsel.

#### xd_scount

`size_t xd_scount(xd_smask m, size_t vl)`

Number of set lanes among 0..vl-1.

#### xd_sredsum

`int32_t xd_sredsum(xd_s a, size_t vl)`

Sum of lanes 0..vl-1 (int32, wrapping).

#### xd_sredmax

`int32_t xd_sredmax(xd_s a, size_t vl)`

Maximum over lanes 0..vl-1.

### float32 (xd_f)

#### xd_fload

`xd_f xd_fload(const float *p, size_t vl)`

Loads vl floats (vl from xd_slen).

#### xd_fstore

`void xd_fstore(float *p, xd_f v, size_t vl)`

Stores lanes 0..vl-1.

#### xd_fdup

`xd_f xd_fdup(float x, size_t vl)`

Broadcast.

#### xd_fadd

`xd_f xd_fadd(xd_f a, xd_f b, size_t vl)`

Lane-wise a + b.

#### xd_fsub

`xd_f xd_fsub(xd_f a, xd_f b, size_t vl)`

Lane-wise a - b.

#### xd_fmul

`xd_f xd_fmul(xd_f a, xd_f b, size_t vl)`

Lane-wise a * b.

#### xd_fmulx

`xd_f xd_fmulx(xd_f a, float k, size_t vl)`

Lane-wise a * k.

#### xd_faddx

`xd_f xd_faddx(xd_f a, float k, size_t vl)`

Lane-wise a + k.

#### xd_fmac

`xd_f xd_fmac(xd_f acc, xd_f a, xd_f b, size_t vl)`

Fused acc[i] += a[i] * b[i] for i < vl; lanes i >= vl unchanged (tail-undisturbed).

#### xd_fmacx

`xd_f xd_fmacx(xd_f acc, float k, xd_f a, size_t vl)`

Fused acc[i] += k * a[i] for i < vl; lanes i >= vl unchanged.

#### xd_fredsum

`float xd_fredsum(xd_f a, size_t vl)`

Sum of lanes 0..vl-1 in an unspecified order (results may differ from a sequential sum in the last bits).

#### xd_fredmax

`float xd_fredmax(xd_f a, size_t vl)`

Maximum over lanes 0..vl-1.

The question

Can an LLM agent write correct, fast kernels for a DSP whose instruction set exists only in an internal manual? Public models have never seen such an ISA, so everything must come from the documentation and from tool feedback.

The stand-in ISA

XDSP is a fictional vector DSP with 73 documented operations (Q15 fixed-point multiply with rounding and saturation, widening multiply-accumulate, de-interleaving complex loads, masks, reductions, float FMA). Each operation is a thin wrapper over one RISC-V Vector 1.0 intrinsic, so programs run bit-exactly on QEMU, but the names, types and manual are new, so a model cannot have memorised them.

Conditions

Public ISA without docs; XDSP without docs; XDSP with the full manual in the prompt; XDSP with retrieved manual sections (BM25); the agent (retrieval, a SEARCH tool, and up to four repair rounds driven by compiler errors, test failures and the manual entries of the operations involved); the agent plus two optimisation rounds driven by measured instruction counts.

Verification

Every program is compiled with clang 18 for RV64GCV and run on QEMU 8.2 (VLEN 128). Tests cover edge sizes, a forced-saturation case and writes past the output buffer. Cost is the dynamic instruction count inside the kernel (from QEMU's execution trace), not cycles. Baselines: scalar C at -O2, clang auto-vectorisation at -O3 (also at the same register width, LMUL = 1), and hand-written XDSP kernels.

Limits

14 kernels and 3 seeds per condition is small; instruction counts are not cycle counts; the ISA is a renamed RISC-V subset, so it is easier than a real VLIW DSP with slot and latency constraints.