From source to artifacts¶

Example files: loom/docs/examples/elementwise-transform/

Loom keeps a program understandable while it moves from reusable source to a target-specific artifact. Linking, specialization, launch configuration, correctness cases, benchmark workloads, and compiler evidence operate on the same program instead of meeting only after code has become opaque.

The useful unit is therefore larger than an instruction stream and smaller than a framework. A Loom module can contain reusable computation, alternative implementations, dispatchable kernels, executable checks, and command programs. An application chooses which roots and facts matter, then asks Loom to link and specialize only that reachable program.

The canonical elementwise-transform example maps a selectable doubling motif over f32 buffers, then follows that program through linking, target specialization, native kernel emission, and command-program lowering.

One program, several forms¶

.loom is the canonical, human-readable source form. .loombc preserves the same linkable program in bytecode. Neither form implies one target or one deployment strategy.

motif.loom + kernel.loom + model.loom + configuration + target profile
                                  │
                          link and specialize
                                  │
                  lower the selected executable roots
                                  │
           native code + launch contracts + compiler evidence

The same source can be compiled ahead of time, linked and specialized when a model is loaded, or JIT-compiled for a concrete invocation. Those choices change when facts become available; they do not require different kernel languages.

Modules are ordinary composition boundaries. Public symbols can become roots for another module or an embedding, private symbols remain implementation details, and unreachable helpers and providers can disappear after selection. The linker resolves declared symbol relationships; merely placing a module in a library does not make its entire contents live.

Exact calls and selectable implementations¶

Loom separates naming one implementation from requesting an implementation that satisfies a contract.

Source construct	Meaning
`func.def` and `func.call`	Define and call an exact helper by symbol.
`func.template`	Provide a visible implementation of a named contract.
`func.apply`	Request an implementation of that contract at compile time.

An exact bit-manipulation helper is naturally a func.call. An operation whose best implementation depends on element format, shape facts, subgroup size, or target capabilities is naturally a func.apply. During specialization, Loom matches the reachable providers, their predicates and requirements, and the known facts. The selected provider becomes an ordinary callable boundary that normal inlining and optimization can remove.

This is the foundation of a scalable library: callers name semantic demands; libraries provide reusable implementations; target selection remains a fact of the final compilation rather than a global property baked into every helper. This example defines an exact helper, a wave32-specific provider, and a portable provider for the same guide.elementwise_transform contract:

Source: loom/docs/examples/elementwise-transform/motif.loom

motif.loom

// These targetless helpers and providers deliberately omit `public`. An
// explicitly linked library can satisfy reachable dependencies from them
// without adding their implementation symbols to the public module interface.
func.def inline @double(%value: f32) -> (f32) {
  %doubled = scalar.addf %value, %value : f32
  func.return %doubled : f32
}

func.template<guide.elementwise_transform> requires [#target.subgroup.size<32>] priority(20) @wave32_elementwise_transform(%value: f32) -> (f32) {
  %result = func.call @double(%value) : (f32) -> (f32)
  func.return %result : f32
}

func.template<guide.elementwise_transform> priority(1) @portable_elementwise_transform(%value: f32) -> (f32) {
  %result = func.call @double(%value) : (f32) -> (f32)
  func.return %result : f32
}

All three definitions omit public and are therefore private. Supplying the motif as an explicit library makes those implementations available while Loom forms the requested program, but it does not add them to the module's public interface. Provider selection and normal reachability can then remove the implementations the selected root does not need.

The wave32 provider is eligible only when the selected target establishes a subgroup size of 32. The fallback stays targetless, so the motif remains useful to targets that know nothing about AMDGPU.

The complete operation inventory lives in the generated func dialect reference.

A kernel owns two contracts¶

A kernel.def contains a launch configuration region and a device body. They answer different questions.

The launch configuration region receives workload values and computes the physical workgroup grid, required workgroup size, and optional cluster size. The body receives the arguments carried through the device launch ABI and uses workgroup, workitem, subgroup, and memory operations to perform the work.

