Loom

https://rocm.github.io/hrx-system/loom/
Admin
Loom is a source-first compiler for kernels and asynchronous device programs, keeping program...
Tokens:214,711
Snippets:2,951
Trust Score:9.7
Update:1 month ago
Show doc for...
Tokens:
Raw
### Run the module linking example

Source: https://rocm.github.io/hrx-system/loom/integration/module-composition

Executes the link_modules example using the Bazel build system.

```bash
python dev.py bazel run //loom/binding/c/example:link_modules
```

--------------------------------

### vector.iota usage examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/iota

Examples demonstrating the construction of vectors with static and dynamic shapes using base and step values.

```mlir
%lanes = vector.iota %c0 step %c1 : vector<16xindex>
```

```mlir
%offsets = vector.iota %base step %stride : vector<4xi32>
```

```mlir
%dyn = vector.iota %c0 step %c1 : vector<[%n]xindex>
```

--------------------------------

### vector.store.compress usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/store-compress

Demonstrates a basic store operation using a vector, a view, and a mask.

```mlir
vector.store.compress %v, %view[%row, %col], %mask : vector<4xf32>, view<[%m]x[%n]xf32, %layout>, vector<4xi1>
```

--------------------------------

### Prefetch operation example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/view/ops/prefetch

Demonstrates the syntax for the view.prefetch operation with read intent and L2 locality.

```mlir
view.prefetch %view[%row, %col] {intent = read, locality = l2} : view<[%M]x[%N]xf32, %layout>
```

--------------------------------

### Matrix Multiply-Accumulate Examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/mma

Examples demonstrating the vector.mma operation with different vector types for lhs, rhs, and accumulator operands.

```mlir
%r = vector.mma %lhs, %rhs, %init : vector<8xf16>, vector<8xf16>, vector<8xf32>
```

```mlir
%r = vector.mma %lhs, %rhs, %init : vector<6xi32>, vector<6xi32>, vector<8xf32>
```

--------------------------------

### Unconditional branch examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/cfg/ops/br

Examples showing basic unconditional branching and branching with forwarded arguments.

```mlir
cfg.br ^done
```

```mlir
cfg.br ^join(%value: i32)
```

--------------------------------

### vector.dot4f8 usage examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/dot4f8

Examples showing the syntax for vector.dot4f8 with different floating-point kind configurations.

```mlir
%r = vector.dot4f8<fp8bf8> %lhs, %rhs, %acc : vector<4xi32>, vector<4xf32>
```

```mlir
%r = vector.dot4f8<bf8fp8> %lhs, %rhs, %acc : vector<[%N]xi32>, vector<[%N]xf32>
```

--------------------------------

### Create a buffer view

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/buffer/ops/view

Example of forming a typed view from a buffer and an offset.

```mlir
%view = buffer.view %buffer[%offset] : buffer -> view<[%M]xf32, %layout>
```

--------------------------------

### func.return operation examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/func/ops/return

Examples showing the func.return operation with no operands, a single operand, and multiple operands.

```MLIR
func.return
```

```MLIR
func.return %r : f32
```

```MLIR
func.return %a, %b : tensor<[%M]xf32>, index
```

--------------------------------

### kernel.launch.yield usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/launch-yield

Basic syntax for the kernel.launch.yield terminator operation.

```mlir
kernel.launch.yield
```

--------------------------------

### Usage of kernel.subgroup.vote.all

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/subgroup-vote-all

Example demonstrating the syntax for invoking the subgroup vote all operation.

```mlir
%all = kernel.subgroup.vote.all %p : i1
```

--------------------------------

### buffer.fill operation example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/buffer/ops/fill

Demonstrates the syntax for filling a buffer with a specific scalar pattern.

```MLIR
buffer.fill %pattern, %target[%target_offset], %byte_length : bf16
```

--------------------------------

### Basic scf.for loop examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scf/ops/for

Simple loop structures using index or byte-offset domains.

```mlir
scf.for %iv = [%c0 to %n step %c1] {
  scf.yield
}
```

```mlir
scf.for %byte_offset = [%zero to %byte_length step %one] {
  scf.yield
}
```

--------------------------------

### Floating-point addition examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/addf

Basic usage of the scalar.addf operation with and without fast-math flags.

```mlir
%result = scalar.addf %lhs, %rhs : f32
```

```mlir
%result = scalar.addf<fast> %lhs, %rhs : f32
```

--------------------------------

### View instruction encoding output

Source: https://rocm.github.io/hrx-system/loom/workflows/oracles/llvm-mc

