Add experimental EZ convenience API for simplified Metalium programming

[olofj]

Ticket

N/A — new experimental API

Problem description

The TT-Metalium programming API is powerful but demands significant boilerplate for even simple programs. A basic "read tiles, compute, write back" example requires manually managing MeshDevice lifecycle, constructing DeviceLocalBufferConfig + ReplicatedBufferConfig + MeshBuffer, assembling CircularBufferConfig objects, building TensorAccessorArgs, wiring DataMovementConfig/ComputeConfig, orchestrating MeshWorkload + EnqueueMeshWorkload, and cleaning up — typically 6-8 disconnected setup steps before any domain logic appears.

This is a steep learning curve for newcomers and makes the programming examples — the primary teaching material — harder to follow than they need to be. The intent of each program gets buried under API ceremony.

What's changed

A new tt::tt_metal::experimental::ez namespace provides two thin RAII wrappers (~690 lines of library code) that let programs express what they want instead of how to wire it up:

DeviceContext — owns device lifecycle, buffer creation, and data transfer:

DeviceContext ctx(0);
auto buf = ctx.dram_tile_buffer(n_tiles);
ctx.write(buf, data);
ctx.run(program);
auto result = ctx.read<bfloat16>(buf);

ProgramBuilder — fluent builder for programs with deferred kernel creation and method chaining:

auto program = ProgramBuilder(core)
    .cb(CBIndex::c_0).cb(CBIndex::c_1).cb(CBIndex::c_16)
    .reader("reader.cpp", {src0, src1})
        .runtime_args({src0->address(), n_tiles})
    .compute("compute.cpp")
        .named_args({{"cb_in0", CBIndex::c_0}, {"cb_in1", CBIndex::c_1}})
        .defines({{"ACTIVATION", "relu"}})
        .runtime_args({n_tiles})
    .writer("writer.cpp", {dst})
        .runtime_args({dst->address(), n_tiles})
    .build();

Kernel creation is deferred until build(). This means per-kernel configuration — defines(), named_args(), runtime_args() — can be called on the KernelRef returned by reader()/writer()/compute(), and each applies only to that kernel. Forwarding methods on KernelRef allow seamless chaining back to the builder without any intermediate commit step.

The API is intentionally minimal — it wraps existing Metalium primitives without hiding them. ctx.device() and builder.build() provide escape hatches to the full API at any point.

All 21 programming examples have been converted. The reduction is most dramatic in the simpler examples — exactly the ones newcomers encounter first — where up to 45% of the code was pure boilerplate. Even the complex multi-core matmul examples see meaningful cleanup. Educational comments about hardware concepts (tile layout, circular buffers, NoC mechanics, SFPU pipelines) are preserved throughout.

Example	Before	After	Reduction
eltwise_sfpu	174	95	45%
add_2_integers_in_compute	176	102	42%
eltwise_binary	192	127	33%
vecadd_multi_core	304	220	27%
matmul_single_core	261	198	24%
matmul_multicore_reuse	413	334	19%
Total (all 21 examples)	4,300	3,168	26%

Importantly, the lines removed are not comments or logic — they are API plumbing. The domain logic and educational content are preserved.

Key ProgramBuilder features exercised across the examples:

• reader()/writer() auto-generate TensorAccessorArgs from buffer lists
• .on(core) places kernels on specific cores for multi-core programs
• .runtime_args(lambda) computes per-core args without manual core iteration
• .defines() and .named_args() on KernelRef for per-kernel compile-time configuration
• .semaphore() and .physical_core() for synchronization and NoC multicast
• .cb(index, l1_buffer) for L1-backed circular buffers
• .cb(CircularBufferConfig) raw overload for shared CB index patterns

Commit structure:

1. Library + unit tests (12 test cases covering all API surface)
2. Simple single-core example conversions (10 examples)
3. Data movement and compute chain conversions (5 examples)
4. Multi-core example conversions (3 examples)
5. Matmul example conversions (3 examples — most complex)
6. Lazy kernel creation, per-kernel named_args/defines, eliminate done() (13 test cases)

Checklist

• Post-commit and Blackhole CI — cannot trigger workflows as an external contributor; maintainers please run.
• New/Existing tests provide coverage for changes
• All 13 unit tests and 19 example binaries verified on Blackhole p100a hardware

[olofj]

I just realized that the named_args() pattern of needing to be before the kernel (i.e. program()) statement feels backwards, especially since the runtime_args() are after. Looking into refactoring that, but the rest can be reviewed.