### StableHLO dot_general Algorithm Examples Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20240315-algorithm-support.md Examples demonstrating the structure and values of the DotAlgorithm attribute for the dot_general operation in StableHLO. These examples illustrate different precision configurations, component counts, and accumulation strategies. ```text // Inputs are casted to tf32, and then accumulated in f32: { lhs_type = tf32, rhs_type = tf32, accum_type = f32, lhs_component_count = 1, rhs_component_count = 1, num_primitive_ops = 1, allow_imprecise_accum = false } // bf16_6x: each input is decomposed to 3 bf16 components, then 6 dot operations are done on those components, and the result is accumulated in f32. { lhs_type = bf16, rhs_type = bf16, accum_type = f32, lhs_component_count = 3, rhs_component_count = 3, num_primitive_ops = 6, allow_imprecise_accum = false } // Inputs are (casted to) f8e5m2, and we accumulate in f32, but for some steps we may accumulate in lower precision. { lhs_type = f8e5m2, rhs_type = f8e5m2, accum_type = f32, lhs_component_count = 1, rhs_component_count = 1, num_primitive_ops = 1, allow_imprecise_accum = true } ``` -------------------------------- ### VHLO Op Versioning Example (TableGen) Source: https://github.com/openxla/stablehlo/blob/main/docs/vhlo.md This example demonstrates how different versions of a StableHLO op are represented in the VHLO dialect using TableGen syntax. It shows how an op's evolution is captured by creating new versions (e.g., VHLO_MyOpV1, VHLO_MyOpV2) with distinct version ranges and attribute sets, ensuring that older versions remain accessible. ```tablegen def VHLO_MyOpV1 : VHLO_Op<"my_op_v1", "0.9.0", "0.10.0"> { let arguments = (ins VHLO_AnyType:$operand ); let results = (outs VHLO_AnyType:$result); } def VHLO_MyOpV2 : VHLO_Op<"my_op_v2", "0.11.0", "current"> { let arguments = (ins VHLO_AnyType:$operand, VHLO_AnyAttr:$attr // New attribute added to StableHLO in 0.11.0 ); let results = (outs VHLO_AnyType:$result); } ``` -------------------------------- ### Install StableHLO Nightly Wheels Source: https://github.com/openxla/stablehlo/blob/main/README.md Installs StableHLO using pip from nightly wheels available on the GitHub Releases page. ```shell pip install stablehlo -f https://github.com/openxla/stablehlo/releases/expanded_assets/dev-wheels ``` -------------------------------- ### Install JAX and StableHLO dependencies Source: https://github.com/openxla/stablehlo/blob/main/docs/tutorials/jax-export.ipynb Installs the required Python packages including JAX, Flax, Transformers, and TensorFlow for model export workflows. ```bash !pip install -U jax jaxlib flax transformers tf-nightly ``` -------------------------------- ### ExampleAdd.cpp Source: https://github.com/openxla/stablehlo/blob/main/examples/c++/README.md A simple C++ example demonstrating how to use the StableHLO library to add two numbers and interpret a StableHLO program with concrete inputs. ```APIDOC ## C++ Example: Interpreting StableHLO Addition ### Description This example shows how to use the StableHLO library in C++ to perform addition on a StableHLO program. It demonstrates interpreting a given MLIR module containing StableHLO operations with specified concrete inputs. ### Method N/A (This is a code example, not an API endpoint) ### Endpoint N/A ### Parameters N/A ### Request Example ```cpp // Assume 'module' is an MLIR module with function "main" containing StableHLO // operations. llvm::outs() << "Program:\n " << module << "\n"; // Create concrete inputs to be used for interpreting "main". auto inputValue1 = mlir::DenseElementsAttr::get( tensorType, block_builder.getFloatAttr(tensorType.getElementType(), static_cast(10))); auto inputValue2 = mlir::DenseElementsAttr::get( tensorType, block_builder.getFloatAttr(tensorType.getElementType(), static_cast(20))); llvm::outs() << "Inputs: " << inputValue1 << ", " << inputValue2 << "\n"; mlir::stablehlo::InterpreterConfiguration config; auto results = evalModule(module, {inputValue1, inputValue2}, config); llvm::outs() << "Output: " << (*results)[0]; ``` ### Response #### Success Response (Output) - **Output** (tensor<3x4xf32>) - The result of the StableHLO addition operation. #### Response Example ```mlir Program: module @test_module { func.func @main(%arg0: tensor<3x4xf32>, %arg1: tensor<3x4xf32>) -> tensor<3x4xf32> { %0 = stablehlo.add %arg0, %arg1 : tensor<3x4xf32> %1 = stablehlo.add %arg0, %arg1 : tensor<3x4xf32> return %1 : tensor<3x4xf32> } } Inputs: dense<1.000000e+01> : tensor<3x4xf32>, dense<2.000000e+01> : tensor<3x4xf32> Output: dense<3.000000e+01> : tensor<3x4xf32> ``` ``` -------------------------------- ### Install PyTorch and XLA dependencies Source: https://github.com/openxla/stablehlo/blob/main/docs/tutorials/pytorch-export.ipynb Installs the necessary packages including torch, torchvision, and torch_xla to enable model export functionality. ```bash !pip install torch_xla==2.5.0 torch==2.5.0 torchvision==0.20.0 tensorflow-cpu ``` -------------------------------- ### Install StableHLO and TensorFlow dependencies Source: https://github.com/openxla/stablehlo/blob/main/docs/tutorials/savedmodel-embed.ipynb Installs the necessary Python packages for StableHLO and TensorFlow CPU support via pip. ```bash !pip install stablehlo -f https://github.com/openxla/stablehlo/releases/expanded_assets/dev-wheels !pip install tensorflow-cpu ``` -------------------------------- ### StableHLO Sparse Tensor Encoding Examples Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20230210-sparsity.md Examples demonstrating how to define sparse tensor encodings in StableHLO, including CSR, CSC, and BSR formats. These examples illustrate the mapping between tensor dimensions and storage levels. ```stablehlo tensor<10x20xf32, #CSR> #CSR = #sparse_tensor.encoding< (i, j) -> (i : dense, j : compressed) > ``` ```stablehlo tensor<10x20xf32, #CSC> #CSC = #sparse_tensor.encoding< (i, j) -> (j : dense, i : compressed) > ``` ```stablehlo tensor<10x20xf32, #BSR> #BSR = #sparse_tensor.encoding< (i, j) -> ( i floordiv 2 : dense , j floordiv 3 : compressed ) ``` -------------------------------- ### Execute math generation scripts Source: https://github.com/openxla/stablehlo/blob/main/build_tools/math/README.md Examples of running Python scripts to generate StableHLO math patterns and corresponding MLIR test files. ```bash python build_tools/math/generate_ChloDecompositionPatternsMath.py python build_tools/math/generate_tests.py ``` -------------------------------- ### Example StableHLO Operation (MLIR) Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md An illustrative example of a StableHLO operation, specifically 'select_and_scatter', shown in MLIR format. It demonstrates the consumption of input values, functions, and attributes, along with the op signature. ```mlir ... ``` -------------------------------- ### Install Dependencies and Run Tests Source: https://github.com/openxla/stablehlo/blob/main/stablehlo/integrations/python/stablehlo/savedmodel/README.md Instructions for setting up the necessary environment and running the provided test script for the StableHLO to TensorFlow SavedModel conversion. Requires TensorFlow CPU and the StableHLO Python package. ```shell pip install tensorflow-cpu PYTHONPATH="./build/python_packages/stablehlo" python3 ../../tests/stablehlo_to_tf_saved_model_test.py ``` -------------------------------- ### StableHLO Recv Operation Example Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20250421-source-target-pairs.md Demonstrates the usage of the stablehlo.recv operation. It defines source-target pairs and a channel handle to facilitate data reception between devices. ```mlir %results0, %results1 = "stablehlo.recv"(%token) { source_target_pairs = dense<[[0, 1], [1, 2]]> : tensor<2x2xi64>, channel_handle = #stablehlo.channel_handle, is_host_transfer = false } : (!stablehlo.token) -> (tensor<2x2xi64>, !stablehlo.token) ``` -------------------------------- ### Install functional_algorithms dependencies Source: https://github.com/openxla/stablehlo/blob/main/build_tools/math/README.md Commands to install the required functional_algorithms package using pip or mamba/conda. ```bash pip install functional_algorithms ``` ```bash mamba install -c conda-forge functional_algorithms ``` -------------------------------- ### Example StableHLO Quantized Test Case Source: https://github.com/openxla/stablehlo/blob/main/docs/quantization.md A typical example of a StableHLO quantized test case demonstrating the use of uniform_quantize, add, and uniform_dequantize operations. It includes input data, golden output, and a custom_call for checking equality, validating quantization behavior. ```mlir func.func @main() -> tensor<11xf32> { %operand_0 = stablehlo.constant dense<...> : tensor<11xf32> %operand_1 = stablehlo.constant dense<...> : tensor<11xf32> %golden = stablehlo.constant dense<...> : tensor<11xf32> %0 = stablehlo.uniform_quantize %operand_0 : (tensor<11xf32>) -> tensor<11x!quant.uniform> %1 = stablehlo.uniform_quantize %operand_1 : (tensor<11xf32>) -> tensor<11x!quant.uniform> %2 = stablehlo.add %1, %0 : tensor<11x!quant.uniform> %result = stablehlo.uniform_dequantize %2 : (tensor<11x!quant.uniform>) -> tensor<11xf32> %4 = stablehlo.custom_call @check.eq(%golden, %result) : (tensor<11xf32>, tensor<11xf32>) -> tensor return %3 : tensor<11xf32> } ``` -------------------------------- ### StableHLO All-Reduce Operation Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Documentation for the StableHLO all_reduce operation, detailing its inputs, outputs, constraints, and providing an MLIR example. ```APIDOC ## ALL_REDUCE /stablehlo/docs/ops/all_reduce.md ### Description Performs an all-reduce operation across replicas. ### Method N/A (This is a conceptual operation description, not an API endpoint) ### Endpoint N/A ### Parameters #### Inputs - **operands** (variadic tensors) - Input tensors for the operation. - **replica_groups** (tensor) - Defines the groups of replicas for the reduction. - **channel_id** (constant si64) - Identifier for the communication channel. - **use_global_device_ids** (constant i1) - Flag to use global device IDs. - **computation** (function) - The reduction computation to perform. #### Outputs - **results** (variadic tensors) - Output tensors after the all-reduce operation. ### Constraints - (C1) `is_unique(replica_groups)`. - (C2) `size(replica_groups)` is defined based on usage (cross_replica, cross_replica_and_partition, flattened_ids). - (C3) `0 <= replica_groups < size(replica_groups)`. - (C4) If `use_global_device_ids = true`, then `channel_id > 0`. - (C5) `computation` has a specific type signature. - (C6) `shape(results...) = shape(operands...)`. - (C7) `element_type(results...) = E`. ### Request Example ```mlir // Example for all_reduce operation %result:2 = "stablehlo.all_reduce"(%operand0, %operand0) ({ ^bb0(%arg0: tensor, %arg1: tensor): %0 = "stablehlo.add"(%arg0, %arg1) : (tensor, tensor) -> tensor "stablehlo.return"(%0) : (tensor) -> () }) { replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>, channel_handle = #stablehlo.channel_handle } : (tensor<4xi64>, tensor<4xi64>) -> (tensor<4xi64>, tensor<4xi64>) ``` ### Response #### Success Response (200) - **results** (variadic tensors) - The results of the all-reduce operation. #### Response Example ```json { "results": [ // Tensor data for results ] } ``` ``` -------------------------------- ### Define DotAlgorithm Attributes Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Examples of DotAlgorithm configurations used to specify precision types, accumulation types, and decomposition strategies for dot operations in StableHLO. ```text // Inputs are casted to tf32, and then accumulated in f32: {lhs_precision_type = tf32, rhs_precision_type = tf32, accumulation_type = f32, lhs_component_count = 1, rhs_component_count = 1, num_primitive_operations = 1, allow_imprecise_accumulation = false} // bf16_6x: each input is decomposed to 3 bf16 components, then 6 dot operations are done on those components, and the result is accumulated in f32. {lhs_precision_type = bf16, rhs_precision_type = bf16, accumulation_type = f32, lhs_component_count = 3, rhs_component_count = 3, num_primitive_operations = 6, allow_imprecise_accumulation = false} ``` -------------------------------- ### CMake: Setup Sanitizers and Python Bindings Source: https://github.com/openxla/stablehlo/blob/main/CMakeLists.txt Includes and configures sanitizers for the build using the SetupSanitizers module. It also sets up Python development packages and configures the nanobind domain for StableHLO if Python bindings are enabled. ```cmake include(SetupSanitizers) setup_sanitizers() if(STABLEHLO_ENABLE_BINDINGS_PYTHON) include(MLIRDetectPythonEnv) mlir_configure_python_dev_packages() if(NOT MLIR_BINDINGS_PYTHON_NB_DOMAIN) set(MLIR_BINDINGS_PYTHON_NB_DOMAIN "stablehlo" CACHE STRING "nanobind domain for StableHLO" FORCE) endif() endif() ``` -------------------------------- ### Implement StableHLO Main Function Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md An example of a StableHLO function named @main that performs matrix operations on tensors. It demonstrates input/output signatures and the usage of stablehlo operations like reshape, dot, and add. ```mlir func.func @main( %image: tensor<28x28xf32>, %weights: tensor<784x10xf32>, %bias: tensor<1x10xf32> ) -> tensor<1x10xf32> { %0 = "stablehlo.reshape"(%image) : (tensor<28x28xf32>) -> tensor<1x784xf32> %1 = "stablehlo.dot"(%0, %weights) : (tensor<1x784xf32>, tensor<784x10xf32>) -> tensor<1x10xf32> %2 = "stablehlo.add"(%1, %bias) : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32> %3 = "stablehlo.constant"() {value = dense<0.0> : tensor<1x10xf32>} : () -> tensor<1x10xf32> %4 = "stablehlo.maximum"(%2, %3) : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32> "func.return"(%4): (tensor<1x10xf32>) -> () } ``` -------------------------------- ### Manage StableHLO Serialization with C API Source: https://context7.com/openxla/stablehlo/llms.txt Provides examples of using the StableHLO C API to query version information, handle compatibility requirements, and perform serialization/deserialization of portable artifacts. ```c #include "stablehlo/integrations/c/StablehloDialectApi.h" #include "mlir-c/IR.h" void stringCallback(MlirStringRef str, void* userData) { char** result = (char**)userData; *result = malloc(str.length + 1); memcpy(*result, str.data, str.length); (*result)[str.length] = '\0'; } int main() { MlirContext ctx = mlirContextCreate(); int apiVersion = stablehloGetApiVersion(); char* currentVersion = NULL; stablehloGetCurrentVersion(stringCallback, ¤tVersion); const char* moduleStr = "func.func @main(%arg0: tensor) -> tensor { return %arg0 : tensor }"; char* artifact = NULL; MlirLogicalResult result = stablehloSerializePortableArtifactFromStringRef( mlirStringRefCreateFromCString(moduleStr), mlirStringRefCreateFromCString(currentVersion), stringCallback, &artifact ); if (mlirLogicalResultIsSuccess(result)) { char* deserialized = NULL; stablehloDeserializePortableArtifact(mlirStringRefCreateFromCString(artifact), stringCallback, &deserialized); } free(currentVersion); free(artifact); mlirContextDestroy(ctx); return 0; } ``` -------------------------------- ### StableHLO dot_general Operation Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Demonstrates the usage of the stablehlo.dot_general operation with specified dot dimensions, precision configuration, and algorithm details. It takes two tensors as input and produces a resulting tensor. ```mlir // %lhs: [ // [[1, 2], // [3, 4]], // [[5, 6], // [7, 8]] // ] // %rhs: [ // [[1, 0], // [0, 1]], // [[1, 0], // [0, 1]] // ] %result = "stablehlo.dot_general"(%lhs, %rhs) { dot_dimension_numbers = #stablehlo.dot< lhs_batching_dimensions = [0], rhs_batching_dimensions = [0], lhs_contracting_dimensions = [2], rhs_contracting_dimensions = [1] >, precision_config = [#stablehlo, #stablehlo], algorithm = #stablehlo.dot_algorithm< lhs_precision_type = tf32, rhs_precision_type = tf32, accumulation_type = f32, lhs_component_count = 1, rhs_component_count = 1, num_primitive_operations = 1, allow_imprecise_accumulation = false > } : (tensor<2x2x2xi64>, tensor<2x2x2xi64>) -> tensor<2x2x2xi64> // %result: [ // [[1, 2], // [3, 4]], // [[5, 6], // [7, 8]] // ] ``` -------------------------------- ### StableHLO TanOp Example Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20240312-tanop.md Demonstrates the usage of the proposed 'stablehlo.tan' operation with a 2x2 tensor of 64-bit floating-point numbers. It shows the input operand and the expected output result after applying the tangent function element-wise. ```mlir // %operand: [ // [0.0, 1.57079632], // [0, pi/2] // [3.14159265, 4.71238898] // [pi, 3pi/2] // ] %result = "stablehlo.tan"(%operand) : (tensor<2x2xf64>) -> tensor<2x2xf64> // %result: [ // [0.0, 1.63312e+16], // [0.0, 5.44375e+15] // ] ``` -------------------------------- ### StableHLO All-to-All Operation Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Demonstrates the usage of the stablehlo.all_to_all operation, which redistributes tensor slices across replicas. It specifies dimensions for splitting and concatenation, along with replica group configurations. ```mlir %result:2 = "stablehlo.all_to_all"(%operand1, %operand2) { split_dimension = 1 : i64, concat_dimension = 0 : i64, split_count = 2 : i64, replica_groups = dense<[[0, 1]]> : tensor<1x2xi64> // channel_id = 0 } : (tensor<2x4xi64>, tensor<2x4xi64>) -> (tensor<4x2xi64>, tensor<4x2xi64>) ``` -------------------------------- ### JAX Dot Product with Float8E5M2FNUZ Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20230321-fp8_fnuz.md An example of using JAX to perform a dot product operation using the Float8E5M2FNUZ data type, demonstrating the resulting MLIR and the impact of lower precision on the output. ```python import jax import jax.numpy as jnp x = jnp.arange(16, dtype=jnp.float8_e5m2fnuz) # MLIR for iota # module @jit_iota { # func.func public @main() -> tensor<16xf8E5M2FNUZ> { # %0 = stablehlo.iota dim = 0 : tensor<16xf8E5M2FNUZ> # return %0 : tensor<16xf8E5M2FNUZ> # } # } print(x) # MLIR for dot_general # module @jit_matmul { # func.func public @main(%arg0: tensor<16xf8E5M2FNUZ>, %arg1: tensor<16xf8E5M2FNUZ>) -> tensor { # %0 = "stablehlo.dot_general"(%arg0, %arg1) {dot_dimension_numbers = #stablehlo.dot, precision_config = [#stablehlo, #stablehlo]} : (tensor<16xf8E5M2FNUZ>, tensor<16xf8E5M2FNUZ>) -> tensor # return %0 : tensor # } # } result = x @ x print(result) ``` -------------------------------- ### StableHLO Gather Operation Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Demonstrates the usage of the stablehlo.gather operation to extract slices from an operand tensor based on start indices. It specifies dimension numbers, slice sizes, and whether indices are sorted. The operation takes an operand tensor and start indices tensor as input and produces a resulting tensor. ```mlir // %operand: [ // [ // [[1, 2], [3, 4], [5, 6], [7, 8]], // [[9, 10],[11, 12], [13, 14], [15, 16]], // [[17, 18], [19, 20], [21, 22], [23, 24]] // ], // [ // [[25, 26], [27, 28], [29, 30], [31, 32]], // [[33, 34], [35, 36], [37, 38], [39, 40]], // [[41, 42], [43, 44], [45, 46], [47, 48]] // ] // ] // %start_indices: [ // [ // [[0, 0], [1, 0], [2, 1]], // [[0, 1], [1, 1], [0, 9]] // ], // [ // [[0, 0], [2, 1], [2, 2]], // [[1, 2], [0, 1], [1, 0]] // ] // ] %result = "stablehlo.gather"(%operand, %start_indices) { dimension_numbers = #stablehlo.gather< offset_dims = [3, 4], collapsed_slice_dims = [1], operand_batching_dims = [0], start_indices_batching_dims = [1], start_index_map = [2, 1], index_vector_dim = 3>, slice_sizes = array, indices_are_sorted = false } : (tensor<2x3x4x2xi32>, tensor<2x2x3x2xi64>) -> tensor<2x2x3x2x2xi32> // %result: [ // [ // [ // [[1, 2], [3, 4]], // [[3, 4], [5, 6]], // [[13, 14], [15, 16]] // ], // [ // [[33, 34], [35, 36]], // [[35, 36], [37, 38]], // [[41, 42], [43, 44]] // ] // ], // [ // [ // [[1, 2], [3, 4]], // [[13, 14], [15, 16]], // [[21, 22], [23, 24]] // ], // [ // [[43, 44], [45, 46]], // [[33, 34], [35, 36]], // [[27, 28], [29, 30]] // ] // ] // ] ``` -------------------------------- ### StableHLO Get Tuple Element Operation Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Shows the stablehlo.get_tuple_element operation used to extract a specific element from a tuple. The operation requires a tuple and an index, returning the element at that index. Note: This operation is being considered for deprecation. ```mlir // %operand: ([1.0, 2.0], (3)) %result = "stablehlo.get_tuple_element"(%operand) <{index = 0 : i32}> : (tuple, tuple>>) -> tensor<2xf64> // %result: [1.0, 2.0] ``` -------------------------------- ### StableHLO Get Dimension Size Operation Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Illustrates the stablehlo.get_dimension_size operation, which returns the size of a specified dimension of an operand tensor. The operation takes a tensor and a dimension index as input and outputs a scalar tensor representing the dimension size. ```mlir // %operand: [[1, 2, 3], [4, 5, 6]] %result = "stablehlo.get_dimension_size"(%operand) { dimension = 1 : i64 } : (tensor<2x3xi64>) -> tensor // %result: 3 ``` -------------------------------- ### Dynamic Slice Operation in MLIR Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Demonstrates the stablehlo.dynamic_slice operation which extracts a slice from an operand tensor based on start indices and slice sizes. It's useful for creating sub-tensors dynamically. The example shows slicing a 4x4 tensor into a 2x2 tensor. ```mlir // %operand: [ // [0, 0, 1, 1], // [0, 0, 1, 1], // [0, 0, 0, 0], // [0, 0, 0, 0] // ] // %start_indices0: -1 // %start_indices1: 3 %result = "stablehlo.dynamic_slice"(%operand, %start_indices0, %start_indices1) { slice_sizes = array } : (tensor<4x4xi32>, tensor, tensor) -> tensor<2x2xi32> // %result: [ // [1, 1], // [1, 1] // ] ``` -------------------------------- ### Dynamic Update Slice Operation in MLIR Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Illustrates the stablehlo.dynamic_update_slice operation, which updates a slice of an operand tensor with new values from an update tensor. The operation defines how the result tensor is constructed based on the operand, update, and start indices. The example updates a portion of a 4x4 tensor. ```mlir // %operand: [ // [1, 1, 0, 0], // [1, 1, 0, 0], // [1, 1, 1, 1], // [1, 1, 1, 1] // ] // %update: [ // [1, 1], // [1, 1] // ] // %start_indices0: -1 // %start_indices1: 3 %result = "stablehlo.dynamic_update_slice"(%operand, %update, %start_indices0, %start_indices1) : (tensor<4x4xi32>, tensor<2x2xi32>, tensor, tensor) -> tensor<4x4xi32> // %result: [ // [1, 1, 1, 1], // [1, 1, 1, 1], // [1, 1, 1, 1], // [1, 1, 1, 1] // ] ``` -------------------------------- ### Build StableHLO Programs with C++ API Source: https://context7.com/openxla/stablehlo/llms.txt Demonstrates how to programmatically construct an MLIR module, define function signatures, add operations like stablehlo.add, and interpret the result using the StableHLO reference interpreter. ```cpp #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/Verifier.h" #include "stablehlo/dialect/StablehloOps.h" #include "stablehlo/reference/Api.h" #include "stablehlo/reference/Configuration.h" int main() { mlir::MLIRContext context; // Create module mlir::OwningOpRef module = mlir::ModuleOp::create(mlir::UnknownLoc::get(&context)); module->getContext()->loadDialect(); module->getContext()->loadDialect(); // Define tensor type: tensor<3x4xf32> auto tensorType = mlir::RankedTensorType::get( {3, 4}, mlir::Float32Type::get(&context)); // Create function signature: (tensor<3x4xf32>, tensor<3x4xf32>) -> tensor<3x4xf32> auto funcType = mlir::FunctionType::get( &context, {tensorType, tensorType}, {tensorType}); // Create function auto function = mlir::func::FuncOp::create( mlir::UnknownLoc::get(&context), "main", funcType, {}); function.setVisibility(mlir::func::FuncOp::Visibility::Public); module->push_back(function); // Create function body mlir::Block* block = function.addEntryBlock(); mlir::OpBuilder builder = mlir::OpBuilder::atBlockEnd(block); mlir::Location loc = builder.getUnknownLoc(); // Add stablehlo.add operation llvm::SmallVector args(block->args_begin(), block->args_end()); auto addOp = mlir::stablehlo::AddOp::create(builder, loc, args, {}); // Add return mlir::func::ReturnOp::create(builder, loc, addOp->getResult(0)); // Verify the module assert(mlir::succeeded(mlir::verify(module.get()))); // Interpret with concrete values auto getConstant = [&](float val) { return mlir::DenseElementsAttr::get( tensorType, builder.getFloatAttr(tensorType.getElementType(), val)); }; mlir::stablehlo::InterpreterConfiguration config; auto results = evalModule(*module, {getConstant(10.0), getConstant(20.0)}, config); return 0; } ``` -------------------------------- ### CMake: Include Directories and Definitions Source: https://github.com/openxla/stablehlo/blob/main/CMakeLists.txt Configures include paths and definitions for the project, pointing to LLVM, MLIR, and current source/binary directories. It also sets up library directories and definitions from LLVM. ```cmake include_directories(${LLVM_INCLUDE_DIRS}) include_directories(${MLIR_INCLUDE_DIRS}) include_directories(${CMAKE_CURRENT_SOURCE_DIR}) include_directories(${CMAKE_CURRENT_BINARY_DIR}) link_directories(${LLVM_BUILD_LIBRARY_DIR}) add_definitions(${LLVM_DEFINITIONS}) ``` -------------------------------- ### Establish execution dependencies with stablehlo.after_all Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Creates a data dependency token to ensure that specified input operations complete before subsequent dependent operations execute. ```mlir // %input0: !stablehlo.token // %input1: !stablehlo.token %result = "stablehlo.after_all"(%input0, %input1) : (!stablehlo.token, !stablehlo.token) -> !stablehlo.token ``` -------------------------------- ### StableHLO Sequential Execution Example Source: https://github.com/openxla/stablehlo/blob/main/docs/spec.md Demonstrates a simple StableHLO program for sequential execution. It defines a main function that takes no inputs and returns a f64 tensor computed by adding two constants. The execution order is dataflow-aligned. ```mlir func.func @main() -> tensor { %0 = stablehlo.constant dense<1.0> : tensor %1 = stablehlo.constant dense<2.0> : tensor %2 = stablehlo.add %0, %1 : tensor return %2 : tensor } ``` -------------------------------- ### Execute StableHLO Programs with Reference Interpreter Source: https://context7.com/openxla/stablehlo/llms.txt Demonstrates how to parse a StableHLO module and execute it using the reference interpreter with concrete NumPy input tensors. It covers dialect registration, module parsing, input attribute creation, and result retrieval. ```python from mlir import ir from mlir.dialects import stablehlo import numpy as np PROGRAM = """ func.func @test(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> { %0 = stablehlo.add %arg0, %arg0 : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32> func.return %0 : tensor<2x2xf32> } """ with ir.Context() as context: stablehlo.register_dialect(context) stablehlo.register_interpreter_dialect(context) module = ir.Module.parse(PROGRAM) input_data = np.array([[1.0, 2.0], [3.0, 4.0]], dtype=np.float32) input_attr = ir.DenseElementsAttr.get(input_data) results = stablehlo.eval_module(module, [input_attr]) output = np.array(results[0]) print(f"Input:\n{input_data}") print(f"Output (input + input):\n{output}") ``` -------------------------------- ### StableHLO Result Accuracy Attribute Examples Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20241015-result-accuracy.md Provides examples of how to use the '#stablehlo.result_accuracy' attribute with different configurations, including default values, highest accuracy, and numerical tolerance specifications. ```text #stablehlo.result_accuracy #stablehlo.result_accuracy #stablehlo.result_accuracy ``` -------------------------------- ### JAX Float8 Dot Product Example Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20240808-f8E4M3_f8E3M4.md Demonstrates a simple dot product computation using Float8E4M3 data type in JAX. It shows the MLIR generated for the operation and the resulting array, highlighting potential precision differences due to the lower-precision format. ```python >>> import jax >>> import jax.numpy as jnp >>> x = jnp.arange(8, dtype=jnp.float8_e4m3) module @jit_iota { func.func public @main() -> tensor<8xf8E4M3> { %0 = stablehlo.iota dim = 0 : tensor<8xf8E4M3> return %0 : tensor<8xf8E4M3> } } >>> x Array([0, 1, 2, 3, 4, 5, 6, 7], dtype=float8_e4m3) >>> x @ x module @jit_matmul { func.func public @main(%arg0: tensor<8xf8E4M3> {mhlo.sharding = ""}, %arg1: tensor<8xf8E4M3> {mhlo.sharding = ""}) -> tensor { %0 = "stablehlo.dot_general"(%arg0, %arg1) {dot_dimension_numbers = #stablehlo.dot, precision_config = [#stablehlo, #stablehlo]} : (tensor<8xf8E4M3>, tensor<8xf8E4M3>) -> tensor return %0 : tensor } } Array(144, dtype=float8_e4m3) ``` -------------------------------- ### Custom Call for Buffer Operations (MLIR) Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20250729-buffer.md Demonstrates the usage of 'stablehlo.custom_call' for buffer operations in MLIR. It includes examples for creating uninitialized buffers ('CreateBuffer'), initializing buffers ('Pin'), and deallocating buffers ('Unpin'), showcasing their respective inputs and outputs. ```mlir %uninitialized_buffer = "stablehlo.custom_call"() { call_target_name = "CreateBuffer", api_version = 4 : i32, } : () -> memref<4xf64> %initialized_buffer = "stablehlo.custom_call"(%init_value) { call_target_name = "Pin", api_version = 4 : i32, } : (tensor<4xf64>) -> memref<4xf64> %dealloc_buffer = "stablehlo.custom_call"(%initialized_buffer) { call_target_name = "Unpin", api_version = 4 : i32, } : (memref<4xf64>) -> tensor<4xf64> ``` -------------------------------- ### MLIR Collective Broadcast Example Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20231017-collective-broadcast.md This MLIR snippet demonstrates the usage of the stablehlo.collective_broadcast operation. It shows how to configure replica groups and channel handles for broadcasting a tensor across replicas. The example specifies the number of replicas and partitions, the initial operand values, and the expected result after the broadcast. ```mlir // num_replicas: 4 // num_partitions: 1 // %operand@(0, 0): [[1, 2]] // %operand@(1, 0): [[3, 4]] // %operand@(2, 0): [[5, 6]] // %operand@(3, 0): [[7, 8]] %result = "stablehlo.collective_broadcast"(%operand) { replica_groups = dense<[[2, 1]]> : tensor<1x2xi64>, channel_handle = #stablehlo.channel_handle } : (tensor<1x2xi64>) -> tensor<1x2xi64> // %result@(0, 0): [[0, 0]] // %result@(1, 0): [[5, 6]] // %result@(2, 0): [[5, 6]] // %result@(3, 0): [[0, 0]] ``` -------------------------------- ### StableHLO TanOp Specification Source: https://github.com/openxla/stablehlo/blob/main/rfcs/20240312-tanop.md Details the semantics, inputs, outputs, constraints, and provides an example for the StableHLO TanOp. ```APIDOC ## stablehlo.tan ### Description Performs element-wise tangent operation on the `operand` tensor and produces a `result` tensor. ### Method N/A (This is an IR operation specification, not an API endpoint) ### Endpoint N/A ### Parameters #### Inputs - **operand** (tensor) - Required - A tensor of floating-point, complex type, or per-tensor quantized tensor. ### Outputs - **result** (tensor) - The resulting tensor after applying the tangent operation, with the same type as the operand. ### Constraints - **C1**: `baseline_type(operand) = baseline_type(result)`. ### Request Example ```mlir // %operand: [ // [0.0, 1.57079632], // [0, pi/2] // [3.14159265, 4.71238898] // [pi, 3pi/2] // ] %result = "stablehlo.tan"(%operand) : (tensor<2x2xf64>) -> tensor<2x2xf64> // %result: [ // [0.0, 1.63312e+16], // [0.0, 5.44375e+15] // ] ``` ### Response #### Success Response (N/A - This is an IR operation) - **result** (tensor) - The computed tangent values. #### Response Example (See Request Example for MLIR representation of the output) ``` -------------------------------- ### Create Dynamic Update Slice Operation Source: https://github.com/openxla/stablehlo/blob/main/docs/generated/StablehloBuilder.md Constructs a stablehlo.dynamic_update_slice operation. It requires the operand, update, and start indices. ```C++ MlirOp DynamicUpdateSlice(MlirOp &operand, MlirOp &update, ArrayRef start_indices); ``` -------------------------------- ### Serialize and Deserialize StableHLO Artifacts Source: https://context7.com/openxla/stablehlo/llms.txt Demonstrates how to convert MLIR files into portable bytecode format for deployment and how to restore them back to human-readable text. The serialization process supports target versioning to ensure compatibility. ```bash stablehlo-translate --serialize model.mlir --target=1.0.0 > model_v1.mlir.bc stablehlo-translate --deserialize model.mlir.bc > restored_model.mlir ``` -------------------------------- ### Create Dynamic Slice Operation Source: https://github.com/openxla/stablehlo/blob/main/docs/generated/StablehloBuilder.md Constructs a stablehlo.dynamic_slice operation. It requires the operand, start indices, and slice sizes. ```C++ MlirOp DynamicSlice(MlirOp &operand, ArrayRef start_indices, ::llvm::ArrayRef slice_sizes); ``` -------------------------------- ### Create StableHLO Programs with Python API Source: https://context7.com/openxla/stablehlo/llms.txt Demonstrates how to create StableHLO programs programmatically using MLIR Python bindings. It involves registering the dialect, creating constants, and performing operations like addition. ```Python from mlir import ir import mlir.dialects.stablehlo as stablehlo import mlir.dialects.func as func from mlir.ir import Context, Location, InsertionPoint, Module import numpy as np with Context() as ctx, Location.unknown(): stablehlo.register_dialect(ctx) module = Module.create() with InsertionPoint(module.body): @func.func() def main(): # Create constant tensors a_value = ir.DenseElementsAttr.get(np.zeros(shape=[3, 4], dtype=np.int64)) b_value = ir.DenseElementsAttr.get(np.ones(shape=[3, 4], dtype=np.int64)) a = stablehlo.constant(a_value) b = stablehlo.constant(b_value) # Perform element-wise addition result = stablehlo.add(a, b) return result assert main.func_op.verify() print(str(module)) ``` -------------------------------- ### Save and Load StableHLO Artifacts Source: https://github.com/openxla/stablehlo/blob/main/docs/tutorials/pytorch-export.ipynb Demonstrates how to save a StableHLO program to a directory and reload it for execution. These artifacts provide stable, version-compatible serialization. ```python from torch_xla.stablehlo import StableHLOGraphModule # Save to tmp stablehlo_program.save('/tmp/stablehlo_dir') # Reload and execute reloaded = StableHLOGraphModule.load('/tmp/stablehlo_dir') print(reloaded(sample_input[0])) ``` -------------------------------- ### Create Gather Operation Source: https://github.com/openxla/stablehlo/blob/main/docs/generated/StablehloBuilder.md Constructs a stablehlo.gather operation. It requires the operand, start indices, dimension numbers, and slice sizes. ```C++ MlirOp Gather(MlirOp &operand, MlirOp &start_indices, ::mlir::stablehlo::GatherDimensionNumbersAttr dimension_numbers, ::llvm::ArrayRef slice_sizes, /*optional*/bool indices_are_sorted = false); ``` -------------------------------- ### Create stablehlo::RealDynamicSliceOp Source: https://github.com/openxla/stablehlo/blob/main/docs/generated/StablehloBuilder.md Creates a new stablehlo.real_dynamic_slice operation. This operation extracts a slice from a tensor with dynamic start indices and limits. ```C++ MlirOp RealDynamicSlice(Type resultType, MlirOp &operand, MlirOp &start_indices, MlirOp &limit_indices, MlirOp &strides); ``` -------------------------------- ### Create Dynamic Gather Operation Source: https://github.com/openxla/stablehlo/blob/main/docs/generated/StablehloBuilder.md Constructs a stablehlo.dynamic_gather operation. It requires the operand, start indices, slice sizes, and dimension numbers. ```C++ MlirOp DynamicGather(MlirOp &operand, MlirOp &start_indices, MlirOp &slice_sizes, ::mlir::stablehlo::GatherDimensionNumbersAttr dimension_numbers, /*optional*/bool indices_are_sorted = false); ```