Example of the output format when using the --show-encoding flag.

```text
v_cvt_pk_f32_fp8_e32 v[6:7], v4.l
  encoding: [0x04,0xdd,0x0c,0x7e]
```

--------------------------------

### Shuffle operation example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/subgroup-shuffle

Demonstrates the syntax for performing a shuffle operation with the xor mode.

```MLIR
%r, %valid = kernel.subgroup.shuffle<xor> %v, %offset, %width : f32, i32, i32
```

--------------------------------

### Insert bits into a vector

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/bitfield-insert

Example of inserting a 4-bit field starting at offset 0 into a vector of 16-bit integers.

```mlir
%packed = vector.bitfield.insert %lo into %zero {offset = 0, width = 4} : vector<16xi32>, vector<16xi8>
```

--------------------------------

### Retrieve subgroup size in MLIR

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/subgroup-size

Example showing the syntax for invoking the kernel.subgroup.size operation to get the subgroup size as an index.

```mlir
%size = kernel.subgroup.size : index
```

--------------------------------

### Dispatching a kernel with project configuration

Source: https://rocm.github.io/hrx-system/loom/guide/command-programs

Demonstrates the syntax for calling a project configuration and dispatching a kernel using the resulting indices.

```MLIR
%x, %y, %z = func.call pure @project_block$config(%element_count) : (index) -> (index, index, index)
kernel.dispatch @project_block[%x, %y, %z](%weights, %input, %output) : [index, index, index](view<8xf32>, buffer, buffer)
```

--------------------------------

### Create target profile and compiler

Source: https://rocm.github.io/hrx-system/loom/integration/jit-kernel

Sets up the AMDGPU profile and compiler instance for the JIT process.

```c
  loomc_amdgpu_profile_options_t profile_options = {
      .type = LOOMC_STRUCTURE_TYPE_AMDGPU_PROFILE_OPTIONS,
      .structure_size = sizeof(profile_options),
      .identifier = loomc_make_cstring_view("guide-amdgpu"),
      .identity =
          {
              .target = loomc_make_cstring_view(target),
          },
  };
  if (loomc_status_is_ok(status)) {
    status = loomc_target_profile_create_amdgpu(
        state->target_environment, &profile_options, loomc_allocator_system(),
        &state->target_profile);
  }
  if (loomc_status_is_ok(status)) {
    status = loomc_compiler_create(state->context, NULL,
                                   loomc_allocator_system(), &state->compiler);
  }
```

--------------------------------

### buffer.compare usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/buffer/ops/compare

Performs a lexicographical comparison of two buffer ranges starting at specified offsets for a given byte length.

```mlir
%order = buffer.compare %lhs[%lhs_offset], %rhs[%rhs_offset], %byte_length
```

--------------------------------

### Load i64 from buffer example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/ireevm/ops/buffer-load-i64

Demonstrates the syntax for loading an i64 value from a buffer at a given offset.

```mlir
%value = ireevm.buffer.load.i64 %buffer[%offset] : ireevm.ref<ireevm.buffer>, i64 -> i64
```

--------------------------------

### Configure invocation options

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/config_8h.html

Demonstrates how to define configuration bindings and options, including a JSON configuration string and policy flags.

```c
const char* json_config = "{\"tile_m\":\"128\",\"tile_n\":\"64\"}";
loomc_config_binding_t overrides[] = {
{
.key = loomc_make_cstring_view("tile_m"),
.value = loomc_make_cstring_view("256"),
},
};
loomc_config_options_t config = {
.bindings = overrides,
.binding_count = 1,
.json_object = loomc_make_cstring_view(json_config),
.flags = LOOMC_CONFIG_POLICY_FLAG_REQUIRE_RESOLVED,
};
```

--------------------------------

### Load i32 from buffer example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/ireevm/ops/buffer-load-i32

Demonstrates the syntax for loading an i32 value from a buffer at a given offset.

```mlir
%value = ireevm.buffer.load.i32 %buffer[%offset] : ireevm.ref<ireevm.buffer>, i64 -> i32
```

--------------------------------

### Runtime Loading Facts Output

Source: https://rocm.github.io/hrx-system/loom/integration/jit-kernel

Example output showing workload configuration and artifact details.

```text
workload=1009 workgroups=16x1x1 workgroup_size=64x1x1
artifact=guide_jit_kernel.hsaco format=amdgpu-hsaco bytes=<target-dependent>
```

--------------------------------

### loomc_module_function_get_export_info

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/module_8h_source.html

