# Adding A Backend
Any vendor or project can add a backend without asking for approval. The
contract is public, the suite is automatic, and the certificate is the result.
## Process
1. Implement `VyreBackend` for your runtime.
2. Run `ConformanceSuite::new().run(&backend)`.
3. Inspect the returned failures.
4. Fix the backend or runtime behavior that caused each failure.
5. Run the suite again until it returns zero failures.
6. Publish the generated certificate for the level you passed.
No communication with the vyre maintainers is required. A backend that passes
the suite at a level conforms at that level. A backend that fails any required
test does not conform at that level.
## Minimal WGSL Backend
The backend receives complete WGSL and must return exact output bytes. In a
real backend, `VendorRuntime` is your runtime wrapper around device creation,
shader compilation, buffer upload, dispatch, synchronization, and readback.
```rust
use vyre_conform::{ConformanceSuite, Convention, DispatchConfig, VyreBackend};
pub struct VendorBackend {
runtime: VendorRuntime,
}
impl VendorBackend {
pub fn new(runtime: VendorRuntime) -> Self {
Self { runtime }
}
}
impl VyreBackend for VendorBackend {
fn name(&self) -> &str {
"vendor-runtime"
}
fn max_convention(&self) -> Convention {
Convention::V1
}
fn dispatch(
&self,
wgsl: &str,
input: &[u8],
output_size: usize,
config: DispatchConfig,
) -> Result<Vec<u8>, String> {
let output = self.runtime.run_compute_shader(
wgsl,
input,
output_size,
config.workgroup_size,
config.workgroup_count,
config.lookup_data.as_deref(),
).map_err(|err| {
format!(
"vendor dispatch failed: {err}. Fix: compile the supplied WGSL, bind input/output/params buffers according to the requested convention, dispatch the requested workgroups, and read back exactly {output_size} bytes."
)
})?;
if output.len() != output_size {
return Err(format!(
"vendor dispatch returned {} bytes, expected {output_size}. Fix: size the readback buffer from output_size and return no padding bytes.",
output.len()
));
}
Ok(output)
}
}
fn certify(runtime: VendorRuntime) -> Result<(), Vec<vyre_conform::ParityFailure>> {
let backend = VendorBackend::new(runtime);
let suite = ConformanceSuite::new();
let failures = suite.run(&backend);
if failures.is_empty() {
Ok(())
} else {
Err(failures)
}
}
```
## Native IR Backend
Backends that do not consume WGSL should override `dispatch_program()`. The
`program` argument is a serialized `ir::Program` produced by
`Program::to_bytes()`. Native backends deserialize the IR, validate it, lower it
to the target kernel format, execute it with the supplied input/config, and
return exactly `output_size` bytes.
```rust
impl VyreBackend for VendorBackend {
fn name(&self) -> &str {
"vendor-runtime"
}
fn dispatch(
&self,
wgsl: &str,
input: &[u8],
output_size: usize,
config: DispatchConfig,
) -> Result<Vec<u8>, String> {
self.runtime.run_wgsl(wgsl, input, output_size, config).map_err(|err| {
format!(
"vendor WGSL dispatch failed: {err}. Fix: compile the supplied WGSL or override dispatch_program for native IR tests."
)
})
}
fn dispatch_program(
&self,
program: &[u8],
input: &[u8],
output_size: usize,
config: DispatchConfig,
) -> Result<Vec<u8>, String> {
let ir = vyre::ir::Program::from_bytes(program)?;
let output = self.runtime.run_ir(ir, input, output_size, config).map_err(|err| {
format!(
"vendor IR dispatch failed: {err}. Fix: lower the serialized vyre IR to the native target, bind resources according to DispatchConfig, and read back exactly {output_size} bytes."
)
})?;
if output.len() != output_size {
return Err(format!(
"vendor IR dispatch returned {} bytes, expected {output_size}. Fix: size native output/readback storage from output_size.",
output.len()
));
}
Ok(output)
}
}
```
`max_convention()` returns the highest calling convention the backend fully
supports. If you omit it, the default is `Convention::V1`. Return
`Convention::V2 { lookup_binding: 3 }` only after your runtime binds the
additional lookup buffer required by V2. The suite must never dispatch a test
whose convention is higher than the value returned here.
The important rule is that `dispatch()` executes the WGSL the suite supplied
and `dispatch_program()` executes the IR the suite supplied. Do not special-case
operation IDs, substitute CPU results, skip kernels, or cache outputs across
inputs. Those behaviors invalidate the certificate.