This distinction keeps a workload such as “process 1009 elements” separate from the physical decision “launch four workgroups of 256 invocations.” An embedding evaluates the compiled launch configuration for the workload; it does not reverse-engineer grid dimensions from a kernel name or duplicate the kernel's scheduling arithmetic.

Workload arguments and device arguments are explicit even when the same value appears in both contracts. That explicit boundary lets Loom specialize launch geometry when facts are known while preserving a correct dynamic path when they are not. The generated kernel dialect reference defines the complete launch, execution, collective, and synchronization vocabulary.

The kernel below accepts its element count as a workload and device value, derives a one-subgroup workgroup size from the target profile, and applies the motif contract without naming either provider:

Source: loom/docs/examples/elementwise-transform/kernel.loom

kernel.loom

// The targetless kernel owns its physical launch and applies the transform
// contract without naming one provider.
kernel.def @elementwise_transform_f32(%element_count: index) {
  %unit = index.constant 1 : index
  %workgroup_size = target.subgroup.size : index
  %rounding = index.sub %workgroup_size, %unit : index
  %rounded_count = index.add %element_count, %rounding : index
  %workgroups = index.div %rounded_count, %workgroup_size : index
  kernel.launch.config workgroups(%workgroups, %unit, %unit) workgroup_size(%workgroup_size, %unit, %unit) : index
} launch(%element_count: index, %input: buffer, %output: buffer) {
  %workgroup = kernel.workgroup.id<x> : index
  %lane = kernel.workitem.id<x> : index
  %workgroup_size = kernel.workgroup.size<x> : index
  %element_index = index.madd %workgroup, %workgroup_size, %lane : index
  %input_noalias, %output_noalias = buffer.assume.noalias %input, %output : buffer, buffer
  %base = index.constant 0 : offset
  %input_view = buffer.view %input_noalias[%base] : buffer -> view<[%element_count]xf32>
  %output_view = buffer.view %output_noalias[%base] : buffer -> view<[%element_count]xf32>
  %in_bounds = index.cmp ult, %element_index, %element_count : index
  scf.if %in_bounds {
    %value = view.load %input_view[%element_index] : view<[%element_count]xf32> -> f32
    %result = func.apply<guide.elementwise_transform>(%value) : (f32) -> (f32)
    view.store %result, %output_view[%element_index] : f32, view<[%element_count]xf32>
  }
  kernel.return
}

Ask when a value becomes known¶

Most confusion about specialization disappears once every value has a clear binding stage.

Value category	Binding stage	Role
Configuration values	Link or compile time	Resolve `config.decl` through a matching definition and seed facts used by specialization.
Target facts	Compile invocation	Describe capabilities such as subgroup size, supported types, and resource limits for the selected executable root.
Kernel workload values	Launch-configuration evaluation	Compute the physical launch for one workload without pretending the values were compile-time constants.
Kernel launch arguments	Device issue	Carry scalars and buffers through the selected target ABI into the kernel body.
Command specialization arguments	Command-program specialization	Shape a reusable command artifact and its aggregate launch requirements.
Command buffer bindings	Command-program issue	Attach concrete parameter, transient, input, and output storage without recompiling the schedule.

Configuration is not a hidden global flag. A module declares artifact-level choices such as a model's layer count or weight encoding, another input defines them, and ordinary symbol dependency analysis keeps their effect visible:

config.decl @model.layer_count : %value: index where [range(%value, 1, 256)]
config.decl @model.weight_encoding : encoding<schema>

Per-launch element counts stay workload values instead of becoming global configuration. Likewise, target requirements are compile-time selection constraints rather than runtime branches. Reusable code generally stays targetless; a provider names a target requirement only when its algorithm actually depends on one.