Gets export metadata for a function.

```APIDOC
## loomc_module_function_get_export_info

### Description
Gets export metadata for a function.

### Signature
`loomc_status_t loomc_module_function_get_export_info(const loomc_module_t *module, const loomc_module_function_t *function, loomc_module_function_export_info_t *out_info)`
```

--------------------------------

### Create a Vulkan-style SPIR-V target profile

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/spirv_2profile_8h.html

Demonstrates initializing feature, limit, and environment facts to create a reusable SPIR-V target profile.

```c
loomc_spirv_feature_fact_t facts[] = {
{
.feature = LOOMC_SPIRV_FEATURE_FLOAT16,
.state = LOOMC_TARGET_FACT_STATE_TRUE,
.provenance = loomc_make_cstring_view("vulkaninfo:shaderFloat16"),
},
};
loomc_spirv_limit_fact_t limits[] = {
{
.limit = LOOMC_SPIRV_LIMIT_MAX_WORKGROUP_SIZE_X,
.state = LOOMC_TARGET_FACT_STATE_TRUE,
.value = 1024,
.provenance =
loomc_make_cstring_view("vulkaninfo:maxComputeWorkGroupSize[0]"),
},
};
loomc_spirv_environment_fact_t environment[] = {
{
.environment = LOOMC_SPIRV_ENVIRONMENT_MAX_SPIRV_VERSION,
.state = LOOMC_TARGET_FACT_STATE_TRUE,
.value = LOOMC_SPIRV_VERSION_1_3,
.provenance = loomc_make_cstring_view("vulkaninfo:apiVersion"),
},
};
loomc_spirv_profile_options_t options = {
.type = LOOMC_STRUCTURE_TYPE_SPIRV_PROFILE_OPTIONS,
.structure_size = sizeof(loomc_spirv_profile_options_t),
.identifier = loomc_make_cstring_view("offline-vulkan13"),
.preset = LOOMC_SPIRV_PROFILE_PRESET_VULKAN_1_3_BDA,
.feature_facts = facts,
.feature_fact_count = 1,
.limit_facts = limits,
.limit_fact_count = 1,
.environment_facts = environment,
.environment_fact_count = 1,
};
loomc_target_profile_t* profile = NULL;
loomc_result_t* result = NULL;
loomc_status_t status = loomc_target_profile_create_spirv(
target_environment, &options, loomc_allocator_system(), &profile,
&result);
if (!loomc_status_is_ok(status)) return status;
if (!loomc_result_succeeded(result)) {
// Inspect structured diagnostics; profile is NULL on profile failure.
}
loomc_result_release(result);
```

--------------------------------

### loomc_module_get_global_at

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/module_8h.html

Gets global metadata by global ordinal.

```APIDOC
## loomc_status_t loomc_module_get_global_at(const loomc_module_t *module, loomc_host_size_t global_ordinal, loomc_module_global_t *out_global)

### Description
Gets global metadata by global ordinal.

### Parameters
- **module** (const loomc_module_t *) - Module to inspect.
- **global_ordinal** (loomc_host_size_t) - Ordinal in the module's unfiltered global metadata sequence.
- **out_global** (loomc_module_global_t *) - Receives global metadata.

### Returns
- **loomc_status_t** - OK when the ordinal names a supported global, NOT_FOUND when it does not, or another non-OK status for API misuse.
```

--------------------------------

### loomc_module_get_function_at

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/module_8h.html

Gets function metadata by public function ordinal.

```APIDOC
## loomc_status_t loomc_module_get_function_at(const loomc_module_t *module, loomc_host_size_t function_ordinal, loomc_module_function_t *out_function)

### Description
Gets function metadata by public function ordinal.

### Parameters
- **module** (const loomc_module_t *) - Module to inspect.
- **function_ordinal** (loomc_host_size_t) - Ordinal in the module's unfiltered function metadata sequence.
- **out_function** (loomc_module_function_t *) - Receives function metadata.

### Returns
- **loomc_status_t** - OK when the ordinal names a supported function, NOT_FOUND when it does not, or another non-OK status for API misuse.
```

--------------------------------

### Example Linked .loom Module

Source: https://rocm.github.io/hrx-system/loom/workflows/link-and-package

A sample .loom module demonstrating function definitions, calls, and constant scaling operations.

```loom
func.def @project_layer(%value: i32) -> (i32) {
  %result = func.call @scale(%value) : (i32) -> (i32)
  func.return %result : i32
}

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

// This library satisfies the exact contract retained in the partial artifact.
func.def pure @scale(%value: i32) -> (i32) {
  %factor = scalar.constant 2 : i32
  %result = scalar.muli %value, %factor : i32
  func.return %result : i32
}
```

