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.

Compose explicit modules

Every .loom file verifies independently. An exact call into another module therefore needs a local declaration that states the complete contract. The link invocation or embedding supplies the libraries that may satisfy it:

Source: loom/docs/examples/module-composition/root.loom

root.loom

// The declaration states the exact callable contract required from a library.
func.decl @project_layer(%value: i32) -> (i32)

func.def public @entry(%value: i32) -> (i32) {
  %result = func.call @project_layer(%value) : (i32) -> (i32)
  func.return %result : i32
}

The layer has its own declared dependency on @scale. The checked workflow links only the available layer, writes a standalone partial .loombc, reloads it, and then supplies the kernel library:

loom/docs/examples/module-composition/run.sh build/module-composition

Source contract	Selection behavior
Exact declaration	Resolve the compatible direct definition or unique exported library definition from the explicitly supplied source universe.
template.decl plus template.apply	Select an eligible family implementation from the supplied libraries using facts, requirements, and priority.
Runtime import attributes on func.decl	Preserve an external ABI contract for runtime resolution instead of satisfying it from Loom libraries.

Library paths and insertion order do not participate in symbol identity. The application or build graph chooses the source universe; Loom declarations state what that program requires. A partial link may preserve reachable declarations for a later boundary, while a closed link rejects any exact requirement that remains unresolved.

The linking workflow follows partial and transitive composition in detail. The source-module chapter defines the language contract.

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.
template.decl	Declare a stable compile-time family contract in the module that uses or implements it.
template.def	Provide one selectable implementation of that family.
template.apply	Request an implementation of that family at compile time.
template.call	Call one named template implementation without candidate ranking.

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 template.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
}

template.decl @guide.elementwise_transform(%value: f32) -> (f32)

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

template.def<@guide.elementwise_transform> priority(1) @portable_elementwise_transform(%value: f32) -> (f32) {
  %result = func.call @double(%value) : (f32) -> (f32)
  template.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 inventories live in the generated func and template dialect references.

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, uses a portable 64-thread workgroup, and applies the motif contract without naming either provider. The target profile still selects the eligible motif implementation independently of that launch geometry:

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

kernel.loom

// This reusable source kernel owns portable physical launch geometry and
// applies the transform contract without naming one provider.
template.decl @guide.elementwise_transform(%value: f32) -> (f32)

kernel.def export("elementwise_transform_f32") @elementwise_transform_f32(%element_count: index) {
  %unit = index.constant 1 : index
  %workgroup_size = index.constant 64 : 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 = template.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	Link or compile product boundary	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 source control flow and launch computation; direct counts and scalar command data must become exact before portable preparation.
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 pure launch-count computation; its buffer bindings remain replaceable when the materialized program is issued. Exact counts become direct commands, while intentionally dynamic counts use an explicit view<3xi32> produced before an indirect dispatch.

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. Its portable workgroup geometry makes the physical count exact without loading the kernel implementation. 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
}

Command deployment products

loom-compile --format=loom-command materializes each selected command root as a portable .loomcmd and writes the shared executable-entry manifest. The target-specific kernel executable remains a separate artifact so an embedding can cache, replace, or prebuild it independently.

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 the selected program against the explicit provider universe, specializes that closure for the generic GFX11 profile, and emits both kernel and command products. It 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 closed target-specialized elementwise-transform.loom module, the GFX11 HSACO, a command manifest, one portable .loomcmd, and the captured target Low IR. The --target=amdgpu:gfx11-generic link argument makes subgroup facts available before template-provider pruning, so the wave32 implementation is selected without pulling the portable alternative into the product. The documentation build invokes the 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_target_context_0_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(64, 1, 1) @elementwise_transform_f32() asm {
  %kernarg = live_in<amdgpu.kernarg_segment_ptr> : reg<amdgpu.sgpr x2>
  %workgroup = live_in<amdgpu.workgroup_id.x> : reg<amdgpu.sgpr>
  %lane = live_in<amdgpu.workitem_id.x> : reg<amdgpu.vgpr>
  %element_count = s_load_dwordx2_offset_only %kernarg
  %4 = s_load_dwordx4_offset_only %kernarg {offset = 8}
  %element_index = v_lshl_add_u32_shift_imm %workgroup, %lane, 6
  %6 = v_mov_b32 0
  %7 = slice %element_count[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %8 = v_mov_b32_copy %7
  %9 = slice %element_count[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %10 = v_mov_b32_copy %9
  %11 = v_cmp_lt_u32_src0_inline %10 {lhs = 0}
  %12 = v_cmp_lt_u32 %element_index, %8
  %13 = v_cmp_eq_i32_src0_inline %10 {lhs = 0}
  %14 = s_and_b64 %13, %12
  %in_bounds = s_or_b64 %11, %14
  %16, %17 = s_and_saveexec_b64 %in_bounds
  low.cond_br %17, ^_bb1, ^_bb2 : reg<amdgpu.scc>
^_bb1:
  %18 = v_lshlrev_b32_src0_inline %lane, 2
  %19 = s_mov_b32 0
  %20 = concat(%workgroup, %19) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %21 = s_mov_b32 8
  %22 = s_lshl_b64 %20, %21
  %23 = slice %4[0] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %24 = slice %4[1] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %25 = slice %22[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %26 = slice %22[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %27 = s_add_u32 %23, %25
  %28 = s_addc_u32 %24, %26
  %29 = concat(%27, %28) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %value = global_load_b32_saddr %18, %29
  %doubled = v_add_f32 %value, %value
  %32 = concat(%workgroup, %19) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  %33 = s_lshl_b64 %32, %21
  %34 = slice %4[2] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %35 = slice %4[3] : reg<amdgpu.sgpr x4> -> reg<amdgpu.sgpr>
  %36 = slice %33[0] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %37 = slice %33[1] : reg<amdgpu.sgpr x2> -> reg<amdgpu.sgpr>
  %38 = s_add_u32 %34, %36
  %39 = s_addc_u32 %35, %37
  %40 = concat(%38, %39) : (reg<amdgpu.sgpr>, reg<amdgpu.sgpr>) -> reg<amdgpu.sgpr x2>
  global_store_b32_saddr %18, %doubled, %40
  low.br ^_bb2
^_bb2:
  s_mov_b64_exec %16
  return
}

amdgpu.target<gfx11-generic> @__loom_target_context_0_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 = 2}) @elementwise_transform() asm {
  %23 = resource<command_input> {index = 0, source_type = buffer} : reg<cmd.binding>
  %24 = resource<command_input> {index = 1, source_type = buffer} : reg<cmd.binding>
  %25 = resource<command_input> {index = 0, source_type = index} : reg<cmd.executable>
  %26 = resource<command_input> {index = 0, source_type = index} : reg<cmd.entry>
  %27 = cmd.constant.u64 0
  %28 = cmd.constant.u64 -1
  %29 = cmd.constant.u32 16
  %30 = cmd.constant.u32 1
  %31 = cmd.constant.b64 1009
  cmd.dispatch.direct<%25, %26>[%29, %30, %30](%31, %23, %27, %28, %24, %27, %28)
  return
}

The GFX11 tab comes directly from the installed-tool workflow above. The command-program tab is the readable Low view of the .loomcmd produced from the same closed program: launch configuration is projected as pure caller IR, folded to an exact direct count, and lowered while the device implementation remains unopened. The deployment artifact is binary and target-neutral; the HSACO independently supplies its logical kernel entry.

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.