diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/05_build_validate_luti2.md b/content/learning-paths/mobile-graphics-and-gaming/luti/05_build_validate_luti2.md new file mode 100644 index 0000000000..6395d7d52c --- /dev/null +++ b/content/learning-paths/mobile-graphics-and-gaming/luti/05_build_validate_luti2.md @@ -0,0 +1,118 @@ +--- +title: Build and validate the LUTI2 decoding example +description: Build and run the first LUTI2 example, compare its results with plain C, and inspect the generated SME2 instructions. +weight: 6 + +### FIXED, DO NOT MODIFY +layout: learningpathall +--- + +## Build and validate the example + +You've inspected how `code/example_1_luti_decoding.c` implements the same calculation in +plain C and with SME2 LUTI2. Now build the complete example and verify that both +implementations produce the same output. + +Run the following commands from the `code` directory. + +### Build and run on macOS + +Build and run the executable on an SME2-supported device: + +```bash +make example_1_luti_decoding +./example_1_luti_decoding +``` + +### Cross-compile and run on Android + +On macOS or Linux, use LLVM 22 and the NDK r29 installation selected by `ANDROID_NDK_HOME`. The build host doesn't need SME2 support. + +Build the standalone Android executable: + +```bash +make example_1_luti_decoding_android +``` + +With an Android device connected through `adb`, push the executable to the device: + +```bash +adb push example_1_luti_decoding_android /data/local/tmp/example_1_luti_decoding_android +``` + +Open an `adb` shell, make the file executable, and run it: + +```bash +adb shell +cd /data/local/tmp +chmod 755 example_1_luti_decoding_android +./example_1_luti_decoding_android +``` + +After running the file, enter `exit` to return to the build host's shell. + +### Check the result + +On an SME2-capable device, the program prints the matrix shape, lookup table, decoded right-hand side (RHS) samples, and a matrix preview. + +For a 512-bit SVL, the output begins with: + +```output +SVL = 512 bits; matrix shape M=16, K=4, N=64 +2-bit LUT mapping: +bits idx signed raw byte + 00 0 -3 0xFD + 01 1 -1 0xFF + 10 2 1 0x01 + 11 3 3 0x03 +``` + +The RHS decoding preview and C matrix preview follow. The final validation line is similar to: + +```output +PASS: LUTI2 SME2 matches plain C matmul. +``` + +If SME2 is unavailable, the runner exits without performing the calculations: + +```output +SKIP: No support for SME2 on this device. +``` + +A `SKIP` result doesn't validate the calculation. You can still inspect the generated instructions on the build host. + +## Inspect the generated SME2 instructions + +For the native macOS executable, run: + +```bash +make disassemble-example-1 +``` + +For the Android executable, run the following command on the macOS or Linux build host: + +```bash +make disassemble-example-1-android +``` + +The Makefile selects the host's LLVM disassembler and displays only LUTI2 and `SMOPA` instructions. Disassembly doesn't require SME2 hardware. + +The output is similar to: + +```output +100000cec: c08c8024 luti2 { z4.b - z7.b }, zt0, z1[0] +100000cf0: a0840000 smopa za0.s, p0/m, p0/m, z0.b, z4.b +100000cf4: a0850001 smopa za1.s, p0/m, p0/m, z0.b, z5.b +100000cf8: a0860002 smopa za2.s, p0/m, p0/m, z0.b, z6.b +100000cfc: a0870003 smopa za3.s, p0/m, p0/m, z0.b, z7.b +``` + +Addresses vary by build. Confirm that one LUTI2 is followed by four `SMOPA` instructions targeting `ZA0` through `ZA3`. + +## What you've accomplished and what's next + +You've built the first example, compared the plain C and SME2 results, and +confirmed that the generated code contains LUTI2 followed by four `SMOPA` +instructions. + +Next, you'll learn to program LUTIs through practical SME2 examples. diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/06_programming_luti_sme2.md b/content/learning-paths/mobile-graphics-and-gaming/luti/06_programming_luti_sme2.md new file mode 100644 index 0000000000..d10799dd53 --- /dev/null +++ b/content/learning-paths/mobile-graphics-and-gaming/luti/06_programming_luti_sme2.md @@ -0,0 +1,284 @@ +--- +title: Apply LUTI to SME2 matrix kernels +description: Select LUTI destination groups and source segments for FP16 FMOPA and two-stage SDOT kernels. +weight: 7 + +### FIXED, DO NOT MODIFY +layout: learningpathall +--- + +## The four-step LUTI recipe + +The examples in `example_2_luti_programming.c` show a recipe-based approach to programming with lookup-table instructions (LUTIs). + +The examples cover the following combinations and are based on KleidiAI matrix multiplication micro-kernels: + +| Example | Decode | Arithmetic | Main concept | +|---|---|---|---| +| FP16 LUTI4 + FMOPA | LUTI4 to `float16` | GEMM using FMOPA | Use LUTI and source segments | +| LUTI4 -> LUTI2 -> SDOT | LUTI4, then LUTI2 to `int8` | GEMV using SDOT | Use multiple lookup tables and source segments | + + +For every LUTI call, answer the following 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? | Select the correct 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. + + +## 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). +It uses 4-bit codes as indices that map to `float16` values, and shows how the source-index segment is interpreted relative to the destination-group size using generic SVL terminology. + +Open `example_2_luti_programming.c` and find `arm_lp_gemm_luti4` to follow along: + +```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 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. + +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: + +```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 | + +--------+--------+--------+--------+ + +``` +This encoding reduces the bytes required to store weights. + +The `arm_lp_gemv_luti2_luti4` function in `example_2_luti_programming.c` implements this two-stage decode: + +```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 and what's next + +You've seen how destination-group size determines the meaning of a source segment. You've also learned how the same four-step method applies once or repeatedly in a multi-stage decode. + +Next, you'll build and validate both programming examples. diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/07_build_validate_luti_programming.md b/content/learning-paths/mobile-graphics-and-gaming/luti/07_build_validate_luti_programming.md new file mode 100644 index 0000000000..220849e4f6 --- /dev/null +++ b/content/learning-paths/mobile-graphics-and-gaming/luti/07_build_validate_luti_programming.md @@ -0,0 +1,78 @@ +--- +title: Build and validate the LUTI programming examples +description: Build and run the LUTI4 and two-stage LUTI examples, then validate their output against reference results. +weight: 8 + +### FIXED, DO NOT MODIFY +layout: learningpathall +--- + +## Build and validate the examples + +You've inspected how the examples apply the four-step LUTI method to FP16 +FMOPA and two-stage SDOT kernels. Now build both examples and validate their +results against the reference implementations. + +### Build and run on macOS + +Build and run the executable on an SME2-supported device: + +```bash +make example_2_luti_programming +./example_2_luti_programming +``` + +### Cross-compile the examples and run on Android + +On macOS or Linux, build the AArch64 Android executable with LLVM 22 and the NDK r29 installation selected by `ANDROID_NDK_HOME`: + +```bash +make example_2_luti_programming_android +``` + +With an Android device connected through `adb`, push the executable to the device: + +```bash +adb push example_2_luti_programming_android /data/local/tmp/example_2_luti_programming_android +``` + +Open an `adb` shell, make the file executable, and run it: + +```bash +adb shell +cd /data/local/tmp +chmod 755 example_2_luti_programming_android +./example_2_luti_programming_android +``` + +After it finishes, enter `exit` to return to the build host's shell. + +The executable performs the definitive runtime check. If SME2 isn't available, it prints `SKIP: No support for SME2 on this device.` +The program exits successfully without running either set of examples. +You can still disassemble the executable on the build host. + +### Check the result + +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`). + +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. + +## What you've accomplished and what's next + +You've built the LUTI4 and two-stage LUTI examples and confirmed that both +kernels match their reference results. + +Next, you'll compare the ZT0-based LUTI path with Z-register table forms and +SME2.1 strided destination groups. diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/08_additional_luti_features.md b/content/learning-paths/mobile-graphics-and-gaming/luti/08_additional_luti_features.md new file mode 100644 index 0000000000..0ba0f440a2 --- /dev/null +++ b/content/learning-paths/mobile-graphics-and-gaming/luti/08_additional_luti_features.md @@ -0,0 +1,117 @@ +--- +title: Identify additional LUTI features +description: Compare ZT0 and Z-register lookup tables, SME2.1 destination forms, and their compile-time feature macros. +weight: 9 + +### FIXED, DO NOT MODIFY +layout: learningpathall +--- + +## Compare feature paths + +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 lookup-table instruction (LUTI) features use a different table source or add specialized forms: + +| 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 | Adds Z-register table forms of LUTI2 and LUTI4 which produce one Z-register result without using `ZT0`. | + +{{% notice Note %}} `FEAT_LUT` and `FEAT_SME2p1` aren't currently implemented on any hardware. The following code excerpts are illustrative. Use the compiler feature macros to guard these paths and prepare your kernel for when hardware support for these features becomes available. {{% /notice %}} + +### Use Z-register tables with FEAT_LUT + +`FEAT_LUT` expands the functionality of `FEAT_SME2`, allowing the use of scalable Z registers as the lookup-table source: + +```text +table Z register(s) + packed-index Z register + | + LUTI2/LUTI4 + | + one result Z register +``` + +This provides an alternative to using the `ZT0` register. Use this form for vector kernels that don't need the `ZT0` register and need multiple LUTs. + +For example, the two-stage decode loop in the previous example reloads `ZT0` using `svldr_zt()` for its LUTI4 and LUTI2 tables. + +With `FEAT_LUT`, both lookup tables are kept in separate Z registers and loaded once before the loop with the appropriate predicate. This removes the need to call `svldr_zt()` inside the loop. + +The following excerpt illustrates one segment of the two-stage decode: + +```c +// Load 16-entries LUTI4-Table and 4-entries LUTI2-Table before the loop. +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; + + // First stage: Look up the codewords using the LUTI4 table. + // No need to call svldr_zt(). + svuint8_t codewords = svluti4_lane_u8( + luti4_table_z, packed_indices, /* segment */ 0); + + // Second stage: Decode signed values using the LUTI2 table. + // No need to call svldr_zt(). + svint8_t values = svreinterpret_s8_u8( + svluti2_lane_u8(luti2_table_z, codewords, /* segment */ 0)); + + // Consume values, then continue with the next input block. + .. +``` + +Each Z-register-table LUTI instruction produces one destination Z register. +The table entries are the one or two SVL-sized Z registers. The logical LUT is still the four entries for LUTI2 or sixteen entries for LUTI4. + +### Place results with FEAT_SME2p1 + +The SME2 examples use consecutive destination registers. `FEAT_SME2p1` extends the `ZT0` SME2 forms with strided destination pairs and quads: + +```text +Consecutive pair: { z0, z1 } +Strided pair: { z0, z8 } + +Consecutive quad: { z0, z1, z2, z3 } +Strided quad: { z0, z4, z8, z12 } +``` + +For example, the following LUTI2 instruction writes a strided destination quad: +```asm +luti2 {z0.b, z4.b, z8.b, z12.b}, zt0, z1[0] +``` + +This form gives the register allocator more placement options. It doesn't change +the index width, table contents, or number of results produced by the +instruction. + +## Guard feature paths with compiler feature macros + +Use Arm C Language Extensions (ACLE) macros to guard the currently standardized intrinsic families: + +```c +#if defined(__ARM_FEATURE_SME2) +// Original ZT0-based LUTI2 and LUTI4 forms. +#endif + +#if defined(__ARM_FEATURE_LUT) && \ + (defined(__ARM_FEATURE_SVE2) || defined(__ARM_FEATURE_SME2)) +// Z-register table forms. +#endif + +#if defined(__ARM_FEATURE_SME2p1) +// SME2.1 forms, including strided destination groups. +#endif +``` + +These macros describe the compiler target. Runtime dispatch separately checks +that the operating system exposes the required feature on the current processor. + +## What you've learned + +You've identified the table source, execution state, output shape, and feature macro for the main LUTI variants. Use these distinctions when selecting an instruction form and when adding compile-time and runtime feature checks to production code. + +You can now choose between `ZT0`-based SME2 forms, Z-register table forms with `FEAT_LUT`, and strided destination groups with `FEAT_SME2p1`, depending on the register pressure and algorithm structure of your kernel. diff --git a/content/learning-paths/mobile-graphics-and-gaming/luti/_index.md b/content/learning-paths/mobile-graphics-and-gaming/luti/_index.md index 93ce1325e4..0f2c24fdac 100644 --- a/content/learning-paths/mobile-graphics-and-gaming/luti/_index.md +++ b/content/learning-paths/mobile-graphics-and-gaming/luti/_index.md @@ -1,5 +1,5 @@ --- -title: Decode low-bit weights with Arm SME2 LUTI +title: Decode low-bit weights with Arm SME2 lookup-table instructions description: Learn how to decode packed low-bit weights using Arm SME2 LUTI2 and LUTI4, and validate the results against a plain C implementation. minutes_to_complete: 60