--------------------------------

### Integer Negation Example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/negi

Performs negation on a 32-bit integer input.

```MLIR
%result = scalar.negi %input : i32
```

--------------------------------

### Choose generic or exact specialization

Source: https://rocm.github.io/hrx-system/loom/workflows/compile-artifacts

Demonstrates the difference between targeting a generic profile for portability and an exact profile for specific hardware features.

```bash
loom-compile kernel.loom \
  --backend=amdgpu-hal \
  --target=gfx11-generic \
  --output=kernel-gfx11.hsaco
```

```bash
loom-compile kernel.loom \
  --backend=amdgpu-hal \
  --target=gfx1151 \
  --output=kernel-gfx1151.hsaco
```

--------------------------------

### Build and run a VM binary

Source: https://rocm.github.io/hrx-system/loom/workflows/build-with-bazel

Builds a VM binary using Bazel and executes a specific function within the resulting .vmfb file.

```bash
bazel build //loom/docs/examples/elementwise-transform:elementwise_vm

iree-run-module \
  --module=bazel-bin/loom/docs/examples/elementwise-transform/elementwise_vm.vmfb \
  --function=double_i32 \
  --input=21
```

--------------------------------

### Define a public check.case

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/check/ops/case

An example of a check.case operation with public visibility.

```mlir
check.case public @smoke {
  check.return
}
```

--------------------------------

### Create a context linked with a target environment

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/target_8h.html

Demonstrates initializing a SPIR-V target environment and attaching it to a new Loom context using options structures.

```c
#include "loomc/target/spirv/base.h"
loomc_target_environment_t* target_environment = NULL;
loomc_status_t status = loomc_target_environment_create_spirv(
loomc_allocator_system(), &target_environment);
if (!loomc_status_is_ok(status)) return status;
loomc_context_target_options_t target_options = {
.type = LOOMC_STRUCTURE_TYPE_CONTEXT_TARGET_OPTIONS,
.structure_size = sizeof(loomc_context_target_options_t),
.target_environment = target_environment,
};
loomc_context_options_t context_options = {
.type = LOOMC_STRUCTURE_TYPE_CONTEXT_OPTIONS,
.structure_size = sizeof(loomc_context_options_t),
.next = &target_options,
};
loomc_context_t* context = NULL;
status = loomc_context_create(&context_options, loomc_allocator_system(),
&context);
loomc_target_environment_release(target_environment);
if (!loomc_status_is_ok(status)) return status;
// `context` can now parse SPIR-V target records and create target
// pipelines backed by the prepared target environment.
```

--------------------------------

### scalar.maxsi usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/maxsi

Basic usage of the scalar.maxsi operation with 32-bit integers.

```MLIR
%result = scalar.maxsi %lhs, %rhs : i32
```

--------------------------------

### pool.store usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/pool/ops/store

Demonstrates the syntax for storing tile data into a pool with specified page and offset parameters.

```MLIR
pool.store %pool, %pid, %pb, %off, %data : pool<[%BS]>, i32, offset, offset, tile<[16, 128]xf16>
```

--------------------------------

### Floating-point multiplication example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/mulf

Basic usage of the scalar.mulf operation with f32 types.

```MLIR
%result = scalar.mulf %lhs, %rhs : f32
```

--------------------------------

### Launch program with specializations and bindings

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/command/ops/program-launch

Invokes a program with explicit specialization values and buffer bindings.

```MLIR
command.program.launch @decode[%token_count](%parameters, %transient) : [index](buffer, buffer)
```

--------------------------------

### scalar.fmai usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/fmai

Basic usage of the fmai operation with 64-bit integers.

```MLIR
%result = scalar.fmai %a, %b, %c : i64
```

--------------------------------

### buffer.pack usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/buffer/ops/pack

Demonstrates packing two byte ranges with specific alignment requirements into a single slab.

```mlir
%total, %header_offset, %payload_offset = buffer.pack [align(16) %header_bytes, align(256) %payload_bytes] : offset
```

--------------------------------

### scalar.minnumf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/minnumf

Basic usage of the minnumf operation with f32 operands.

```mlir
%result = scalar.minnumf %lhs, %rhs : f32
```

--------------------------------

### Launch static program

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/command/ops/program-launch

Invokes a program with parameters and a single buffer binding.

```MLIR
command.program.launch @static_program(%parameters) : (buffer)
```

