Nine pages under LLM/ document the upstream radiance engine at 8da043c: - engine shape, plugin model, RAD2 container, ABI and environment - every serve flag grouped with defaults and the law behind each knob - the HTTP surface, request fields and refusal classes - rad-convert, the recipe language, quantisers, imports, reuse, rad-export, KL scoring - the (arch id, quant) family table and the recipe of record per model - kernel/quantiser plugin authoring, encodings, build and checks - architecture plugin authoring: declare, step, probe, lane and media rules - rad-info/radbench/radtune and the shell harnesses - the container contract, host prerequisites, readiness and restart triage Also completes the dangling Related line on the frank-archon endpoint page.
20 KiB
type, title, description, tags, timestamp
| type | title | description | tags | timestamp | ||||
|---|---|---|---|---|---|---|---|---|
| howto | Radiance Plugins — kernel libraries and quantisers | Writing radiance kernel and quantiser plugins: the ABI, kernel rows, hooks, encodings, and the build path. |
|
2026-10-30 |
Radiance plugins — kernels and quantisers
Three kinds of .so extend the engine, each exporting a few C symbols and linking nothing of
radiance's (they compile against the installed abi/ headers):
| kind | value | contract |
|---|---|---|
| kernel library | RAD_PLUGIN_KERNEL (1) |
ops for some hardware: a card family, the CPU, an oracle |
| architecture | RAD_PLUGIN_ARCH (2) |
one model family at one weight format — see Radiance Architecture Plugins |
| quantiser | RAD_PLUGIN_QUANT (3) |
weights → encoding planes, for rad-convert |
In-tree plugins use the same interface a third party does: each builds with rad_add_plugin and is
found and ranked by what it declares, never by its name. ctest -L oot rebuilds every bundled
plugin out of tree against an install.
Identity and ABI
#include "rad_abi.h" /* required; pulls rad_types.h and rad_encoding.h */
uint32_t rad_plugin_abi_version(void) { return RAD_ABI_VERSION; }
const RadPluginInfo* rad_plugin_info(void) {
static const RadPluginInfo i = {
RAD_PLUGIN_KERNEL, "mykernels", "0.1", "a GEMM or two", "gfx1100",
};
return &i;
}
RAD_ABI_VERSION(31 at this commit) is bumped on any shape change in a header a plugin compiles against. Return the macro; a mismatch is refused at load with both numbers printed.build_targetis printed, not trusted: the targets the library was built for ("gfx1100","gfx11-generic", several separated by spaces/commas),"host", or""for portable. The engine checks the library's actual code objects instead (§"A library for a new card").- Optional exports, resolved by exact name:
rad_plugin_open(returnRAD_E_UNSUPPORTEDto decline this machine and unload),rad_plugin_close,rad_kernel_concurrent,rad_kernel_collective,rad_plugin_reference,rad_kernel_reference,rad_kernel_chain.
What a kernel library exports
int rad_kernel_schema_count(void);
const RadOpSchema* rad_kernel_schema_at(int i);
int rad_kernel_count(void);
const RadKernelInfo* rad_kernel_at(int i);
A row declares one op implementation and what it can serve:
static const RadConstraint c_gemm[] = {
RAD_CLE("M", 64), /* M <= 64 */
RAD_CDIV("K", 128), /* K divisible by 128 */
RAD_CIN("dtype", "fp8a8 w4a8") /* dtype in that set */
};
static const RadKernelInfo k[] = {{
.name = "mygemm_m64", .op = "gemm_nt_q", .family = "gemm",
.computes = "row-major A times column-major B, fp8 in, bf16 out",
.shape = "M<=64, K%128==0",
.dtypes = "fp8a8, w4a8",
.domain = RAD_DOMAIN_DEVICE,
.priority = 50,
.constraints = c_gemm, .n_constraints = 3,
.launch = my_launch,
}};
- Operators:
RAD_CEQ,RAD_CLE,RAD_CGE,RAD_CDIV,RAD_CIN. The engine takes bucket boundaries fromLE/GE/EQand resolves once per band;DIVandINare not intervals. - Declare only the band you serve. The engine believes the claim; a kernel that says
M <= 2048but is correct to 64 is wrong at concurrency. device.*keys are matched against what the engine measured about the ranks' cards, not against op parameters:RAD_CEQ(RAD_FACT_P2P, 1)means the kernel touches another card's memory. libr4d pairs every collective with a_hosttwin requiringdevice.p2p == 0, so exactly one matches any machine.- Priority breaks ties within one plugin, highest first, then declaration order. Between
plugins the hierarchy is absolute: the first library in
--kernelsorder with any matching row supplies the op, however specialised a later library's row is. That makes a user's library an override, not a competitor. - A derivative parameter (
RAD_DERIVED) is supplied by the kernel; a constraint on a key the geometry omits does not hold during matching, and the winner'sRAD_C_EQwrites it into the resolved geometry.RAD_PROLE_SEQ_CHUNKmarks the parameter the scheduler must cut steps on;RAD_PROLE_CAPACITYmarks a count the kernel only bounds itself by.
Operands and extents
RadArgs.t is positional in the schema's order; each RadTensor carries the extent the caller
narrowed it to, usually smaller than the buffer. Size loops and grids from those extents, never from
a declared maximum. Where an op takes cu_seqlens, derive the sequence count from its own length
(n = len - 1) and each sequence's rows from cu[n] and cu[n+1].
An absent optional operand is a null RadTensor.data, not a shorter operand count.
rad_arg_in(a, i) fetches one correctly; rad_tensor_rows, rad_tensor_lastdim,
rad_tensor_row_pitch, rad_args_geti_or, rad_args_getf_or, rad_args_gets_or are the other
everyday helpers in rad_plugin.h.
Optional hooks
| hook | meaning |
|---|---|
init / fini |
per-resolved-instance state; RadArgs.instance hands it back. init runs once per (op, band, domain) after selection and before load |
scratch |
bytes of scratch one launch needs; sized at declare from the band's upper bound |
layout / relayout |
the kernel stores this weight in its own order, and how the loader produces it from the canonical planes |
unrelayout |
the inverse, so a tool holding stored bytes can compare against a dense oracle |
opd_shape |
size and fill every operand, so a tool can call the kernel cold |
replaces |
the chain of conventional ops this kernel computes in one launch |
tunables / tune_valid |
the instantiation axes, and which combinations are legal at a shape |
replaces must be exact and is compared byte for byte against the chain — a kernel that is only
close publishes nothing. tunables declares axes and legal values, not named variants; the core
takes the cross product, and deflt must be launchable (the loader refuses the plugin otherwise).
rad-tune measures survivors and writes the cache the selector reads; the chosen point arrives as
ordinary parameters (rad_args_geti("sk")), with no variant index.
layout answers: RAD_OK (own form, relayout will produce it), RAD_E_UNSUPPORTED (stored as
is — the common case, and the only form that can be read in place from the file/mapped tier),
RAD_E_DTYPE (this encoding is not mine — a declare-time error), RAD_E_SHAPE, or RAD_LAYOUT_ANY
(reads the bytes without interpreting them; claims no storage). A weight only LAYOUT_ANY readers
touch is stored as it is.
What a launch may do
On a card the RadStream is the engine's own AQL queue, not a hipStream_t. The engine binary
defines and serves these HIP entry points:
hipLaunchKernel(and<<<>>>,hipLaunchKernelGGL);hipMemcpyAsync,hipMemcpy2DAsync,hipMemsetAsync;hipStreamSynchronize(not from a launch — a launch must not synchronise);hipMemset(runs on the engine's queue and waits);hipGetLastError/hipPeekAtLastError, with HIP's semantics.
Nothing else in HIP takes an engine stream. hipModuleLaunchKernel, hipExtLaunchKernel,
cooperative launches, hipLaunchHostFunc, and libraries that launch through them (rocBLAS,
hipBLASLt, MIOpen, hipRTC) would hand HIP a queue it does not know. Header-only device code
(Composable Kernel templates, rocPRIM primitives) compiled into your .hip works. hipEventRecord
and hipStreamWaitEvent on an engine stream are refused by name.
The engine also reads the code object, which limits a kernel to: no hostcall (printf/assert
are refused at the first launch), no dynamic stack, ≤ 768 bytes of kernel arguments, ≤ 64
KiB static+dynamic LDS per dispatch, and an uncompressed fat binary. Device memory a kernel
needs comes from the scratch hook, never hipMalloc — the engine's allocator knows what is left.
rad_plugin.h's conversions define the result: the reference contract is f32 accumulation, IEEE
round-to-nearest-even, no FMA and no reassociation. A plugin whose bf16 narrowing truncates
cannot be checked against the reference. rad_load_f32/rad_store_f32/rad_load_bits/
rad_store_bits read and write any core dtype from a plane.
Colours of agreement
rad_kernel_concurrent(i)— kernelitouches only its operands and scratch, so independent launches may overlap.rad_kernel_collective(i)— it exchanges data with other ranks; the engine keeps bulk link traffic away.2promises more: the launch ends the hold itself once the other ranks' bytes have landed, through alink_gateword it reaches at system scope; returningRAD_LINK_ENDEDsays it did.rad_kernel_chain(i)— the kernel can wait in its own prologue instead of behind the packet processor's barrier bit. Chained forms are recognised byr4dchainin the name and use aRAD_CHAIN_REC_BYTESrecord atRAD_CHAIN_KARG_OFF.rad_plugin_reference/rad_kernel_reference— an oracle library (or an oracle row beside a device row). A reference library ranks below every unnamed library, is refused for serving unless--debug-accept-reference-kernels, and is whatrad-kbenchchecks against.
Defining a new op
Only when docs/OPS.md has no name for it. The vocabulary document owns the rules: what a schema
must declare, and that the first plugin in hierarchy order to declare an op fixes its schema — a
later plugin declaring it differently is refused at load, both named. ctest -R ops_doc checks the
document against the registries, so the rules live there only. An op in docs/OPS.md is describable
by libref, and therefore checkable, even when your kernel cannot describe itself; an op you invent
runs only on a library that implements it.
A schema is positional data, not code:
static const RadParamSpec params[] = {
{"M", RAD_P_INT, RAD_REQUIRED, RAD_PROLE_NONE}, /* the ranged extent */
{"n", RAD_P_INT, RAD_REQUIRED, RAD_PROLE_NONE},
{"eps", RAD_P_F64, RAD_REQUIRED, RAD_PROLE_NONE},
{"dtype", RAD_P_STR, RAD_REQUIRED, RAD_PROLE_NONE},
};
static const RadOperandSpec opds[] = {
{"x", RAD_OPD_IN, 0}, {"w", RAD_OPD_WEIGHT, 1}, {"y", RAD_OPD_OUT, 0},
};
static const RadOpSchema schemas[] = {
{"my_op", params, 4, opds, 3, "what it computes, for a human"},
};
rad_kernel_schema_count/at return this table. RAD_REQUIRED/RAD_OPTIONAL/RAD_DERIVED and the
operand roles (IN, OUT, INOUT, WEIGHT, WTAB) are the whole vocabulary; the kernel's
constraints and launch do the rest.
Quantiser plugins
abi/rad_quant.h is the whole contract:
int rad_quant_count(void);
const RadQuantizerInfo* rad_quant_at(int i);
typedef struct RadQuantizerInfo {
const char* name; /* what a recipe names */
const char* doc;
const RadQuantOption* options; /* key, RAD_P_INT|F64|STR, default text, doc */
int n_options;
RadQuantEncodingFn encoding; /* the encoding this quantiser would write */
RadQuantRowBlockFn row_block; /* rows per quantize() call; 0 = whole weight */
RadQuantizeFn quantize; /* logical rows -> planes */
RadQuantDecodeFn decode; /* optional, only for a scheme of its own */
} RadQuantizerInfo;
The flow for each weight a recipe rule picks: ask encoding (given name, shape, and — at declare —
no movement group), let the architecture declare against that encoding, then hand the weight over in
row blocks, so a 51 GB table streams. Other calibration data is the quantiser's own option
(calib=DIR); the output is canonical planes only, so every kernel that accepts the encoding and
the core's decoder can read it — that is how the converter measures error without knowing the method.
A quantiser's own scheme must be decoded by its decode hook (the oracle, rad-kbench and the
error report call it), and must return RAD_E_UNSUPPORTED for an encoding it did not write. Serving a
new encoding also needs a kernel library whose layout accepts it and an architecture that declares
the weight at it — a new format is a new kernel row plus a plugin, never a converter-only change.
Two plugins naming the same quantiser are refused, both named.
The encoding model (abi/rad_encoding.h)
An encoding is the code type, an optional codebook, the grouping of scales/zero points/minimums, and
a transform. It is data, not an enum, so a plugin can invent one. The core can size, slice and
place any encoding from the struct alone, and decodes the plain and affine schemes itself:
q = code codes TILED, block [1x1], or [1xV] with a grid
q = grid[code][k] grid TABLE [entries x V]
q = table[code] table TABLE [1 x entries]
q = -q if the sign bit is set signs TILED u1
q = q - zero zero TILED
v = q * scale * scale.1 * ... every scale level
v = v - min * min.1 * ... a minimum in the value's domain
w = v . T^-1 the transform, per row over columns
So NVFP4 is fp4 codes, e4m3 scale per 16, f32 scale.1 per tensor; Q4_K is u4 codes, u6 scale and
u6 min per 32 each times an f16 per 256; IQ2_XXS is u8 codes per 8 into a 256×8 grid, a sign bit per
value and two scale levels.
- Planes are TILED (cover the weight, element
(i,j)covering a[br,bc]block;0extent means the whole extent) or TABLE (a codebook, grid, rotation parameter; every slice carries it). - Sub-byte elements are a bit stream along the row, lowest bits first; every row starts on a byte.
- Transforms:
"",fwhtN(Hadamard before quantisation; the1/Nlives in the stored scale),perm(GPTQ act-order, a TABLEt.perm),givens(ParoQuant rotations:t.pairs,t.angle, optionalt.scale, per-partt.part, composedt.dense).
A kernel only rearranges planes (relayout); choosing numbers is the quantiser's job. That split
is why a container never has to be reconverted when a kernel library changes its layout.
Build, install, search path
cmake_minimum_required(VERSION 3.21)
project(mykernels LANGUAGES CXX HIP) # CXX alone for a host-only plugin
find_package(radiance REQUIRED)
rad_plugin_defaults() # C++20, optimising build unless asked otherwise
rad_add_plugin(mykernels kernels SOURCES rows.cpp HIP_SOURCES kernels.hip)
The kind is kernels, architectures or quantizers — the directory the engine loads it from.
Only HIP_SOURCES compile as HIP: for RAD_GPU_TARGETS or a GPU_TARGETS list, with -DNDEBUG
outside Debug, uncompressed, against the HIP runtime the installed radiance found. find_package
also sets RADIANCE_ABI_VERSION (to fail at configure time rather than at dlopen),
RADIANCE_HOME, RADIANCE_KERNELS_DIR, RADIANCE_ARCHITECTURES_DIR,
RADIANCE_QUANTIZERS_DIR. cmake --install puts the plugin under the home of the radiance you built
against; RAD_PLUGIN_INSTALL_HOME overrides that.
$RADIANCE_HOME is a search path: homes separated by :, the first winning; each may hold
kernels/, architectures/, quantizers/, and an earlier home's same-name file shadows a later
one. Your plugins go ahead of an installation without copying it:
RADIANCE_HOME=$HOME/radiance-plugins:/opt/radiance/share/radiance radiance --model m.rad
# in the container:
-v /my/plugins:/plugins -e RADIANCE_HOME=/plugins:/opt/radiance/share/radiance
The kernel hierarchy is --kernels (or $RADIANCE_KERNELS): --kernels mykernels,libr4d puts
yours before libr4d. An entry matches a plugin's name or its file (libr4d, r4d, libr4d.so, a
path). An entry matching nothing is warned about by name. With no list, every library found loads,
reference last.
A library for a new card
The engine's own device code is compiled for every AMD family's generic target
(gfx9-generic, gfx9-4-generic, gfx10-1-generic, gfx10-3-generic, gfx11-generic,
gfx12-generic) beside the build's targets, so one engine binary starts anywhere. What it computes
there depends on the installed kernel library.
- Configure with
-DRAD_GPU_TARGETS=gfx942(orgfx11-generic). At load the engine asks the runtime which targets the card accepts, in order (chip, then family:gfx1036, thengfx10-3-generic), with the card's modes (gfx942:sramecc+:xnack-), and takes the best fit in the fat binary. - A library with nothing for this card is left out, not failed: its device rows leave selection,
its host rows stay, and the log says what it carries versus what the card runs. One installation
can carry a library per card.
rad_plugin_openmay decline explicitly withRAD_E_UNSUPPORTED. - The engine uses discrete cards if any exist (all of the first one's architecture), integrated only
when there is nothing else;
ROCR_VISIBLE_DEVICESchooses explicitly. On integrated graphics, device memory is host memory and the budget leaves an eighth of host RAM to the host. - fp8 bytes are OCP E4M3 (
e4m3fn, bias 7, max 448) wherever the engine reads them: container weights, an fp8 KV cache, activations. A card whose fp8 ise4m3fnuz(gfx942) converts in itsrelayouthook (the codes keep their bits except0x80and decode to half the value, so the scale doubles) and keeps activations/KV in OCP bytes, or converts at its own boundaries. - Tensor parallelism is ops you supply.
all_reduce,all_gatherand the fused collectives are rows like any other; the engine inserts none and assumes no link.
Checking your work
rad-kbench --kernels mylib,libr4d,libref --report mylib.md --bench
- It replays a recorded fixture of
libref's answers over real model geometries ($RADIANCE_HOME/kernels.rkb, pluskernels_moe.rkb), running every library with a row for a case. Correctness uses sampled values and the full-tensor norm and seeded projections; completeness is the op-by-library matrix; speed is timed on the same bytes. - Red zones: every device allocation carries a poisoned 256-byte margin, verified after the launch.
- Fusion: a row with
replacesis run against its named chain, out of your library, compared withmemcmp. - A skip is named, not failed — usually a missing
opd_shape.--strictalso fails a stale fixture. rad-avx-check --ref libref.so --avx mylib.sochecks a host library without a card, op by op, at every ISA level and with parallel regions forced onto several threads (for a library exportingavx_isa_force/avx_par_force).rad-schemas mylib.so libref.soprints every schema disagreement at once (it exits 1 on any), and works on a plugin the engine cannot load.--debug-selectionprints which row took each band and which constraint refused each other candidate;/graphon a running server shows the declared graph and which declarations the run phase issued.rad-info --pluginslists every plugin on the search path as the engine loads it: kind, version, declarations, reference status, whether its device code runs on this card, and its file.
A worked example to copy
plugin/spechub/ is the smallest real kernel plugin in the tree: two ops (mtp_couple_propose,
mtp_couple_verify), each with a device row and a host reference row. Read it in this order:
exports.cpp— the whole export block: info, ABI, schema count/at, kernel count/at, plusrad_kernel_referenceclassifying device rows 0 and host rows 1 so the reference gate does not refuse the device rows.spechub.cpp— the rows, their constraints and their argument readers.spechub.hip/ref_spechub.cpp— the device launch and the host oracle.CMakeLists.txt— therad_add_plugincall, the optionalenable_language(HIP), and-ffp-contract=offbecause an FMA in the ordered FP64 mass sums would move a token off the host reference.
Related
- Radiance Architecture Plugins — the third plugin kind, one level up.
- Radiance Conversion — how a quantiser is driven by a recipe.
- Radiance Tools —
rad-kbench,rad-tune,rad-schemas,rad-avx-checkin full. - Radiance Serve Options —
--kernels,--radiance-home, the tuning aids.