diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/01_overview.md b/content/learning-paths/mobile-graphics-and-gaming/luti/01_overview.md
new file mode 100644
index 0000000000..2dd3bf2101
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/01_overview.md
@@ -0,0 +1,101 @@
+---
+title: Understand why LUTI helps low-bit matrix multiplication
+description: Understand how packed low-bit weights reduce memory traffic and how LUTI expands their indices inside an Arm vector data path.
+weight: 2
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+## Overview
+We are at a point where Large Language Model(LLM) inference on the CPU is practical on mobile and edge devices. This is due in large part to the rise of low-bit AI models.
+Low-bit AI models store weights in compact packed format that must be efficiently expanded before matrix multiplication. You will see how 2-bit and 4-bit weights are stored, why conventional unpacking adds cycles, and how lookup-table instructions (LUTI) removes the unpacking step.
+
+By the end, you should be able to explain why packed low-bit indices are efficient for storage and memory traffic, how the indices are laid out in memory and how to use LUTI instructions to efficiently expand it.
+
+## Why use sub-byte weights?
+
+Large language model (LLM) inference on mobile and edge devices is often limited by memory capacity and bandwidth. During inference, model weights must be transferred from memory to the CPU, contributing to latency and energy use.
+
+Quantization reduces this traffic by storing weights in lower-precision formats. Weight-only quantization maps each 32-bit floating-point (`fp32`) weight to a compact logical code and stores shared metadata, such as a scale or zero point, for each block.
+
+{{% notice Note %}} The terms `4-bit` and `2-bit` specify the number of bits assigned to each logical code. These codes do not necessarily denote the numerical datatypes `int4` or `int2`. A `4-bit` or `2-bit` code might represent a signed integer, an unsigned integer, or an index into a codebook, depending on the quantization format. {{% /notice %}}
+
+Physical packing is the storage layout that places several low-bit codes into each byte. If you ignore the metadata, four `2-bit` codes or two `4-bit` codes can be stored in one byte.
+
+
+ 
+
+
+
+ Figure 1. Packing low-bit weight codes into a scalable vector register. An int8 value occupies 1 byte, whereas one byte can hold four 2-bit codes or two 4-bit codes. Consequently, a 128-bit (16 bytes) vector can contain 16 int8 values, 32 packed 4-bit codes, or 64 packed 2-bit codes.
+
+
+
+This approach trades reconstruction accuracy for lower memory use. Its value also depends on decoding the packed codes efficiently. LUTI addresses that work by expanding low-bit codes directly into arithmetic-ready vector values.
+
+## Understand the LUTI operation
+
+Matrix multiplication kernels do not usually operate on packed 2-bit or 4-bit codes.
+Before arithmetic, the codes must be decoded into values the computation can consume.
+
+Conceptually, the operation is:
+```c
+index = get_lut_index(packed_code);
+expanded_value = lookup_table[index];
+```
+
+Armv9-A LUTI instructions perform lookup-table operations that map low-bit indices to expanded values. LUTI2 and LUTI4 operate on 2-bit and 4-bit indices, respectively.
+
+- `LUTI2` uses each 2-bit index to select one of four lookup-table values.
+- `LUTI4` uses each 4-bit index to select one of sixteen lookup-table values.
+
+The lookup table defines the expanded value associated with each code according to the quantization scheme.
+
+### LUT for 2-bit codes
+
+For example, a 2-bit lookup table might contain:
+
+| Packed Code | LUT Index | Expanded Value |
+|---|---|---|
+| `0b00` | lut[0] | `-2` |
+| `0b01` | lut[1] | `-1` |
+| `0b10` | lut[2] | `0` |
+| `0b11` | lut[3] | `1` |
+
+
+### From packed 2-bit codes to 8-bit values
+The key benefit of LUTI is that the matrix multiplication kernel can load weights in their compact form. A source vector of packed weights therefore carries more values per memory load than a vector containing already expanded 8-bit, 16-bit, or 32-bit values.
+
+For 2-bit codes, LUTI uses the packed indices in a source vector to select lookup-table entries and writes the resulting values to destination vector registers. Figure 2 shows how 2-bit codes expand into 8-bit values.
+
+
+ 
+
+
+
+ Figure 2. LUTI maps packed 2-bit codes to lookup-table indices, selects the relevant 8-bit values, and writes them to a destination vector. Compared with storing expanded 8-bit values, the packed format allows each memory load to supply four times as many weights.
+
+
+
+## Identify LUTI responsibilities
+
+Keep these boundaries in mind when using LUTI:
+
+- The lookup table defines the meaning of each packed code
+- LUTI2 and LUTI4 expand indices; they don't calculate quantization metadata
+- Scaling, zero-point correction, bias, activation, clamping, and requantization remain separate operations
+- Expansion happens in the vector path, close to the arithmetic that consumes the values
+
+## What you've learned and what's next
+You've learned how LUTI uses packed low-bit codes as indices and expands them into values for subsequent arithmetic.
+
+Next, you'll set up the compiler and SME2 hardware needed to build and run the examples.
diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/02_luti_and_sme2.md b/content/learning-paths/mobile-graphics-and-gaming/luti/02_luti_and_sme2.md
new file mode 100644
index 0000000000..7ff19020b2
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/02_luti_and_sme2.md
@@ -0,0 +1,148 @@
+---
+title: Use LUTI with SME2
+description: Trace packed indices through the SME2 ZT0 table, streaming Z registers, and ZA matrix accumulators.
+weight: 3
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+## SME and SME2 architectural features
+
+SME extends the Armv9-A architecture and accelerates matrix-heavy computations, such as outer products and matrix multiplication.
+SME introduces __Streaming SVE mode__ and the __scalable ZA matrix-storage__ array. ZA accumulates matrix outer products and multi-vector dot products.
+
+SME2 builds on SME and adds __multi-vector instructions__ and the fixed __512-bit ZT0 lookup-table register__.
+
+For this learning path, the important points to note is that LUTI uses packed low-bit codes from Z source registers, reads the corresponding look-up table entries from `ZT0`, and writes expanded operands into Z destination registers.
+The expanded operands can then be consumed by SME2 matrix instructions such as `SDOT` or `SMOPA`, with results accumulated in `ZA` array.
+
+## ZT0 lookup-table register
+
+SME2 provides a fixed 512-bit architectural register named `ZT0`. It contains 64 bytes, arranged as sixteen 32-bit table entries.
+
+
+ 
+ 
+
+
+
+ Figure 3. ZT0 lookup-table organization and use by LUTI2 and LUTI4. ZT0 entries form a single linear table of entries 0–15. LUTI2 selects among entries 0–3 using 2-bit indices, while LUTI4 can select among all entries 0–15 using 4-bit indices. According to the destination element size, LUTI copies the low 8, 16, or 32 bits of the selected 32-bit table entry into the destination Z registers.
+
+
+
+LUTI instructions use packed low-bit indices from a source Z register (`Zn`) to select the corresponding `ZT0` register entry.
+
+The relevant `ZT0` entries are expanded to chosen output element destination width and written to output Z registers (`Zd`).
+
+LUTI2 and LUTI4 can populate one, two, or four destination Z registers. The number of destination registers and the expanded element width (`.B`, `.H`, or `.S`) determine how many packed source bits fill the destinations.
+
+The element suffix specifies the expanded destination width:
+- `.B` produces 8-bit elements.
+- `.H` produces 16-bit elements.
+- `.S` produces 32-bit elements.
+
+## Streaming mode
+
+ZT0-based SME2 LUTI instructions need both streaming mode and ZA enabled. `SMSTART` enables the required state, and `SMSTOP` disables it.
+
+Streaming mode changes the execution context in three ways:
+
+- Vector and predicate lengths use the *streaming vector length* (SVL), which can differ from the non-streaming vector length
+- Streaming instructions, including the multi-register LUTI2, SDOT, and SMOPA forms, become available
+- `PSTATE.ZA` controls access to both the ZA matrix-storage array and `ZT0`
+
+Efficient kernels enter streaming mode before repeated loops and exit afterwards. Streaming mode does not automatically stream matrix data from memory. The kernel still loads only the current computation tile.
+
+## Follow the SME2 LUTI data path
+
+A simplified SME2 LUTI sequence is:
+
+```text
+SMSTART
+* load ZT0 once
+* perform LUTI2
+SMSTOP
+```
+
+A detailed LUTI SME2 flow is:
+ 1. Enter SME streaming mode with `SMSTART` and load the LUT into the sixteen 32-bit `ZT0` register.
+ 2. Load LHS activations and packed RHS data for the current computation tile.
+ 3. Use LUTI2 or LUTI4 to expand the packed RHS indices from `ZT0` into Z registers.
+ 4. Feed the expanded RHS elements and LHS activations to SME2 instructions, such as `SDOT` or `SMOPA`.
+ 5. Accumulate partial matrix products in the SME2 `ZA` array.
+ 6. Convert, clamp, and store the completed output tile as required by the kernel.
+ 7. Exit SME streaming mode with `SMSTOP`. This disables the ZA and ZT0 state after the kernel completes.
+
+LUTI replaces the explicit unpack/decode portion of the data path. It does not replace the matrix multiply instruction that consumes the expanded values.
+
+### Example: kernel with LUTI2
+
+This example shows how LUTI2 expands packed 2-bit RHS weights. It simplifies register allocation, predication, addressing, and loop control to focus on the LUTI data flow.
+
+{{% notice Note %}} This example uses a 512-bit streaming vector length (SVL). The SVL is a CPU specific property. {{% /notice %}}
+
+__1. Define and pass the LUT__
+
+`ZT0` contains sixteen 32-bit entries. LUTI2 uses entries 0–3; entries 4–15 are unused and contain zero.
+```c
+static const int32_t lut_i8_i2[16] = {-2, -1, 0, 1,};
+```
+
+For a `.B` LUTI result, the low 8 bits of the selected 32-bit entry form the destination element.
+
+__2. Load the LUT into ZT0__
+
+Enter streaming mode, initialize ZA, and load ZT0. The lookup table does not change across the inner matrix loop, so load it once before the loop.
+
+```asm
+smstart // Enable Streaming SVE mode and ZA/ZT0 state
+zero {za} // Zero initialize accumulators for this output tile
+ldr zt0, [x_lut] // load LUT into fixed 512-bit ZT0 table
+```
+
+__3. Load LHS and packed RHS__
+
+Load the LHS activations and packed RHS 2-bit indices for the current computation tile, not the entire matrix.
+
+```asm
+ld1rqb {z0.b}, ... , [x_lhs] // load LHS
+ld1b {z16.b-z19.b}, ... , [x_packed_rhs] // load RHS packed 2-bit indices
+```
+ - The `z0.b` register receives the LHS activations.
+ - The `z16.b`–`z19.b` registers receive the packed RHS 2-bit indices.
+
+`ld1rqb` is the SVE load-and-replicate-quadword operation. For a 512-bit SVL, you can view `z0` as four 128-bit regions. `ld1rqb` replicates the 16-byte LHS block across those regions.
+
+__4. LUTI2 expands the packed indices__
+```asm
+luti2 { z24.b - z27.b }, zt0, z16[0] // unpack 2-bit indices
+luti2 { z4.b - z7.b }, zt0, z17[0]
+
+luti2 { z8.b - z11.b }, zt0, z18[0]
+luti2 { z12.b - z15.b }, zt0, z19[0]
+```
+Each LUTI2 instruction reads packed 2-bit indices from one source Z register (`z16` to `z19`) and expands them into four `.B` destination registers for a 512-bit SVL.
+
+__5. Feed the expanded vectors directly to SDOT__
+
+The expanded vectors can now feed SME2 matrix instructions such as `SDOT` or `SMOPA`, which accumulate the results in `ZA`.
+
+## What you've learned and what's next
+
+You've learned how LUTI operates within SME2: a micro-kernel loads the lookup table into `ZT0`, uses packed low-bit indices in Z registers to expand the RHS values, and feeds those expanded values into SME2 instructions.
+
+Next, you'll apply these concepts to a complete low-bit matrix multiplication kernel and examine how LUTI can replace explicit unpacking and decoding in the inner computation loop.
+
+## Further reading
+- [Arm SME2 Introduction](https://developer.arm.com/community/arm-community-blogs/b/architectures-and-processors-blog/posts/part4-arm-sme2-introduction)
+- [SME2 lookup table Armv9-A Documentation](https://developer.arm.com/documentation/109246/0101/SME-Overview/SME-and-SME2/SME2-lookup-table)
+- [Introduction to streaming and non-streaming mode](https://arm-software.github.io/acle/main/acle.html#controlling-the-use-of-streaming-mode)
diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/03_environment_setup.md b/content/learning-paths/mobile-graphics-and-gaming/luti/03_environment_setup.md
new file mode 100644
index 0000000000..18b1b897c9
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/03_environment_setup.md
@@ -0,0 +1,193 @@
+---
+title: Set up the SME2 environment
+description: Build and run the SME2 LUTI examples natively on macOS or cross-compile them for an Android device with SME2 support.
+weight: 4
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+## Choose an SME2 execution environment
+
+You can run the examples using one of the following route:
+
+- Build and run natively on an arm64 macOS® device with an M4 processor or later
+- Cross-compile on macOS or Linux and run on an Android™ phone with SME2 support
+
+See the [list of devices with native SME2 support](https://learn.arm.com/learning-paths/cross-platform/multiplying-matrices-with-sme2/1-get-started/#devices) before selecting a target device.
+
+The examples prefer LLVM Clang 22 or later because they use recent Arm C Language Extensions (ACLE) intrinsics and SME2 assembly syntax. The Makefile checks the appropriate compiler version before building.
+
+## Set up native macOS development
+
+Confirm that the system uses the `arm64` architecture:
+
+```bash
+uname -m
+```
+
+The expected output is:
+
+```output
+arm64
+```
+
+Check that SME2 is available to applications and inspect the maximum streaming vector length (SVL):
+
+```bash
+sysctl -n hw.optional.arm.FEAT_SME2
+sysctl -n hw.optional.arm.sme_max_svl_b
+```
+
+The first command must print `1`. The second command reports the maximum SVL in bytes. For example, `64` bytes corresponds to an SVL of 512 bits.
+
+Install Homebrew LLVM and confirm its version:
+
+```bash
+brew install llvm
+/opt/homebrew/opt/llvm/bin/clang --version
+```
+
+The Makefile selects Homebrew LLVM when it is installed. Otherwise, it falls
+back to Apple Clang from the active Xcode Command Line Tools.
+
+## Set up Android cross-compilation
+To run the LUTI examples codes on Android, install Android Native Development Kit (Android NDK):
+
+{{< tabpane code=true >}}
+ {{< tab header="macOS host" language="bash">}}
+wget https://dl.google.com/android/repository/android-ndk-r29-darwin.zip
+unzip android-ndk-r29-darwin.zip
+ {{< /tab >}}
+ {{< tab header="Linux host" language="bash">}}
+wget https://dl.google.com/android/repository/android-ndk-r29-linux.zip
+unzip android-ndk-r29-darwin.zip
+ {{< /tab >}}
+{{< /tabpane >}}
+
+For easier access and execution of Android NDK tools, add these to the PATH and set the NDK_PATH variable:
+{{< tabpane code=true >}}
+ {{< tab header="macOS host" language="bash">}}
+export NDK_PATH=$HOME/Library/Android/android-ndk-r29/
+export ANDROID_NDK_HOME=$NDK_PATH
+ {{< /tab >}}
+ {{< tab header="Linux host" language="bash">}}
+export NDK_PATH=$HOME/Android/android-ndk-r29/
+export ANDROID_NDK_HOME=$NDK_PATH
+ {{< /tab >}}
+{{< /tabpane >}}
+
+Install Android Debug Bridge (`adb`) if it isn't already available:
+
+{{< tabpane code=true >}}
+ {{< tab header="macOS host" language="bash">}}
+brew install android-platform-tools
+ {{< /tab >}}
+ {{< tab header="Linux host" language="bash">}}
+sudo apt update
+sudo apt install adb
+ {{< /tab >}}
+{{< /tabpane >}}
+
+Enable developer options and USB debugging on the Android phone, connect it to the host, and accept the debugging prompt on the phone. Verify the connection:
+
+```bash
+adb devices -l
+```
+
+Confirm that the device uses the `arm64-v8a` application binary interface (ABI):
+
+```bash
+adb shell getprop ro.product.cpu.abi
+```
+
+The expected output is:
+
+```output
+arm64-v8a
+```
+
+You can also inspect the CPU feature list:
+
+```bash
+adb shell "grep -m1 '^Features' /proc/cpuinfo"
+```
+
+The executable performs the definitive runtime check. If SME2 isn't available to Android applications, it prints `SKIP: No support for SME2 on this device; SME2 tests were not run.`
+
+## Download and explore the code examples
+
+__[!REVIEW - ADD THE PUBLIC CODE-EXAMPLE ARCHIVE URL BEFORE PUBLISHING]__
+
+Download and extract the published code-example archive, then change to its `code` directory.
+
+The directory contains these source and build files:
+
+```text
+code/
+├── Makefile
+├── example_1_luti_sme2.c
+├── luti_sme2_programming.c
+└── luti_sme2_programming_test.c
+```
+
+## Build and run on macOS
+
+From the `code` directory, clean previous outputs and build the native executable:
+
+```bash
+make clean
+make
+```
+
+Running `make macos` performs the same native build explicitly. Both commands
+use `-march=native+sme2+nosve2+nosve`. The `+nosve2+nosve` modifiers prevent
+the compiler from emitting non-streaming SVE or SVE2 instructions in the
+macOS executable.
+
+Run the introductory plain C and SME2 comparison:
+
+```bash
+./sme2_luti
+```
+
+Run the additional LUTI programming examples:
+
+```bash
+./sme2_luti --learning
+```
+
+## Build and run on Android
+
+From the same `code` directory, cross-compile for Android:
+
+```bash
+make clean
+make android
+```
+Connect your Android device to your development machine using a cable.
+Approve the connection on your phone and use adb to copies the executable to `/data/local/tmp/sme2_luti_android`:
+
+```bash
+adb push sme2_luti_android /data/local/tmp/sme2_luti_android
+```
+
+Start a new shell to access the device’s system from your development machine and runs the executable:
+
+```
+adb shell
+./data/local/tmp/sme2_luti_android
+```
+
+Run the additional LUTI programming examples with:
+
+```bash
+./data/local/tmp/sme2_luti_android --learning
+```
+
+## What you've accomplished and what's next
+
+You've prepared a compatible Clang compiler, verified an SME2-capable target,
+and built the same executable for native macOS or AArch64 Android.
+
+Next, you'll use `example_1_luti_sme2.c` to compare plain C shifts, masks, and scalar lookups with SME2 `LUTI2` expansion and matrix accumulation.
\ No newline at end of file
diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/04_lut_decoding_c_sme2.md b/content/learning-paths/mobile-graphics-and-gaming/luti/04_lut_decoding_c_sme2.md
new file mode 100644
index 0000000000..09eae32227
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/04_lut_decoding_c_sme2.md
@@ -0,0 +1,344 @@
+---
+title: Compare plain C decoding with SME2 LUTI2
+description: Compare low-bit matrix multiplication in plain C with an SME2 implementation that uses LUTI2 and SMOPA.
+weight: 5
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+## Compare plain C and SME2
+
+In this section, you compare two ways to expand packed 2-bit right-hand side
+(RHS) values for signed 8-bit matrix multiplication.
+
+The plain C reference extracts each index with a shift and mask. It then uses
+the index to select a signed 8-bit value from the lookup table. The SME2
+implementation uses `LUTI2` to expand one packed vector and `SMOPA` to
+accumulate four adjacent output panels in `ZA0`-`ZA3`.
+
+Both implementations use the same matrix dimensions, packed RHS bytes, and
+lookup table. The program compares their output matrices element by element.
+
+## Set up the example
+
+Open `code/example_1_luti_sme2.c`. The file contains both implementations and
+the validation code. The following snippets highlight the sections to inspect
+before you build the complete example.
+
+### Inspect the matrix shape
+
+The example derives its matrix dimensions from the streaming vector length (SVL).
+SVL is the number of bits in one streaming vector.
+
+Use these relationships when reasoning about the dimensions:
+
+- `svcntb()` returns the number of bytes in one streaming vector.
+- `svcntw()` returns the number of 32-bit words in one streaming vector.
+- One 32-bit word contains four bytes, so `svcntb() = 4 * svcntw()`
+
+The example uses those values to define the matrix shape:
+
+```c
+const size_t m = streaming_vector_words(); // svcntw()
+const size_t n = 4 * m; // svcntb()
+```
+
+Both implementations calculate:
+
+```text
+DST[M, N] = LHS[M, K] x RHS[K, N]
+```
+
+`M` is the number of 32-bit words for one streaming vector: `M = svcntw()`.
+`N` is the number of bytes in one streaming vector: `N = svcntb()`, which is equivalent to `N = 4 * M`.
+
+The example fixes `K = 4` for two related reasons. First, LUTI2 uses 2-bit lookup indices, so one packed RHS byte contains four 2-bit groups.
+Those four groups provide the four RHS values along the K dimension for one output column.
+Second, with `M = svcntw()`, setting `K = 4` makes the `M * K` signed 8-bit LHS block occupy exactly one streaming vector.
+
+For an SVL of 512 bits:
+
+- M = svcntw() = 512 / 32 = 16
+- N = svcntb() = 512 / 8 = 64
+- K = 4
+
+| Block | Calculation | Size |
+|---|---|---|
+| LHS | `M * K = 16 * 4` | 64 signed 8-bit values |
+| Logical RHS | `K * N = 4 * 64` | 256 2-bit codes |
+| Packed RHS | `256 * 2 bits = 512 bits` | 64 bytes |
+| DST | `M * N = 16 * 64` | 1024 signed 32-bit values |
+
+The complete LHS block and packed RHS block each fit in one streaming Z register.
+
+### Low-bit packed format
+
+For each output column, the packed RHS stores the four `K` dimension RHS values in one byte: rhs_packed[col].
+LUTI2 uses 2-bit indices, so the byte is split into four 2-bit groups.
+Each group selects the lookup-table value for one RHS element, RHS[k, col].
+
+```text
+rhs_packed[col]
+bits [7:6] | [5:4] | [3:2] | [1:0]
+RHS: [k3,col] | [k2,col] | [k1,col] | [0k,col]
+```
+
+The example reads the 2-bit groups from the least-significant bits first. For example:
+
+```text
+packed byte: 0xE4 = 11_10_01_00
+ | | | |
+ k3 k2 k1 k0
+
+lookup indices read in k order: 0, 1, 2, 3
+```
+
+The example uses these constants in both implementations:
+
+```c
+enum {
+ K = 4,
+ LUT_INDICES_PER_BYTE = 4,
+};
+```
+
+### Inspect the lookup table
+
+The example uses this mapping:
+
+```text
+2-bit code index signed 8-bit value raw byte
+ 00 0 -3 0xFD
+ 01 1 -1 0xFF
+ 10 2 1 0x01
+ 11 3 3 0x03
+```
+
+```c
+static const int8_t lut_i8_i2[16] = {
+ -3, -1, 1, 3,
+ 0, 0, 0, 0,
+ 0, 0, 0, 0,
+ 0, 0, 0, 0,
+};
+```
+
+The 2-bit indices select only entries 0-3. Entries 4-15 remain zero.
+
+## Inspect the plain C reference
+
+The plain C reference calculates one output element at a time. It reads one
+packed RHS byte for each column. For each `k` position, the inner loop shifts
+the corresponding 2-bit field into bits `[1:0]`, applies the `0x03` mask,
+and uses the result to index `lut_i8_i2`.
+
+```c
+static void plain_c_matmul(
+ const int8_t* lhs,
+ const uint8_t* rhs_packed,
+ int32_t* dst,
+ size_t m,
+ size_t n) {
+
+ // Iterate over the output rows and columns.
+ for (size_t row = 0; row < m; ++row) {
+ for (size_t col = 0; col < n; ++col) {
+
+ // Read the packed byte for this output column.
+ const uint8_t packed_byte = rhs_packed[col];
+
+ // Initialize the accumulator for this output element.
+ int32_t acc_sum = 0;
+
+ for (size_t k_idx = 0; k_idx < K; ++k_idx) {
+ // Extract the 2-bit lookup table index for this k position.
+ const unsigned bit_shift = 2U * (unsigned)k_idx;
+ const uint8_t lut_idx = (uint8_t)((packed_byte >> bit_shift) & 0x3U);
+ const int8_t expanded_byte = lut_i8_i2[lut_idx];
+
+ // Widen the LHS and RHS values before multiplication.
+ const int32_t lhs_value = (int32_t)lhs[row * K + k_idx];
+ const int32_t expanded_rhs_value = (int32_t)expanded_byte;
+
+ acc_sum += lhs_value * expanded_rhs_value;
+ }
+
+ // Store the accumulator in the corresponding output element.
+ dst[row * n + col] = acc_sum;
+ }
+ }
+}
+```
+
+For each output element, the loop implements:
+
+```text
+DST[row, col] = sum(k=0..3) LHS[row, k] * RHS[k, col]
+```
+
+The program uses this implementation as the correctness reference. The
+innermost loop shifts, masks, looks up, and multiplies each value.
+
+## Inspect the SME2 LUTI2 implementation
+
+### Step 1: Store the same lookup values in `ZT0`
+
+The SME2 path uses the same logical lookup values. `ZT0` has a fixed physical
+layout of sixteen 32-bit entries (64 bytes).
+
+```c
+static const int32_t zt0_table[16] __attribute__((aligned(64))) = {
+ -3, -1, 1, 3,
+ 0, 0, 0, 0,
+ 0, 0, 0, 0,
+ 0, 0, 0, 0,
+};
+```
+
+`LUTI2 .B` selects entries 0-3 and copies the low byte of each selected
+32-bit entry. The low bytes of `ZT0` entries 0-3 therefore match the bit
+patterns in the plain C table for the four signed 8-bit values.
+
+### Step 2: Declare a locally streaming SME function
+
+```c
+__arm_new("za", "zt0") __arm_locally_streaming
+```
+
+These attributes tell the compiler to run the function body in streaming
+mode and provide new `ZA` and `ZT0` state for the function.
+
+### Step 3: Follow the SME2 compute path
+
+Review the SME2 implementation below. Inline assembly loads `ZT0`, expands
+the packed RHS, and accumulates four output panels. ACLE intrinsics read `ZA`
+and store the output matrix.
+
+```c
+__arm_new("za", "zt0") __arm_locally_streaming
+static void luti2_sme2_asm_matmul(const int8_t *lhs,
+ const uint8_t *rhs_packed,
+ int32_t *dst, size_t m, size_t n) {
+ __asm__ volatile(
+ "ptrue p0.b\n"
+ "ldr zt0, [%[table]]\n" // Load the lookup table into ZT0
+ "zero {za}\n" // Clear all ZA accumulator state
+
+ // Load one streaming vector from each input.
+ "ld1b {z0.b}, p0/z, [%[lhs]]\n"
+ "ld1b {z1.b}, p0/z, [%[rhs]]\n"
+
+ // Expand one packed source vector into four signed 8-bit vectors.
+ "luti2 {z4.b-z7.b}, zt0, z1[0]\n"
+
+ // Accumulate four adjacent M-by-M output panels in ZA0-ZA3.
+ "smopa za0.s, p0/m, p0/m, z0.b, z4.b\n"
+ "smopa za1.s, p0/m, p0/m, z0.b, z5.b\n"
+ "smopa za2.s, p0/m, p0/m, z0.b, z6.b\n"
+ "smopa za3.s, p0/m, p0/m, z0.b, z7.b\n"
+ :
+ : [lhs] "r"(lhs),
+ [rhs] "r"(rhs_packed),
+ [table] "r"(zt0_table)
+ : "p0", "z0", "z1", "z4", "z5", "z6", "z7",
+ "za", "zt0", "memory");
+
+ // Read four vectors from ZA and store them.
+ for (uint32_t row = 0; row < m; ++row) {
+ svint8x4_t read_tiles = svread_hor_za8_s8_vg4(0, 4 * row);
+ svint32x4_t output_tiles = svreinterpret_s32_s8_x4(read_tiles);
+ svst1_s32_x4(
+ svptrue_c32(),
+ dst + (size_t)row * n,
+ output_tiles);
+ }
+}
+```
+
+The inline assembly loads the lookup table into `ZT0` and loads one vector
+from each input. `LUTI2` expands the packed RHS into four vectors. Four
+`SMOPA` instructions accumulate those vectors into `ZA0`-`ZA3`.
+
+The output loop reads one horizontal row across `ZA0`-`ZA3`. It
+reinterprets the returned bytes as four vectors of `int32_t` accumulators and
+stores them as one contiguous output row.
+
+`LUTI2` does not perform the multiplication. The packed RHS stays compact
+until the matrix kernel needs it. The expanded values then pass directly from
+Z registers to SME2 matrix instructions.
+
+To see the same instruction pattern in production code, inspect the
+[`qai8dxp_qsu2csp` Arm® KleidiAI™ micro-kernel source](https://gitlab.arm.com/kleidi/kleidiai/-/blob/v1.30.0/kai/ukernels/matmul/matmul_clamp_f32_qai8dxp_qsu2cxp/kai_matmul_clamp_f32_qai8dxp1vlx4_qsu2cxp4vlx4_1vlx4vl_sme2_mopa_asm.S).
+
+## Build and validate the example
+
+From the `code` directory, compile `example_1_luti_sme2.c` with SME2 enabled:
+
+```bash
+/opt/homebrew/opt/llvm/bin/clang \
+ -O2 -Wall -Wextra -Werror \
+ -march=native+sme2 \
+ -isysroot "$(xcrun --show-sdk-path)" \
+ example_1_luti_sme2.c \
+ -o example_1_luti_sme2
+```
+
+Run the example:
+
+```bash
+./example_1_luti_sme2
+```
+
+The program first prints the logical lookup table:
+
+```text
+bits idx signed raw byte
+ 00 0 -3 0xFD
+ 01 1 -1 0xFF
+ 10 2 1 0x01
+ 11 3 3 0x03
+```
+
+It then prints sample packed bytes, their indices in `k0` to `k3` order, and
+the signed and hexadecimal decoded values.
+
+After showing a preview of the result matrix, the program reports:
+
+```text
+PASS: LUTI2 SME2 matches plain C matmul.
+```
+
+The program produces this message only after comparing every SME2 result
+against the corresponding plain C result.
+
+## Inspect the generated SME2 instructions
+
+Disassemble the executable and confirm that the SME2 function contains
+`LUTI2` followed by four `SMOPA` instructions.
+
+```bash
+/opt/homebrew/opt/llvm/bin/llvm-objdump -d example_1_luti_sme2 | \
+ grep -E "luti2|smopa"
+```
+
+## Check your understanding
+
+Before continuing, make sure you can explain:
+
+- Why one packed byte represents four values along the `K` dimension.
+- Why the plain C shift counts are 0, 2, 4, and 6.
+- Why the mask is `0x03`.
+- How one `LUTI2` produces four decoded Z registers.
+- Why four `SMOPA` instructions produce four adjacent output panels in
+ `ZA0`-`ZA3`.
+- How the element-by-element comparison validates the SME2 result.
+
+## What you can do now and what's next
+
+You can now decode the same packed 2-bit RHS data in plain C or expand it with
+SME2 `LUTI2`. You can also pass the expanded Z-register values directly to
+`SMOPA` and accumulate signed 32-bit results in `ZA`.
+
+Next, develop intuition for programming LUTI instructions through practical
+SME2 examples.
\ No newline at end of file
diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/05_programming_luti_sme2.md b/content/learning-paths/mobile-graphics-and-gaming/luti/05_programming_luti_sme2.md
new file mode 100644
index 0000000000..41a7f072d1
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/05_programming_luti_sme2.md
@@ -0,0 +1,307 @@
+---
+title: Program SME2 LUTI examples
+description: Select LUTI destination groups and source segments for FP16 FMOPA and two-stage SDOT kernels.
+weight: 6
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+The examples in `luti_sme2_programming.c` show a recipe-based approach to programming
+with LUTI instructions.
+
+The examples cover the following combinations and have their base
+in KleidiAI's matrix multiplication micro-kernels:
+
+| Example | Decode | Arithmetic | Main concept |
+|---|---|---|---|
+| One | LUTI4 to `float16` | GEMM using FMOPA | Use LUTI and source segments |
+| Two | LUTI4, then LUTI2 to `int8` | GEMV using SDOT | Use multiple LUTs and source segments |
+
+## Run the learning tests
+
+From the `code` directory, run:
+
+```bash
+./sme2_luti --learning
+```
+
+In the output, `SVL` is the streaming vector length in bits. `VL_b` is the
+number of byte elements (`SVL / 8`), `VL_h` is the number of half-word
+elements (`SVL / 16`), and `VL_s` is the number of 32-bit word elements
+(`SVL / 32`).
+
+With a 512-bit SVL, the expected output is:
+```output
+FP16 LUTI4 + FMOPA test (M = VL_s, K = 2, N = 4 * VL_s)
+PASS
+LUTI4 -> LUTI2 -> SDOT test (M = 1, K = N = VL_b)
+PASS
+```
+
+Each `PASS` confirms that the kernel matches the reference result.
+
+## The four-step LUTI recipe
+
+For every LUTI call, answer these questions.
+
+| Step | Decision | Result |
+|---|---|---|
+| 1 | What are the packed-index and table-element widths? | Choose the LUTI form and ZT0 register table-entry width. |
+| 2 | What element type does the destination Z register require? | Select `.B`, `.H`, or `.S`. |
+| 3 | How many destination Z registers do you need the lookup to fill? | Choose x1, x2, or x4 to match the target operation. |
+| 4 | How much of the source Z register fills the destination register group? | Source-register segment |
+
+A *source segment* is the portion of one packed source Z register that fills the chosen destination group.
+Its selector is relative to the destination-group size.
+The examples use several source-segment cases to help you develop intuition for selecting the correct segment.
+
+
+## Example 1: LUTI4 for FP16 GEMM using FMOPA
+
+This example is a focused extraction from KleidiAI's
+[FP16 LUTI4 FMOPA micro-kernel](https://gitlab.arm.com/kleidi/kleidiai/-/blob/v1.30.0/kai/ukernels/matmul/matmul_clamp_f32_f16p_qsi4c32p/kai_matmul_clamp_f32_f16p1vlx2_qsi4c32p4vlx2_1vlx4vl_sme2_mopa.c).
+The example uses 4-bit codes as indices that map to `float16` values.
+This example shows how the source-index segment is interpreted relative to the destination-group size.
+This section uses generic SVL terminology to explain the example.
+
+```c
+__arm_new("za", "zt0") __arm_locally_streaming void arm_lp_gemm_luti4(
+ const float16_t* lhs, const uint8_t* rhs_indices, float32_t* out, const uint32_t* zt0_lut) {
+ uint32_t m = svcntw(); // Number of FP32 rows/columns in one ZA tile.
+
+ /* LUTI4 decode
+ * +------+---------------------------+----------------------------------------+
+ * | Step | Decision | Choice |
+ * +------+---------------------------+----------------------------------------+
+ * | 1 | Index and table width | 4-bit index |
+ * | | | 16-bit ZT0 LUT register element |
+ * | 2 | Destination element type | .H, because FMOPA consumes FP16 |
+ * | 3 | Destination group | x1 for rhs_0/rhs_1, then x2 for rhs_23 |
+ * +------+---------------------------+----------------------------------------+
+ */
+
+ // There are more than one ways to reason about the number of input segment. Here, we go
+ // about it from the source register as the reference point.
+
+ // Load the LUT
+ svldr_zt(0, zt0_lut);
+ svzero_za();
+
+ // Load the LHS
+ svbool_t pg = svptrue_b16();
+ svfloat16_t lhs_ip = svld1_f16(pg, lhs);
+
+ // Load one source register of indices
+ svuint8_t s4_indices = svld1_u8(svptrue_b8(), rhs_indices);
+
+ // Case 1: a single-register destination group uses one of four segments.
+
+ // Step 4 : Number of input segments for x1
+ // --------------------------------------------
+ // one byte of index produces two half words after the look up.
+ // VL_b bytes of indices from a source register produces 2 * VL_h half-word elements or
+ // 4 * VL_b bytes.
+ // In other words, to fill VL_b bytes of destination, VL_b / 4 bytes
+ // is needed => 4 input segments with values 0, 1, 2 and 3
+
+ //
+ // source z register: Packed indices
+ // +-------------+-------------+-------------+-------------+
+ // | segment [3] | segment [2] | segment [1] | segment [0] |
+ // +-------------+-------------+-------------+-------------+
+ // |
+ // | LUTI4 .H, segment [0]
+ // v
+ // destination z register with F16 elements
+ // +----------------------------------------------------+
+ // | VL_h elements |
+ // +----------------------------------------------------+
+
+ // Unpredicated LUTI read.
+ svfloat16_t rhs_0 = svluti4_lane_zt_f16(0, s4_indices, /* segment */ 0);
+ svfloat16_t rhs_1 = svluti4_lane_zt_f16(0, s4_indices, /* segment */ 1);
+
+ // Case 2: a two-register destination group using one of two segments.
+
+ // Step 4 : Number of input segments for x2
+ // --------------------------------------------
+ // VL_b bytes of indices from a source register produces 2 * VL_h half-word elements or
+ // 4 * VL_b bytes.
+ // In other words, to fill 2 * VL_b bytes of destination, VL_b / 2 bytes
+ // is needed => 2 segments with values 0 and 1
+ //
+ // source z register: packed indices
+ // +---------------------------+---------------------------+
+ // | segment [1] | segment [0] |
+ // +---------------------------+---------------------------+
+ // |
+ // | LUTI4 .H, segment [1]
+ // v
+ // two-register destination z register group
+ // +---------------------+---------------------+
+ // | destination 1 | destination 0 |
+ // | VL_h elements | VL_h elements |
+ // +---------------------+---------------------+
+ svfloat16x2_t rhs_23 = svluti4_lane_zt_f16_x2(0, s4_indices, /* Segment */ 1);
+
+ svmopa_za32_f16_m(0, pg, pg, lhs_ip, rhs_0);
+ svmopa_za32_f16_m(1, pg, pg, lhs_ip, rhs_1);
+ svmopa_za32_f16_m(2, pg, pg, lhs_ip, svget2_f16(rhs_23, 0));
+ svmopa_za32_f16_m(3, pg, pg, lhs_ip, svget2_f16(rhs_23, 1));
+
+ // Extract out the data from ZA tiles and store
+ for (uint32_t i_m = 0; i_m < m; i_m += 1) {
+ svfloat32x4_t out_row = svread_hor_za32_f32_vg4(0, 4 * i_m);
+ svst1_f32_x4(svptrue_c32(), out + ((size_t)4 * i_m * m), out_row);
+ }
+}
+```
+## Two-stage LUTI4 and LUTI2 for GEMV using SDOT
+
+The SDOT micro-kernel and block size are similar to the KleidiAI's
+[SDOT micro-kernel](https://gitlab.arm.com/kleidi/kleidiai/-/blob/v1.30.0/kai/ukernels/matmul/matmul_clamp_f32_qai8dxp_qsi4cxp/kai_matmul_clamp_f32_qai8dxp1x4_qsi4cxp4vlx4_1x4vl_sme2_sdot.c),
+with two LUTs added to implement a two-stage decode of vector-quantized
+weights.
+
+### Define the terms
+
+In this context, vector quantization uses a 4-bit code as an index to select an 8-bit codeword that represents a four-weight pattern. The sixteen `ZT0` entries represent sixteen unique patterns.
+
+The two-stage decode uses LUTI4 to select an 8-bit codeword, then LUTI2 to map the codeword to four 8-bit numerical values.
+
+This encoding reduces the bytes required to store weights.
+```text
+ packed 4-bit pattern ID
+ +-----------------------------+
+ | index[3:0] |
+ +-----------------------------+
+ |
+ | LUTI4: select one 8-bit codeword from ZT0
+ v
+ packed 8-bit data with 2-bit symbols
+ +--------+--------+--------+--------+
+ | sym[3] | sym[2] | sym[1] | sym[0] |
+ | 2 bits| 2 bits| 2 bits| 2 bits|
+ +--------+--------+--------+--------+
+ | | | |
+ +--------+--------+--------+
+ | LUTI2: expand each 2-bit symbol from ZT0
+ v
+ +--------+--------+--------+--------+
+ | weight3| weight2| weight1| weight0|
+ | int8 | int8 | int8 | int8 |
+ +--------+--------+--------+--------+
+
+```
+
+```c
+__arm_new("za", "zt0") __arm_locally_streaming void arm_lp_gemv_luti2_luti4(
+ const int8_t* lhs, const uint8_t* rhs_indices, int32_t* out,
+ const uint32_t* zt0_luti4, const uint32_t* zt0_luti2) {
+ const size_t vl_b = svcntb();
+ const size_t lhs_blocks = vl_b / 16;
+ const svbool_t pg8 = svptrue_b8();
+ const svcount_t pn8 = svptrue_c8();
+
+ assert(lhs != NULL);
+ assert(rhs_indices != NULL);
+ assert(out != NULL);
+ assert(zt0_luti4 != NULL);
+ assert(zt0_luti2 != NULL);
+ assert(vl_b % 16 == 0);
+
+ svzero_za();
+ for (size_t i_k = 0; i_k < lhs_blocks; i_k++) {
+ // For a SVL of 512 bits: SVL_b = 64 bytes and SVL_s = 16 words.
+
+ // Replicated read of one 16-byte LHS block for SDOT lanes.
+ svint8_t lhs_ip = svld1rq_s8(pg8, lhs);
+ lhs += 16;
+
+ /* LUTI4 decode (first stage)
+ * +------+---------------------------+----------------------------------------+
+ * | Step | Decision | Choice |
+ * +------+---------------------------+----------------------------------------+
+ * | 1 | Index and table width | 4-bit index |
+ * | | | 8-bit ZT0 LUT register element |
+ * | 2 | Destination element type | NA. Second decode stage addresses this |
+ * | 3 | Destination group | x2 for patterns_01 and patterns_23 |
+ * +------+---------------------------+----------------------------------------+
+ */
+
+ // Two source registers provide the four packed 2-bit vectors needed for the second stage.
+ svuint8x2_t rhs_packed = svld1_u8_x2(pn8, rhs_indices);
+ rhs_indices += 2 * vl_b;
+
+ // Load LUT 1
+ svldr_zt(0, zt0_luti4);
+
+ // Step 4 : Number of input segments for x2
+ // --------------------------------------------
+ // One packed byte contains two 4-bit indices and produces two bytes
+ // after the lookup. SVL_b bytes of indices from one source register
+ // therefore produce 2 * SVL_b bytes, filling the x2 destination
+ // group. This uses one input segment with value 0.
+
+ svuint8x2_t patterns_01 = svluti4_lane_zt_u8_x2(0, svget2_u8(rhs_packed, 0), 0);
+ svuint8x2_t patterns_23 = svluti4_lane_zt_u8_x2(0, svget2_u8(rhs_packed, 1), 0);
+
+ /* LUTI2 decode (second stage)
+ * +------+---------------------------+----------------------------------------+
+ * | Step | Decision | Choice |
+ * +------+---------------------------+----------------------------------------+
+ * | 1 | Index and table width | 2-bit index |
+ * | | | 8-bit ZT0 LUT register element |
+ * | 2 | Destination element type | .B, SDOT consumes int8 |
+ * | 3 | Destination group | x4 for rhs_unpacked |
+ * +------+---------------------------+----------------------------------------+
+ */
+
+ // Load LUT 2
+ svldr_zt(0, zt0_luti2);
+
+ // Step 4 : Number of input segments for x4
+ // --------------------------------------------
+ // One packed byte contains four 2-bit indices and produces four
+ // bytes after the lookup. SVL_b bytes of indices from one source
+ // register therefore produce 4 * SVL_b bytes, filling the x4
+ // destination group. This uses one input segment with value 0.
+
+ svint8x4_t rhs_unpacked = svreinterpret_s8_u8_x4(svluti2_lane_zt_u8_x4(0, svget2_u8(patterns_01, 0), 0));
+
+ // Process k index of 0 to 3.
+ // Process 1VL_s of N
+ svdot_lane_za32_s8_vg1x4(0, rhs_unpacked, lhs_ip, /* Lane */ 0);
+
+ rhs_unpacked = svreinterpret_s8_u8_x4(svluti2_lane_zt_u8_x4(0, svget2_u8(patterns_01, 1), 0));
+
+ // Process k index of 4 to 7
+ // process next 1VL_s of N
+ svdot_lane_za32_s8_vg1x4(0, rhs_unpacked, lhs_ip, /* Lane */ 1);
+
+ rhs_unpacked = svreinterpret_s8_u8_x4(svluti2_lane_zt_u8_x4(0, svget2_u8(patterns_23, 0), 0));
+
+ // Process k index of 8 to 11
+ // Process next 1VL_s of N
+ svdot_lane_za32_s8_vg1x4(0, rhs_unpacked, lhs_ip, /* Lane */ 2);
+
+ rhs_unpacked = svreinterpret_s8_u8_x4(svluti2_lane_zt_u8_x4(0, svget2_u8(patterns_23, 1), 0));
+
+ // Process k index of 12 to 15
+ // Process next 1VL_s or N
+ svdot_lane_za32_s8_vg1x4(0, rhs_unpacked, lhs_ip, /* Lane */ 3);
+
+ // Total processed: 16 K-values, 4VL_s N values
+ }
+
+ svint32x4_t result = svread_za32_s32_vg1x4(0);
+ svst1_s32_x4(svptrue_c32(), out, result);
+}
+```
+
+## What you've learned in this section
+
+You've seen how destination-group size determines the meaning of a source
+segment, and how the same four-step method applies once or repeatedly in a
+multi-stage decode.
\ No newline at end of file
diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/06_additional_luti_features.md b/content/learning-paths/mobile-graphics-and-gaming/luti/06_additional_luti_features.md
new file mode 100644
index 0000000000..41ce92e024
--- /dev/null
+++ b/content/learning-paths/mobile-graphics-and-gaming/luti/06_additional_luti_features.md
@@ -0,0 +1,114 @@
+---
+title: Identify additional LUTI features
+description: Compare ZT0 and Z-register lookup tables, SME2.1 destination forms, and their compile-time feature macros.
+weight: 7
+
+### FIXED, DO NOT MODIFY
+layout: learningpathall
+---
+
+The earlier examples use the original SME2 lookup path: a table in `ZT0`,
+packed indices in Z registers, and one or more Z-register results. Other
+architectural features use a different table source or add specialized forms.
+
+## Compare the feature paths
+
+| Feature path | Table source | Execution state | Distinguishing capability |
+|---|---|---|---|
+| `FEAT_SME2` | Fixed 512-bit `ZT0` | Streaming mode with ZA enabled | `LUTI2` and `LUTI4` can produce one, two, or four Z-register results, subject to the element-width encoding |
+| `FEAT_SME2p1` | Fixed 512-bit `ZT0` | Streaming mode with ZA enabled | Extends the SME2 forms with strided destination pairs and quads |
+| `FEAT_LUT` with `FEAT_SVE2` or `FEAT_SME2` | One or two scalable Z registers | Non-streaming SVE or Streaming SVE, respectively | Add Z-register table forms of `LUTI2` and `LUTI4` which produce one Z-register result without using `ZT0` |
+
+## Use Z-register tables with FEAT_LUT
+
+`FEAT_LUT` expand the functionality of `FEAT_SME2` allowing to use scalable Z registers as the lookup-table source. This provide an alternative to using the `ZT0` register.
+
+```text
+table Z register(s) + packed-index Z register
+ |
+ LUTI2/LUTI4
+ |
+ one result Z register
+```
+
+Use this form for vector kernels that do not need `ZT0` register and where multiple LUT are
+required by the algorithm.
+
+For example, the two-stage vector decode loop in section 5 reloads `ZT0` for its LUTI4 and
+LUTI2 tables. With `FEAT_LUT`, load separate Z-register tables once before
+the loop with the appropriate predicate. For the byte forms shown here, the
+LUTI4 load needs 16 active lanes and the LUTI2 load needs 4.
+
+```c
+const svuint8_t luti4_table_z = svld1_u8(pg_1, luti4_table_storage);
+const svuint8_t luti2_table_z = svld1_u8(pg_2, luti2_table_storage);
+
+for (size_t i_k = 0; i_k < lhs_blocks; ++i_k) {
+ svuint8_t packed_indices = svld1_u8(pg8, rhs_indices);
+ rhs_indices += vl_b;
+```
+ svldr_zt(0, zt0_luti4); svldr_zt(0, zt0_luti2);