--------------------------------

### Load and Invoke a JIT Kernel

Source: https://rocm.github.io/hrx-system/loom/integration/jit-kernel

Demonstrates loading a program, looking up a function by name, and invoking it with workload arguments.

```c
  loomc_status_t status = loomc_launch_config_program_load(
      artifact, NULL, NULL, loomc_allocator_system(), &state->launch_program);
  if (loomc_status_is_ok(status)) {
    jit_kernel_state_reset_result(state);
  }
  loomc_launch_config_function_t function =
      loomc_launch_config_function_invalid();
  if (loomc_status_is_ok(status)) {
    status = loomc_launch_config_program_lookup_function(
        state->launch_program, loomc_make_cstring_view(kKernelExportName),
        &function);
  }

  const uint64_t workload_argument_bits[] = {kElementCount};
  if (loomc_status_is_ok(status)) {
    status = loomc_launch_config_program_invoke(state->launch_program, function,
                                                workload_argument_bits, 1,
                                                out_launch_config);
  }
  return status;
}
```

--------------------------------

### scalar.maxui usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/maxui

Computes the maximum of two 32-bit unsigned integers.

```MLIR
%result = scalar.maxui %lhs, %rhs : i32
```

--------------------------------

### Define Loom build targets in BUILD.bazel

Source: https://rocm.github.io/hrx-system/loom/workflows/build-with-bazel

Example BUILD file demonstrating the use of loom_library, loom_kernel_binary, loom_command_binary, and loom_vm_binary rules.

```starlark
load(
    "@hrx//loom/build_tools/bazel:defs.bzl",
    "loom_command_binary",
    "loom_kernel_binary",
    "loom_library",
    "loom_vm_binary",
)
load("@rules_shell//shell:sh_test.bzl", "sh_test")

package(
    default_visibility = ["//visibility:private"],
    licenses = ["notice"],  # Apache 2.0
)

loom_library(
    name = "motif",
    srcs = ["motif.loom"],
)

loom_library(
    name = "kernel",
    srcs = ["kernel.loom"],
    deps = [":motif"],
)

loom_library(
    name = "model",
    srcs = ["model.loom"],
    deps = [":kernel"],
)

loom_kernel_binary(
    name = "elementwise_kernel",
    target = "@hrx//loom/target/amdgpu:gfx11-generic",
    deps = [":kernel"],
)

loom_command_binary(
    name = "elementwise_command",
    roots = ["@elementwise_transform"],
    target = "@hrx//loom/target/amdgpu:gfx11-generic",
    deps = [":model"],
)

loom_library(
    name = "vm",
    srcs = ["vm.loom"],
)

loom_vm_binary(
    name = "elementwise_vm",
    module_name = "elementwise_support",
    deps = [":vm"],
)
```

--------------------------------

### scalar.log1pf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/log1pf

Basic usage of the log1pf operation in the scalar dialect.

```MLIR
%result = scalar.log1pf %input : f32
```

--------------------------------

### Asynchronous Tensor Store from LDS Example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/async-tensor-store-from-lds

Demonstrates the syntax for initiating an asynchronous tensor store operation using a source LDS tile, a destination global tile, and a descriptor.

```mlir
%copy = kernel.async.tensor.store.from.lds %lds_tile to %global_tile using %desc {cache_scope = device, cache_temporal = non_temporal_writeback} : view<64x64xf32> to view<64x64xf32>, kernel.tensor.lds.descriptor -> kernel.async.token
```

--------------------------------

### Floating-point division example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/divf

Basic usage of the scalar.divf operation for f32 types.

```MLIR
%result = scalar.divf %lhs, %rhs : f32
```

--------------------------------

### scalar.acoshf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/acoshf

Basic usage of the acoshf operation in the scalar dialect.

```MLIR
%result = scalar.acoshf %input : f32
```

--------------------------------

### pool.unpin usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/pool/ops/unpin

Demonstrates the syntax for unpinning a block within a pool.

```mlir
pool.unpin %pool, %bid : pool<[%BS]>, i32
```

--------------------------------

### Vector Clampf Examples

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/clampf

Demonstrates various clamping modes including ordered, number with fastmath flags, and ieee.

```mlir
%result = vector.clampf<ordered> %value, %lower, %upper : vector<16xf32>
```

```mlir
%result = vector.clampf<number, nnan|nsz> %value, %lower, %upper : vector<16xf32>
```

```mlir
%result = vector.clampf<ieee> %value, %lower, %upper : vector<16xf32>
```

--------------------------------

