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

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
| Condition | Pass (any attempt) | Pass first try | First program compiles | Invented identifiers in first program | Vectorised passes | LLM calls / episode | Output tokens / episode | Seconds / episode | Tasks solved (any seed) |
|---|---|---|---|---|---|---|---|---|---|
| Public ISA (RVV intrinsics), no docs | 0/42 = 0% [0-8] | 0/42 = 0% [0-8] | 3/42 = 7% [2-19] | 33/42 = 79% [64-88] | 0/0 | 1.0 | 368 | 12 | 0/14 |
| Unseen ISA, no docs | 11/42 = 26% [15-41] | 11/42 = 26% [15-41] | 15/42 = 36% [23-51] | 25/42 = 60% [44-73] | 0/11 | 1.0 | 259 | 9 | 4/14 |
| Unseen ISA, full manual in prompt | 16/42 = 38% [25-53] | 16/42 = 38% [25-53] | 30/42 = 71% [56-83] | 2/42 = 5% [1-16] | 16/16 | 1.0 | 176 | 8 | 6/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/12 | 1.0 | 209 | 9 | 5/14 |
| Agent with type-aware retrieval | 8/42 = 19% [10-33] | 7/42 = 17% [8-31] | 16/42 = 38% [25-53] | 12/42 = 29% [17-44] | 7/8 | 5.0 | 1088 | 43 | 5/14 |
qwen3-coder:30b
| Condition | Pass (any attempt) | Pass first try | First program compiles | Invented identifiers in first program | Vectorised passes | LLM calls / episode | Output tokens / episode | Seconds / episode | Tasks solved (any seed) |
|---|---|---|---|---|---|---|---|---|---|
| Public ISA (RVV intrinsics), no docs | 2/42 = 5% [1-16] | 2/42 = 5% [1-16] | 2/42 = 5% [1-16] | 40/42 = 95% [84-99] | 0/2 | 1.0 | 385 | 14 | 1/14 |
| Unseen ISA, no docs | 0/42 = 0% [0-8] | 0/42 = 0% [0-8] | 1/42 = 2% [0-12] | 41/42 = 98% [88-100] | 0/0 | 1.0 | 367 | 14 | 0/14 |
| Unseen ISA, full manual in prompt | 13/42 = 31% [19-46] | 13/42 = 31% [19-46] | 31/42 = 74% [59-85] | 2/42 = 5% [1-16] | 13/13 | 1.0 | 305 | 15 | 5/14 |
| Unseen ISA, retrieved manual sections | 0/42 = 0% [0-8] | 0/42 = 0% [0-8] | 0/42 = 0% [0-8] | 35/42 = 83% [69-92] | 0/0 | 1.0 | 266 | 12 | 0/14 |
| Agent: retrieval + search tool + toolchain repair loop | 19/42 = 45% [31-60] | 1/42 = 2% [0-12] | 4/42 = 10% [4-22] | 28/42 = 67% [52-79] | 18/19 | 4.4 | 1902 | 94 | 8/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/21 | 5.2 | 2377 | 126 | 10/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/17 | 1.0 | 295 | 15 | 7/14 |
| Agent with type-aware retrieval | 26/42 = 62% [47-75] | 15/42 = 36% [23-51] | 29/42 = 69% [54-81] | 4/42 = 10% [4-22] | 22/26 | 3.2 | 1767 | 83 | 11/14 |
| Agent with type-aware retrieval + optimisation loop | 24/42 = 57% [42-71] | 16/42 = 38% [25-53] | 28/42 = 67% [52-79] | 10/42 = 24% [13-39] | 24/24 | 4.5 | 2268 | 106 | 10/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).
| Task | no docs | full manual | retrieval | agent | agent + opt | typed retrieval | typed agent | typed agent + opt | best agent insns | vs clang autovec (LMUL=1) | vs hand-written XDSP | vs scalar -O2 |
|---|---|---|---|---|---|---|---|---|---|---|---|---|
| vadd_sat_q15 | 0/3 | 3/3 | - | - | - | 2/3 | 1/3 | - | 1410 | 1.20 | 1.00 | 0.11 |
| dot_q15 | 0/3 | 3/3 | - | - | - | 3/3 | 3/3 | - | 1288 | 0.62 | 1.00 | 0.16 |
| scale_q15 | 0/3 | 3/3 | - | - | - | 3/3 | 2/3 | - | 1155 | 0.69 | 1.00 | 0.11 |
| window_q15 | 0/3 | 1/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| fir_q15 | 0/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| cmag2_q15 | 3/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| cmul_q15 | 0/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| argmax_abs_q15 | 0/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| count_above_i32 | 0/3 | 3/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| decim2_q15 | 2/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| cfar_ca_i32 | 0/3 | 0/3 | - | - | - | 1/3 | 1/3 | - | no vectorised pass | - | - | - |
| axpy_f32 | 0/3 | 3/3 | - | - | - | 3/3 | 1/3 | - | 2818 | 1.36 | 1.00 | 0.34 |
| rms_f32 | 3/3 | 0/3 | - | - | - | 0/3 | 0/3 | - | no vectorised pass | - | - | - |
| dc_remove_f32 | 3/3 | 0/3 | - | - | - | 0/3 | 0/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).
| Task | no docs | full manual | retrieval | agent | agent + opt | typed retrieval | typed agent | typed agent + opt | best agent insns | vs clang autovec (LMUL=1) | vs hand-written XDSP | vs scalar -O2 |
|---|---|---|---|---|---|---|---|---|---|---|---|---|
| vadd_sat_q15 | 0/3 | 3/3 | 0/3 | 2/3 | 2/3 | 3/3 | 3/3 | 3/3 | 1157 | 0.99 | 0.82 | 0.09 |
| dot_q15 | 0/3 | 3/3 | 0/3 | 2/3 | 3/3 | 3/3 | 3/3 | 3/3 | 1163 | 0.56 | 0.90 | 0.14 |
| scale_q15 | 0/3 | 3/3 | 0/3 | 2/3 | 1/3 | 3/3 | 3/3 | 3/3 | 1155 | 0.69 | 1.00 | 0.11 |
| window_q15 | 0/3 | 3/3 | 0/3 | 3/3 | 3/3 | 3/3 | 3/3 | 3/3 | 1411 | 0.68 | 1.00 | 0.11 |
| fir_q15 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 1/3 | 0/3 | 1/3 | 21513 | 0.76 | 2.21 | 0.29 |
| cmag2_q15 | 0/3 | 0/3 | 0/3 | 2/3 | 2/3 | 0/3 | 0/3 | 0/3 | 1666 | 0.92 | 1.18 | 0.16 |
| cmul_q15 | 0/3 | 1/3 | 0/3 | 2/3 | 3/3 | 0/3 | 2/3 | 2/3 | 2563 | 0.86 | 0.95 | 0.10 |
| argmax_abs_q15 | 0/3 | 0/3 | 0/3 | 0/3 | 1/3 | 0/3 | 1/3 | 0/3 | 16002 | 1.02 | 7.28 | 1.02 |
| count_above_i32 | 0/3 | 0/3 | 0/3 | 3/3 | 2/3 | 0/3 | 3/3 | 2/3 | 1799 | 0.44 | 0.78 | 0.29 |
| decim2_q15 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | no vectorised pass | - | - | - |
| cfar_ca_i32 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 1/3 | 1/3 | 1/3 | 18674 | 0.51 | 1.51 | 0.35 |
| axpy_f32 | 0/3 | 0/3 | 0/3 | 3/3 | 3/3 | 3/3 | 3/3 | 3/3 | 2565 | 1.24 | 0.91 | 0.31 |
| rms_f32 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 0/3 | 1/3 | 0/3 | 2062 | 1.33 | 1.00 | 0.40 |
| dc_remove_f32 | 0/3 | 0/3 | 0/3 | 0/3 | 1/3 | 0/3 | 3/3 | 3/3 | 1695 | 0.55 | 0.39 | 0.14 |
Optimisation loop
Episodes in the optimisation condition that passed: 43. The optimisation rounds found a faster correct kernel in 20 of them.
| Model | Task | seed | first correct (insns) | after optimisation | change |
|---|---|---|---|---|---|
| qwen3-coder:30b | axpy_f32 | 1 | 2568 | 2566 | -0% |
| qwen3-coder:30b | axpy_f32 | 1 | 2818 | 2565 | -9% |
| qwen3-coder:30b | axpy_f32 | 2 | 3331 | 3075 | -8% |
| qwen3-coder:30b | axpy_f32 | 2 | 2818 | 2565 | -9% |
| qwen3-coder:30b | cfar_ca_i32 | 1 | 18917 | 18674 | -1% |
| qwen3-coder:30b | cmul_q15 | 2 | 3587 | 2819 | -21% |
| qwen3-coder:30b | count_above_i32 | 1 | 6410 | 6409 | -0% |
| qwen3-coder:30b | count_above_i32 | 2 | 3334 | 1799 | -46% |
| qwen3-coder:30b | dc_remove_f32 | 1 | 4873 | 4365 | -10% |
| qwen3-coder:30b | dot_q15 | 0 | 1288 | 1163 | -10% |
| qwen3-coder:30b | dot_q15 | 1 | 1288 | 1163 | -10% |
| qwen3-coder:30b | dot_q15 | 1 | 1288 | 1163 | -10% |
| qwen3-coder:30b | dot_q15 | 2 | 1288 | 1163 | -10% |
| qwen3-coder:30b | fir_q15 | 1 | 21514 | 21513 | -0% |
| qwen3-coder:30b | vadd_sat_q15 | 0 | 1410 | 1157 | -18% |
| qwen3-coder:30b | vadd_sat_q15 | 1 | 1667 | 1539 | -8% |
| qwen3-coder:30b | vadd_sat_q15 | 1 | 1410 | 1157 | -18% |
| qwen3-coder:30b | vadd_sat_q15 | 2 | 1666 | 1410 | -15% |
| qwen3-coder:30b | vadd_sat_q15 | 2 | 1410 | 1157 | -18% |
| qwen3-coder:30b | window_q15 | 0 | 1795 | 1540 | -14% |
Most frequent invented identifiers (first program, XDSP conditions)
| Identifier | Count |
|---|---|
xd_wzero | 43 |
vld1q_s16 | 19 |
xd_fzero | 16 |
xd_wmsb | 15 |
vint16x4_t | 13 |
saturate_int16 | 13 |
vint32x4_t | 9 |
xd_hload_i16 | 8 |
xd_hstore_i16 | 8 |
xd_hzero | 7 |
v4sf | 6 |
ptrdiff_t | 6 |
vint16_t | 5 |
vx | 5 |
xd_wmadd | 5 |
Baselines (instructions for one call, bench input)
| Task | n | scalar C -O2 | clang -O3 autovec | clang -O3 autovec, LMUL=1 | hand-written XDSP |
|---|---|---|---|---|---|
| vadd_sat_q15 | 1024 | 12661 | 598 | 1174 | 1410 |
| dot_q15 | 1024 | 8196 | 1042 | 2067 | 1288 |
| scale_q15 | 1024 | 10246 | 854 | 1686 | 1155 |
| window_q15 | 1024 | 12296 | 1050 | 2074 | 1411 |
| fir_q15 | 512 | 73735 | 19984 | 28177 | 9735 |
| cmag2_q15 | 1024 | 10243 | 906 | 1802 | 1410 |
| cmul_q15 | 1024 | 25610 | 1507 | 2979 | 2691 |
| argmax_abs_q15 | 1024 | 15662 | 15662 | 15662 | 2197 |
| count_above_i32 | 1024 | 6148 | 2064 | 4113 | 2308 |
| decim2_q15 | 1024 | 10243 | 717 | 1165 | 1283 |
| cfar_ca_i32 | 512 | 53313 | 32756 | 36689 | 12356 |
| axpy_f32 | 1024 | 8194 | 1045 | 2068 | 2818 |
| rms_f32 | 1024 | 5130 | 916 | 1555 | 2063 |
| dc_remove_f32 | 1024 | 12295 | 1695 | 3101 | 4368 |
ONNX -> MLIR -> LLVM -> RISC-V (float FIR, 16 taps, 1024 outputs; axpy, 1024)
| Path | FIR insns | axpy insns | numerics vs onnxruntime |
|---|---|---|---|
| c_autovec | 32779 | - | pass |
| c_autovec_m1 | 48138 | - | pass |
| c_const_taps_autovec | 6863 | - | pass |
| c_const_taps_autovec_m1 | 13646 | - | pass |
| c_scalar | 128005 | - | pass |
| expert_xdsp | 38662 | - | pass |
| mlir_affine_vec4 | 10724 | 4879 | max rel err 7.8e-08 |
| mlir_affine_vec4@lmul1 | 19811 | 4879 | max rel err 7.8e-08 |
| mlir_loops+llvm_O3 | 9060 | 1310 | max rel err 7.8e-08 |
| mlir_loops+llvm_O3@lmul1 | 18019 | 2589 | max 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.