Loom can preserve only information the program supplies. The ranges and divisibility predicates on a config.decl constrain a value even before a composition root binds it to one constant. Control flow also supplies path facts: inside the guarded region in kernel.loom, Loom knows that %element_index < %element_count without an extra assumption. An index.assume is for information the compiler cannot derive, such as a range guaranteed by the producer of a loaded routing index. The buffer.assume.* family carries root alignment, memory-space, aliasing, and identity facts; the example's buffer.assume.noalias makes the input and output independence explicit.

These are optimization inputs, not documentation comments. Shapes, ranges, alignment, aliasing, and target requirements survive linking and specialization so that each later stage can make a stronger decision without rediscovering information that the author or embedding already knew.

The generated config and target references define those two fact sources.

Correctness and measurement stay beside the program¶

A check.case is an executable SSA program that creates inputs, launches code, obtains expected values, and states the comparison policy. A check.benchmark selects named assignments from a case for timing. The benchmark row does not copy the setup or invent a second workload description.

Checks are test-only symbols, not deployment entry points. Keeping them in the same source gives formatters, verifiers, test runners, benchmark runners, and agents one authoritative statement of the workload while allowing published libraries to exclude the harness. The generated check dialect reference covers input generation, fixtures, requirements, expectations, and benchmark records.

Command programs move the same model upward¶

A command.program.def composes kernel launches, parameter views, resources, and serial or concurrent schedule regions into a reusable subgraph. Its leading specialization arguments participate in staged specialization and launch-count evaluation; its buffer bindings remain replaceable when the materialized program is issued.

That separation allows one program to specialize around model structure and target facts while still accepting different weights, cache storage, inputs, and outputs. It is the same idea as a specialized kernel launch, applied to a larger ownership boundary rather than hidden behind a separate graph compiler.

The outer module declares the exact kernel dependency and turns one concrete 1009-element invocation into a reusable command program. The kernel itself retains its dynamic workload contract and can be launched with other counts by other roots:

Source: loom/docs/examples/elementwise-transform/model.loom

model.loom

// The declaration makes the exact dependency explicit in this source module.
kernel.decl @elementwise_transform_f32(%workload_element_count: index) launch(%device_element_count: index, %input: buffer, %output: buffer)

// The command program owns the reusable schedule while buffers remain
// replaceable at issue time.
command.program.def public @elementwise_transform() launch(%input: buffer, %output: buffer) {
  %element_count = index.constant 1009 : index
  kernel.launch @elementwise_transform_f32[%element_count](%element_count, %input, %output) : [index](index, buffer, buffer)
  command.return
}

Current command-program surface

Command-program source, preparation, target-neutral lowering, and immutable artifact infrastructure are present. A general installed-tool materialization path is not yet a published user surface, so the guide does not present an executable command-program tutorial yet.

The generated command dialect reference documents the source constructs that exist today.

Follow one composition to Low¶

The three source listings above are repository .loom files, not prose copies. With the Loom tools on PATH, this command formats them, links and specializes the @elementwise_transform root, compiles its kernel for the generic GFX11 profile, and prints every Loom command it runs:

loom/docs/examples/elementwise-transform/run.sh \
  gfx11-generic build/elementwise-transform/gfx11-generic

The resulting directory contains the specialized elementwise-transform.loom, a VMFB, an HSACO, and the captured target Low IR. The documentation build invokes that same script and regenerates the views below; neither output is checked-in source.

GFX11 kernel LowCommand-program Low