### Define an empty check.case

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/check/ops/case

A basic example of a check.case operation with an empty body.

```mlir
check.case @empty {
  check.return
}
```

--------------------------------

### Constructing a strided layout with static strides

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/encoding/ops/layout-strided

Example showing the use of fully static stride values.

```mlir
%layout = encoding.layout.strided [4096, 1] : encoding<layout>
```

--------------------------------

### Link a reusable index with an invocation-local workspace

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/link_8h.html

Demonstrates the workflow for creating a linker and workspace, configuring link options, and performing a module link operation.

```c
loomc_linker_t* linker = NULL;
loomc_status_t status = loomc_linker_create(
context, NULL, loomc_allocator_system(), &linker);
if (!loomc_status_is_ok(status)) return status;
loomc_workspace_t* workspace = NULL;
status = loomc_workspace_create(NULL, loomc_allocator_system(), &workspace);
if (!loomc_status_is_ok(status)) {
loomc_linker_release(linker);
return status;
}
const loomc_string_view_t roots[] = {
loomc_make_cstring_view("@entry"),
};
loomc_link_options_t link_options = {
.type = LOOMC_STRUCTURE_TYPE_LINK_OPTIONS,
.structure_size = sizeof(loomc_link_options_t),
.link_index = library_index,
.mode = LOOMC_LINK_MODE_LINK,
.root_symbols = roots,
.root_symbol_count = 1,
};
loomc_module_t* module = NULL;
loomc_result_t* result = NULL;
status = loomc_link_module(
linker, workspace, &link_options, &module, &result);
if (loomc_status_is_ok(status) && loomc_result_succeeded(result)) {
// Pass module to compile without a text round trip.
}
loomc_result_release(result);
loomc_module_release(module);
loomc_workspace_release(workspace);
loomc_linker_release(linker);
```

--------------------------------

### Data Fields - g

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/functions_vars_g.html

List of variables starting with 'g' and their associated structures.

```APIDOC
## Data Fields - g

- **get_physical_device_features2** (loomc_spirv_vulkan_function_table_t)
- **get_physical_device_properties2** (loomc_spirv_vulkan_function_table_t)
- **global_ordinal** (loomc_module_global_t)
- **global_symbol** (loomc_module_global_query_options_t)
```

--------------------------------

### Reserve storage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/low/ops/storage-reserve

Demonstrates reserving 16 bytes of private storage with a 4-byte alignment.

```mlir
%slot = low.storage.reserve {byte_alignment = 4, byte_length = 16} : low.storage<private>
```

--------------------------------

### Create a SPIR-V target profile from a live Vulkan physical device

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/vulkan_8h.html

Initializes the function table and profile options, then invokes the creation function to generate a target profile.

```c
loomc_spirv_vulkan_function_table_t functions = {
.type = LOOMC_STRUCTURE_TYPE_SPIRV_VULKAN_FUNCTION_TABLE,
.structure_size = sizeof(loomc_spirv_vulkan_function_table_t),
.get_physical_device_properties2 =
resolved_vkGetPhysicalDeviceProperties2,
.get_physical_device_features2 = resolved_vkGetPhysicalDeviceFeatures2,
.enumerate_device_extension_properties =
resolved_vkEnumerateDeviceExtensionProperties,
};
loomc_spirv_vulkan_profile_options_t options = {
.type = LOOMC_STRUCTURE_TYPE_SPIRV_VULKAN_PROFILE_OPTIONS,
.structure_size = sizeof(loomc_spirv_vulkan_profile_options_t),
.identifier = loomc_make_cstring_view("live-vulkan-device"),
.physical_device = physical_device,
.functions = &functions,
};
loomc_target_profile_t* profile = NULL;
loomc_result_t* result = NULL;
loomc_status_t status = loomc_target_profile_create_spirv_vulkan(
target_environment, &options, loomc_allocator_system(), &profile,
&result);
if (!loomc_status_is_ok(status)) return status;
if (!loomc_result_succeeded(result)) {
// Inspect result diagnostics; profile is NULL on profile failure.
} else {
// Use profile for per-function specialization or further refinement.
loomc_target_profile_release(profile);
}
loomc_result_release(result);
```

--------------------------------

### Define a Wasm target

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/wasm/ops/target

Example of declaring a Wasm target using the simd128 configuration.

```mlir
wasm.target<simd128> @wasm
```

--------------------------------

### buffer.copy syntax example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/buffer/ops/copy

Basic syntax for the buffer.copy operation showing source and target buffers with their respective offsets and the length to copy.

```MLIR
buffer.copy %source[%source_offset], %target[%target_offset], %byte_length
```

--------------------------------

### scalar.sinturnsf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/sinturnsf

Basic syntax for applying the sinturnsf operation to a float value.

```mlir
%result = scalar.sinturnsf %input : f32
```

--------------------------------

### Floating-point subtraction example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/subf

Basic usage of the scalar.subf operation with 32-bit floating-point operands.

```mlir
%result = scalar.subf %lhs, %rhs : f32
```

--------------------------------

### Print configuration schema

Source: https://rocm.github.io/hrx-system/loom/workflows/link-and-package

Command to display the configuration keys, types, and constraints visible to the input program.

```bash
loom-link root.loom --print-config-schema
```

--------------------------------

### scalar.uitofp conversion example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/uitofp

Converts a 32-bit unsigned integer to a 32-bit floating-point value.

```mlir
%result = scalar.uitofp %input : i32 to f32
```

--------------------------------

### Retrieve configuration values

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/config/ops/get

Examples of reading various configuration types using the config.get operation.

```MLIR
%hidden = config.get @model36.model.hidden_size : index
```

```MLIR
%enabled = config.get @model36.features.enable_mtp : i1
```

```MLIR
%dense_encoding = config.get @model36.quant.dense_encoding : encoding<schema>
```

--------------------------------

### Floating-point negation example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/negf

Basic usage of the scalar.negf operation with a 32-bit floating-point value.

```MLIR
%result = scalar.negf %input : f32
```

--------------------------------

### scalar.maxnumf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/maxnumf

Basic syntax for the maxnumf operation using f32 types.

```MLIR
%result = scalar.maxnumf %lhs, %rhs : f32
```

--------------------------------

### Atomic Reduce Operation Example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/atomic-reduce

Demonstrates the syntax for an atomic addition reduction using relaxed ordering and workgroup scope.

```mlir
vector.atomic.reduce<addi> %v, %view[%row, %col][%offsets] {ordering = relaxed, scope = workgroup} : vector<4xi32>, view<[%m]x[%n]xi32, %layout>, vector<4xindex>
```

--------------------------------

### Exercise Loom source through different devices

Source: https://rocm.github.io/hrx-system/loom/workflows/oracles/radv

Use these commands to run a Loom kernel through both Vulkan and native AMDGPU paths for portability testing.

```bash
iree-test-loom kernel.loom \
  --case=@production_cut_case \
  --device=vulkan

iree-test-loom kernel.loom \
  --case=@production_cut_case \
  --device=amdgpu
```

--------------------------------

### scalar.powf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/powf

Basic syntax for the scalar.powf operation using f32 types.

```MLIR
%result = scalar.powf %lhs, %rhs : f32
```

--------------------------------

### scalar.ori usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/ori

Performs a bitwise OR operation on two 32-bit integer operands.

```MLIR
%result = scalar.ori %lhs, %rhs : i32
```

--------------------------------

### Load a launch-config program

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/launch__config_8h.html

Loads and verifies a compiled launch-config program artifact.

```c
loomc_status_t loomc_launch_config_program_load(const loomc_artifact_t *artifact, loomc_artifact_release_fn_t release, void *release_user_data, loomc_allocator_t allocator, loomc_launch_config_program_t **out_program)
```

--------------------------------

### scalar.minui usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/minui

Basic usage of the scalar.minui operation with 32-bit integer operands.

```mlir
%result = scalar.minui %lhs, %rhs : i32
```

--------------------------------

### Basic usage of scf.schedule.fence

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scf/ops/schedule-fence

Demonstrates the syntax for invoking the fence operation.

```mlir
scf.schedule.fence
```

--------------------------------

### Unsigned integer division example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/divui

Performs division on two 32-bit unsigned integers.

```MLIR
%result = scalar.divui %lhs, %rhs : i32
```

--------------------------------

### Pass a framework-owned option dictionary

Source: https://rocm.github.io/hrx-system/loom/reference/c-api/generated/emit_8h.html

Demonstrates how to provide an ordered dictionary of loose string options to the emission process, allowing for identifier overrides.

```c
const loomc_option_entry_t entries[] = {
{loomc_make_cstring_view(LOOMC_EMIT_OPTION_KEY_IDENTIFIER),
loomc_make_cstring_view("default.spv")},
{loomc_make_cstring_view(LOOMC_EMIT_OPTION_KEY_IDENTIFIER),
loomc_make_cstring_view("override.spv")},
};
loomc_option_dict_t dict = {
.type = LOOMC_STRUCTURE_TYPE_OPTION_DICT,
.structure_size = sizeof(loomc_option_dict_t),
.entries = entries,
.entry_count = 2,
};
loomc_emit_options_t emit_options = {
.type = LOOMC_STRUCTURE_TYPE_EMIT_OPTIONS,
.structure_size = sizeof(loomc_emit_options_t),
.next = &dict,
};
```