low.kernel.def retain target<amdgpu.gfx11.generic.core>(@__loom_sealed_target_0) abi_layout({constant_count = 2, direct_arg_count = 1, direct_arg_names = {arg0 = "element_count"}, direct_arg_offsets = [0], direct_arg_parameter_indices = [0], direct_arg_sizes = [8], parameter_count = 3, resource_count = 2, resource_offsets = [8, 16], resource_parameter_indices = [1, 2], uses_kernarg_segment_ptr = true}) workgroup_size(32, 1, 1) @elementwise_transform_f32() {
  %kernarg = low.live_in<amdgpu.kernarg_segment_ptr> : reg<amdgpu.sgpr x2>
  %workgroup = low.live_in<amdgpu.workgroup_id.x> : reg<amdgpu.sgpr>
  %lane = low.live_in<amdgpu.workitem_id.x> : reg<amdgpu.vgpr>
  %element_count = low.op<amdgpu.s_load_dwordx2_offset_only>(%kernarg) {offset = 0} : (reg<amdgpu.sgpr x2>) -> reg<amdgpu.sgpr x2>
  %4 = low.op<amdgpu.s_load_dwordx4_offset_only>(%kernarg) {offset = 8} : (reg<amdgpu.sgpr x2>) -> reg<amdgpu.sgpr x4>
  %element_index = low.op<amdgpu.v_lshl_add_u32.shift_imm>(%workgroup, %lane) {shift = 5} : (reg<amdgpu.sgpr>, reg<amdgpu.vgpr>) -> reg<amdgpu.vgpr>
  %6 = low.const<amdgpu.v_mov_b32> {imm32 = 0} : reg<amdgpu.vgpr>
  %7 = low.slice %element_count[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %8 = low.op<amdgpu.v_mov_b32_copy>(%7) : (reg<amdgpu.sgpr>) -> reg<amdgpu.vgpr>
  %9 = low.slice %element_count[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %10 = low.op<amdgpu.v_mov_b32_copy>(%9) : (reg<amdgpu.sgpr>) -> reg<amdgpu.vgpr>
  %11 = low.op<amdgpu.v_cmp_ult_u32.src0_inline>(%10) {lhs = 0} : (reg<amdgpu.vgpr>) -> reg<amdgpu.sgpr x2>
  %12 = low.op<amdgpu.v_cmp_ult_u32>(%element_index, %8) : (reg<amdgpu.vgpr>, reg<amdgpu.vgpr>) -> reg<amdgpu.sgpr x2>
  %13 = low.op<amdgpu.v_cmp_eq_i32.src0_inline>(%10) {lhs = 0} : (reg<amdgpu.vgpr>) -> reg<amdgpu.sgpr x2>
  %14 = low.op<amdgpu.s_and_b64>(%13, %12) : (reg<amdgpu.sgpr x2>, reg<amdgpu.sgpr x2>) -> reg<amdgpu.sgpr x2>
  %in_bounds = low.op<amdgpu.s_or_b64>(%11, %14) : (reg<amdgpu.sgpr x2>, reg<amdgpu.sgpr x2>) -> reg<amdgpu.sgpr x2>
  %16, %17 = low.op<amdgpu.s_and_saveexec_b64>(%in_bounds) : (reg<amdgpu.sgpr x2>) -> reg<amdgpu.sgpr x2>, reg<amdgpu.scc>
  low.cond_br %17, ^_bb1, ^_bb2 : reg<amdgpu.scc>
^_bb1:
  %18 = low.op<amdgpu.v_lshlrev_b32.src0_inline>(%lane) {imm32 = 2} : (reg<amdgpu.vgpr>) -> reg<amdgpu.vgpr>
  %19 = low.const<amdgpu.s_mov_b32> {imm32 = 0} : reg<amdgpu.sgpr>
  %20 = low.concat(%workgroup, %19) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %21 = low.const<amdgpu.s_mov_b32> {imm32 = 7} : reg<amdgpu.sgpr>
  %22 = low.op<amdgpu.s_lshl_b64>(%20, %21) : (reg<amdgpu.sgpr x2>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %23 = low.slice %4[0] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %24 = low.slice %4[1] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %25 = low.slice %22[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %26 = low.slice %22[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %27 = low.op<amdgpu.s_add_u32>(%23, %25) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr>
  %28 = low.op<amdgpu.s_addc_u32>(%24, %26) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr>
  %29 = low.concat(%27, %28) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %value = low.op<amdgpu.global_load_b32_saddr>(%18, %29) {offset = 0} : (reg<amdgpu.vgpr>, reg<amdgpu.sgpr x2>) -> reg<amdgpu.vgpr>
  %doubled = low.op<amdgpu.v_add_f32>(%value, %value) : (reg<amdgpu.vgpr>, reg<amdgpu.vgpr>) -> reg<amdgpu.vgpr>
  %32 = low.concat(%workgroup, %19) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %33 = low.op<amdgpu.s_lshl_b64>(%32, %21) : (reg<amdgpu.sgpr x2>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %34 = low.slice %4[2] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %35 = low.slice %4[3] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %36 = low.slice %33[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %37 = low.slice %33[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %38 = low.op<amdgpu.s_add_u32>(%34, %36) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr>
  %39 = low.op<amdgpu.s_addc_u32>(%35, %37) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr>
  %40 = low.concat(%38, %39) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  low.op<amdgpu.global_store_b32_saddr>(%18, %doubled, %40) {offset = 0} : (reg<amdgpu.vgpr>, reg<amdgpu.vgpr>, reg<amdgpu.sgpr x2>)
  low.br ^_bb2
^_bb2:
  low.op<amdgpu.s_mov_b64_exec>(%16) : (reg<amdgpu.sgpr x2>)
  low.return
}

amdgpu.target<gfx11-generic> @__loom_sealed_target_0

low.func.def public retain target<cmd.core> abi(command_program) abi_layout({entry_count = 1, executable_count = 1, fixed_buffer_count = 0, rebindable_binding_count = 3}) @elementwise_transform() asm {
  %0 = resource<command_input> {index = 0, source_type = buffer} : reg<cmd.binding>
  %1 = resource<command_input> {index = 1, source_type = buffer} : reg<cmd.binding>
  %2 = resource<command_input> {index = 2, source_type = buffer} : reg<cmd.binding>
  %3 = resource<command_input> {index = 0, source_type = index} : reg<cmd.executable>
  %4 = resource<command_input> {index = 0, source_type = index} : reg<cmd.entry>
  %5 = cmd.constant.u64 0
  %6 = cmd.constant.u64 -1
  %7 = cmd.constant.u64 12
  %8 = cmd.buffer.ref.binding %2, %5, %7
  cmd.dispatch.indirect.static<%3, %4>[%8](%0, %5, %6, %1, %5, %6)
  return
}

The GFX11 tab comes directly from the installed-tool workflow above. The command-program tab shows compiler-owned materialization of the same linked root. Until command-program materialization has a published installed-tool surface, treat that tab as verified compiler output rather than a CLI recipe.

Embedding chooses the deployment policy¶

Loom produces modules, target artifacts, launch contracts, diagnostics, and structured compiler reports. The embedding decides when compilation happens, where artifacts are cached, how executable code is loaded, and how work is issued to a device runtime.

The public loomc C API exposes the in-memory source, link, specialization, compilation, emission, and launch-configuration boundaries needed to build those policies. Executable embedding examples in this guide are included only from checked programs that build against that API; the generated header reference remains authoritative for ownership, lifetime, threading, and failure contracts.

Embed kernel JIT compilation follows one of those checked programs from targetless source through launch evaluation and an in-memory native executable.

Working rules¶

Keep reusable algorithms targetless until a real implementation requirement needs target facts.
Use an exact call for an exact helper and a contract application when the implementation is intentionally substitutable.
Use configuration for compile-time product choices, workload arguments for per-launch shape, and launch arguments for the device ABI.
Put correctness policy in check.case; make benchmarks select those checked workloads instead of reconstructing them.
Keep motifs as functions and templates. Add a kernel ABI where the program actually becomes dispatchable, and a command-program ABI where a reusable subgraph becomes materializable.

Continue with Source modules and canonical text to learn the source and composition contract, or browse the generated language reference for exact syntax and operation contracts. If the Loom tools are not yet on PATH, Acquiring Loom describes the current installation status.