--------------------------------

### Inspect the generated plan

Source: https://rocm.github.io/hrx-system/loom/workflows/benchmark

Use jq to parse and view the summary, benchmarks, and work items from the generated plan file.

```bash
jq '{summary, benchmarks, work_items}' plan.json
```

--------------------------------

### scalar.floorf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/floorf

Basic usage of the scalar.floorf operation with a 32-bit floating-point input.

```MLIR
%result = scalar.floorf %input : f32
```

--------------------------------

### scalar.isnanf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/isnanf

Basic syntax for applying the isnanf operation to a 32-bit float.

```MLIR
%result = scalar.isnanf %input : f32
```

--------------------------------

### Create a Target Pipeline

Source: https://rocm.github.io/hrx-system/loom/integration/jit-kernel

Initializes a target pipeline with specific options to prepare the program for compilation.

```c
  loomc_target_pipeline_options_t pipeline_options = {
      .type = LOOMC_STRUCTURE_TYPE_TARGET_PIPELINE_OPTIONS,
      .structure_size = sizeof(pipeline_options),
      .identifier = loomc_make_cstring_view("guide-prepared-low"),
      .kind = LOOMC_TARGET_PIPELINE_KIND_PREPARED_LOW,
      .source_to_low_max_errors = 20,
  };
  if (loomc_status_is_ok(status)) {
    status = loomc_pass_program_create_from_target_pipeline(
        state->context, &pipeline_options, loomc_allocator_system(),
        &state->pass_program, &state->result);
  }
  if (loomc_status_is_ok(status)) {
    status = require_successful_result(state->result,
                                       "target pipeline preparation failed");
  }
  if (loomc_status_is_ok(status)) {
    jit_kernel_state_reset_result(state);
  }
```

--------------------------------

### Float precision extension example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/extf

Demonstrates the syntax for extending a float from f16 to f32.

```MLIR
%result = scalar.extf %input : f16 to f32
```

--------------------------------

### Atomic reduction operation example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/view/ops/atomic-reduce

Demonstrates the syntax for performing an atomic addition reduction on a 2D view.

```mlir
view.atomic.reduce<addi> %value, %view[%row, %col] {ordering = relaxed, scope = workgroup} : i32, view<[%M]x[%N]xi32, %layout>
```

--------------------------------

### Basic kernel.launch.concurrent usage

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/kernel/ops/launch-concurrent

A minimal example showing the structure of a concurrent launch block with a required yield terminator.

```mlir
kernel.launch.concurrent {
  kernel.launch.yield
}
```

--------------------------------

### Signed integer division example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/divsi

Basic usage of the scalar.divsi operation with 32-bit integers.

```MLIR
%result = scalar.divsi %lhs, %rhs : i32
```

--------------------------------

### scalar.expf usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/expf

Basic usage of the scalar.expf operation with a 32-bit float input.

```MLIR
%result = scalar.expf %input : f32
```

--------------------------------

### Configure and link a module

Source: https://rocm.github.io/hrx-system/loom/workflows/link-and-package

Applies configuration bindings to a module during the linking process to prune providers and resolve requirements.

```bash
loom-link root.loom \
  --library=providers.loom \
  --root=@entry \
  --config=model.hidden_size=4096 \
  --require-resolved-config \
  --output=entry.loom
```

--------------------------------

### Atomic Compare-Exchange Example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/vector/ops/atomic-cmpxchg

Demonstrates the syntax for performing an atomic compare-exchange operation on a vector of 4 i32 elements within a specified view.

```mlir
%old = vector.atomic.cmpxchg %expected, %replacement, %view[%row, %col][%offsets] {success_ordering = acq_rel, failure_ordering = acquire, scope = workgroup} : vector<4xi32>, view<[%m]x[%n]xi32, %layout>, vector<4xindex> -> vector<4xi32>
```

--------------------------------

### scalar.expm1f usage example

Source: https://rocm.github.io/hrx-system/loom/reference/dialects/scalar/ops/expm1f

Basic syntax for applying the expm1f operation to a float input.

```MLIR
%result = scalar.expm1f %input : f32
```
Tokens:
Raw