+** bftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z3_z19_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za16 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** bftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z3_z24_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za16 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** bftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z3_z27_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za16 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z3_z28_0:
+** bftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z3_z28_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za16 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_f16_f16.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_f16_f16.c
new file mode 100644
index 00000000000..7b0aeefe45e
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_f16_f16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target { aarch64_asm_sme-f16f16_ok &&
aarch64_asm_sme-tmop_ok} } } */
+/* { dg-do compile { target { ! { aarch64_asm_sme-f16f16_ok &&
aarch64_asm_sme-tmop_ok } } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop+sme-f16f16"
+
+/*
+** tmopa_lane_za16_f16_f16_0_z0_z3_z20_0:
+** ftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z3_z20_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z3, z20, 0),
+ svtmopa_lane_za16 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with different values.
+** tmopa_lane_za16_f16_f16_1_z2_z3_z20_3:
+** ftmopa za1\.h, {z2\.h - z3\.h}, z3\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_1_z2_z3_z20_3, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (1, z2, z3, z20, 3),
+ svtmopa_lane_za16 (1, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za16_f16_f16_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.h, {\1\.h - \2\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z1_z3_z20_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z1, z3, z20, 0),
+ svtmopa_lane_za16 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_f16_f16_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z3_z19_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za16 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_f16_f16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z3_z24_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za16 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_f16_f16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z3_z27_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za16 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_f16_f16_0_z0_z3_z28_0:
+** ftmopa za0\.h, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z3_z28_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za16 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c
new file mode 100644
index 00000000000..b5381ce37b9
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c
@@ -0,0 +1,83 @@
+/* { dg-do assemble { target { aarch64_asm_sme-f16f16_ok &&
aarch64_asm_sme-tmop_ok} } } */
+/* { dg-do compile { target { ! { aarch64_asm_sme-f16f16_ok &&
aarch64_asm_sme-tmop_ok } } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop+sme-f8f16"
+
+/*
+** tmopa_lane_za16_mf8_mf8_0_z0_z3_z20_0:
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z3_z20_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z3, z20, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z3, z20, 0, fpm0))
+
+/* ZA slice and offset with different values.
+** tmopa_lane_za16_mf8_mf8_1_z2_z3_z20_3:
+** msr fpmr, x0
+** ftmopa za1\.h, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_1_z2_z3_z20_3, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (1, z2, z3, z20, 3, fpm0),
+ svtmopa_lane_za16_fpm (1, z2, z3, z20, 3, fpm0))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za16_mf8_mf8_0_z1_z3_z20_0:
+** msr fpmr, x0
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.h, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z1_z3_z20_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z1, z3, z20, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z1, z3, z20, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z3_z19_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z3, z19, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z3, z19, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z3_z24_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z3, z24, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z3, z24, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z3_z27_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z3, z27, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z3, z27, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z3_z28_0:
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z3_z28_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z3, z28, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z3, z28, 0, fpm0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c
new file mode 100644
index 00000000000..854961df988
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_bf16_bf16_0_z0_z3_z20_0:
+** bftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z3_z20_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_bf16_bf16_3_z2_z3_z20_3:
+** bftmopa za3\.s, {z2\.h - z3\.h}, z3\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_3_z2_z3_z20_3, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_bf16_bf16_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** bftmopa za0\.s, {\1\.h - \2\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z1_z3_z20_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z3_z19_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z3_z24_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z3_z27_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z3_z28_0:
+** bftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z3_z28_0, svbfloat16x2_t,
svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f16_f16.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f16_f16.c
new file mode 100644
index 00000000000..dad5daac214
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f16_f16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_f16_f16_0_z0_z3_z20_0:
+** ftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z3_z20_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_f16_f16_3_z2_z3_z20_3:
+** ftmopa za3\.s, {z2\.h - z3\.h}, z3\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_3_z2_z3_z20_3, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_f16_f16_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.h - \2\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z1_z3_z20_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f16_f16_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z3_z19_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f16_f16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z3_z24_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f16_f16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z3_z27_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f16_f16_0_z0_z3_z28_0:
+** ftmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z3_z28_0, svfloat16x2_t,
svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f32_f32.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f32_f32.c
new file mode 100644
index 00000000000..c61d2f08ed5
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f32_f32.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_f32_f32_0_z0_z3_z20_0:
+** ftmopa za0\.s, {z0\.s - z1\.s}, z3\.s, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z3_z20_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_f32_f32_3_z2_z3_z20_3:
+** ftmopa za3\.s, {z2\.s - z3\.s}, z3\.s, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_3_z2_z3_z20_3, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_f32_f32_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.s - \2\.s}, z3\.s, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z1_z3_z20_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f32_f32_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z3\.s, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z3_z19_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f32_f32_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z3\.s, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z3_z24_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f32_f32_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z3\.s, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z3_z27_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_f32_f32_0_z0_z3_z28_0:
+** ftmopa za0\.s, {z0\.s - z1\.s}, z3\.s, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z3_z28_0, svfloat32x2_t,
svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c
new file mode 100644
index 00000000000..5eca7c6c477
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c
@@ -0,0 +1,83 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop+sme-f8f32"
+
+/*
+** tmopa_lane_za32_mf8_mf8_0_z0_z3_z20_0:
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z3_z20_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z3, z20, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z3, z20, 0, fpm0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_mf8_mf8_3_z2_z3_z20_3:
+** msr fpmr, x0
+** ftmopa za3\.s, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_3_z2_z3_z20_3, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (3, z2, z3, z20, 3, fpm0),
+ svtmopa_lane_za32_fpm (3, z2, z3, z20, 3, fpm0))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_mf8_mf8_0_z1_z3_z20_0:
+** msr fpmr, x0
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z1_z3_z20_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z1, z3, z20, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z1, z3, z20, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z3_z19_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z3, z19, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z3, z19, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z3_z24_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z3, z24, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z3, z24, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z3_z27_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z3, z27, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z3, z27, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z3_z28_0:
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z3_z28_0, svmfloat8x2_t,
svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z3, z28, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z3, z28, 0, fpm0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s16_s16.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s16_s16.c
new file mode 100644
index 00000000000..c8533a76d46
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s16_s16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_s16_s16_0_z0_z3_z20_0:
+** stmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z3_z20_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s16_s16_3_z2_z3_z20_3:
+** stmopa za3\.s, {z2\.h - z3\.h}, z3\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_3_z2_z3_z20_3, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s16_s16_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** stmopa za0\.s, {\1\.h - \2\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z1_z3_z20_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s16_s16_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z3_z19_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s16_s16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z3_z24_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s16_s16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z3_z27_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s16_s16_0_z0_z3_z28_0:
+** stmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z3_z28_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_s8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_s8.c
new file mode 100644
index 00000000000..65a4e976797
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_s8.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_s8_s8_0_z0_z3_z20_0:
+** stmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z3_z20_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s8_s8_3_z2_z3_z20_3:
+** stmopa za3\.s, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_3_z2_z3_z20_3, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s8_s8_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** stmopa za0\.s, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z1_z3_z20_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_s8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z3_z19_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_s8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z3_z24_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_s8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z3_z27_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_s8_0_z0_z3_z28_0:
+** stmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z3_z28_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_u8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_u8.c
new file mode 100644
index 00000000000..8bf14909516
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_u8.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_s8_u8_0_z0_z3_z20_0:
+** sutmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z3_z20_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s8_u8_3_z2_z3_z20_3:
+** sutmopa za3\.s, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_3_z2_z3_z20_3, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s8_u8_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** sutmopa za0\.s, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z1_z3_z20_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_u8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z3_z19_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_u8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z3_z24_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_u8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z3_z27_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_s8_u8_0_z0_z3_z28_0:
+** sutmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z3_z28_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u16_u16.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u16_u16.c
new file mode 100644
index 00000000000..a871cbc1aef
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u16_u16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_u16_u16_0_z0_z3_z20_0:
+** utmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z3_z20_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u16_u16_3_z2_z3_z20_3:
+** utmopa za3\.s, {z2\.h - z3\.h}, z3\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_3_z2_z3_z20_3, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u16_u16_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** utmopa za0\.s, {\1\.h - \2\.h}, z3\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z1_z3_z20_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u16_u16_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z3_z19_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u16_u16_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z3_z24_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u16_u16_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z3\.h, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z3_z27_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u16_u16_0_z0_z3_z28_0:
+** utmopa za0\.s, {z0\.h - z1\.h}, z3\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z3_z28_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_s8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_s8.c
new file mode 100644
index 00000000000..3d06044989a
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_s8.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_u8_s8_0_z0_z3_z20_0:
+** ustmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z3_z20_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u8_s8_3_z2_z3_z20_3:
+** ustmopa za3\.s, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_3_z2_z3_z20_3, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u8_s8_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ustmopa za0\.s, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z1_z3_z20_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_s8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z3_z19_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_s8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z3_z24_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_s8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z3_z27_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_s8_0_z0_z3_z28_0:
+** ustmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z3_z28_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_u8.c
b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_u8.c
new file mode 100644
index 00000000000..bd2519e0ea0
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_u8.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target aarch64_asm_sme-tmop_ok } } */
+/* { dg-do compile { target { ! aarch64_asm_sme-tmop_ok } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop"
+
+/*
+** tmopa_lane_za32_u8_u8_0_z0_z3_z20_0:
+** utmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z3_z20_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z3, z20, 0),
+ svtmopa_lane_za32 (0, z0, z3, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u8_u8_3_z2_z3_z20_3:
+** utmopa za3\.s, {z2\.b - z3\.b}, z3\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_3_z2_z3_z20_3, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (3, z2, z3, z20, 3),
+ svtmopa_lane_za32 (3, z2, z3, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u8_u8_0_z1_z3_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** utmopa za0\.s, {\1\.b - \2\.b}, z3\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z1_z3_z20_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z1, z3, z20, 0),
+ svtmopa_lane_za32 (0, z1, z3, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_u8_0_z0_z3_z19_0:
+** mov (z\d+).d, z19.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z3_z19_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z3, z19, 0),
+ svtmopa_lane_za32 (0, z0, z3, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_u8_0_z0_z3_z24_0:
+** mov (z\d+).d, z24.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z3_z24_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z3, z24, 0),
+ svtmopa_lane_za32 (0, z0, z3, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_u8_0_z0_z3_z27_0:
+** mov (z\d+).d, z27.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z3\.b, \1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z3_z27_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z3, z27, 0),
+ svtmopa_lane_za32 (0, z0, z3, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-Z31.
+** tmopa_lane_za32_u8_u8_0_z0_z3_z28_0:
+** utmopa za0\.s, {z0\.b - z1\.b}, z3\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z3_z28_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z3, z28, 0),
+ svtmopa_lane_za32 (0, z0, z3, z28, 0))
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sve/acle/general-c/ternary_za_uint_dual_single_1.c
b/gcc/testsuite/gcc.target/aarch64/sve/acle/general-c/ternary_za_uint_dual_single_1.c
new file mode 100644
index 00000000000..f1b170bb70a
--- /dev/null
+++
b/gcc/testsuite/gcc.target/aarch64/sve/acle/general-c/ternary_za_uint_dual_single_1.c
@@ -0,0 +1,87 @@
+/* { dg-do compile } */
+
+#include <arm_sme.h>
+
+#pragma GCC target ("arch=armv9-a+sme-tmop")
+
+void
+f1 (uint64_t u64,
+ svfloat32x2_t f32x2, svfloat32_t f32,
+ svfloat16x2_t f16x2, svfloat16_t f16,
+ svint8x2_t s8x2, svint8_t s8,
+ svuint8x2_t u8x2, svuint8_t u8,
+ svint16_t s16, svuint16_t u16)
+ __arm_streaming __arm_inout("za")
+{
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, 0);
+
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8); /* { dg-error {too few
arguments to function 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, 0, 0); /* { dg-error {too many
arguments to function 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (u64, f32x2, f32, u8, 0); /* { dg-error {argument
1 of 'svtmopa_lane_za32_f32_f32' must be an integer constant expression} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, u64); /* { dg-error {argument
5 of 'svtmopa_lane_za32_f32_f32' must be an integer constant expression} } */
+
+ svtmopa_lane_za32_f32_f32 (-1, f32x2, f32, u8, 0); /* { dg-error {passing -1
to argument 1 of 'svtmopa_lane_za32_f32_f32', which expects a value in the
range \[0, 3\]} } */
+ svtmopa_lane_za32_f32_f32 (4, f32x2, f32, u8, 0); /* { dg-error {passing 4
to argument 1 of 'svtmopa_lane_za32_f32_f32', which expects a value in the
range \[0, 3\]} } */
+
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, -1); /* { dg-error {passing -1
to argument 5 of 'svtmopa_lane_za32_f32_f32', which expects a value in the
range \[0, 3\]} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, 4); /* { dg-error {passing 4
to argument 5 of 'svtmopa_lane_za32_f32_f32', which expects a value in the
range \[0, 3\]} } */
+
+ svtmopa_lane_za32_f32_f32 (0, u8, f32, u8, 0); /* { dg-error {incompatible
type for argument 2 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32, f32, u8, 0); /* { dg-error {incompatible
type for argument 2 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f16x2, f32, u8, 0); /* { dg-error
{incompatible type for argument 2 of 'svtmopa_lane_za32_f32_f32'} } */
+
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f16, u8, 0); /* { dg-error
{incompatible type for argument 3 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32x2, u8, 0); /* { dg-error
{incompatible type for argument 3 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, u8, u8, 0); /* { dg-error {incompatible
type for argument 3 of 'svtmopa_lane_za32_f32_f32'} } */
+
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u16, 0); /* { dg-error
{incompatible type for argument 4 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, 0, 0); /* { dg-error {incompatible
type for argument 4 of 'svtmopa_lane_za32_f32_f32'} } */
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, f32, 0); /* { dg-error
{incompatible type for argument 4 of 'svtmopa_lane_za32_f32_f32'} } */
+
+ svtmopa_lane_za32_s8_u8(0, s8x2, u8, u8, 0);
+ svtmopa_lane_za32_s8_u8(0, u8x2, u8, u8, 0); /* { dg-error {incompatible
type for argument 2 of 'svtmopa_lane_za32_s8_u8'} } */
+ svtmopa_lane_za32_s8_u8(0, s8x2, s8, u8, 0); /* { dg-error {incompatible
type for argument 3 of 'svtmopa_lane_za32_s8_u8'} } */
+ svtmopa_lane_za32_u8_s8(0, s8x2, s8, u8, 0); /* { dg-error {incompatible
type for argument 2 of 'svtmopa_lane_za32_u8_s8'} } */
+ svtmopa_lane_za32_u8_s8(0, u8x2, u8, u8, 0); /* { dg-error {incompatible
type for argument 3 of 'svtmopa_lane_za32_u8_s8'} } */
+}
+
+void
+f2 (svfloat32x2_t f32x2, svfloat32_t f32, svuint8_t u8) __arm_streaming
+{
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, 0); /* { dg-error {ACLE
function 'svtmopa_lane_za32_f32_f32' can only be called from a function that
has 'za' state} } */
+}
+
+void
+f3 (svfloat32x2_t f32x2, svfloat32_t f32, svuint8_t u8) __arm_inout("za")
+{
+ svtmopa_lane_za32_f32_f32 (0, f32x2, f32, u8, 0); /* { dg-error {ACLE
function 'svtmopa_lane_za32_f32_f32' can only be called when SME streaming mode
is enabled} } */
+}
+
+#pragma GCC target ("arch=armv9-a+sme-tmop+sme-f8f16")
+
+void
+f4 (svmfloat8x2_t mf8x2, svmfloat8_t mf8, svuint8_t u8, fpm_t fpm)
+ __arm_streaming __arm_inout("za")
+{
+
+ svtmopa_lane_za16_mf8_mf8_fpm (0, mf8x2, mf8, u8); /* { dg-error {too few
arguments to function 'svtmopa_lane_za16_mf8_mf8_fpm'} } */
+ svtmopa_lane_za16_mf8_mf8_fpm (0, mf8x2, mf8, u8, 0, 0, fpm); /* { dg-error
{too many arguments to function 'svtmopa_lane_za16_mf8_mf8_fpm'} } */
+ svtmopa_lane_za16_mf8_mf8_fpm (-1, mf8x2, mf8, u8, 0, fpm); /* { dg-error
{passing -1 to argument 1 of 'svtmopa_lane_za16_mf8_mf8_fpm', which expects a
value in the range \[0, 1\]} } */
+ svtmopa_lane_za16_mf8_mf8_fpm (2, mf8x2, mf8, u8, 0, fpm); /* { dg-error
{passing 2 to argument 1 of 'svtmopa_lane_za16_mf8_mf8_fpm', which expects a
value in the range \[0, 1\]} } */
+ svtmopa_lane_za16_mf8_mf8_fpm (0, mf8x2, mf8, u8, 0, mf8); /* { dg-error
{incompatible type for argument 6 of 'svtmopa_lane_za16_mf8_mf8_fpm'} } */
+}
+
+#pragma GCC target ("arch=armv9-a+sme-tmop+sme-f16f16")
+
+void
+f5 (svfloat16x2_t f16x2, svfloat16_t f16,
+ svuint8_t u8)
+ __arm_streaming __arm_inout("za")
+{
+ svtmopa_lane_za16_f16_f16 (-1, f16x2, f16, u8, 0); /* { dg-error {passing -1
to argument 1 of 'svtmopa_lane_za16_f16_f16', which expects a value in the
range \[0, 1\]} } */
+ svtmopa_lane_za16_f16_f16 (2, f16x2, f16, u8, 0); /* { dg-error {passing 2
to argument 1 of 'svtmopa_lane_za16_f16_f16', which expects a value in the
range \[0, 1\]} } */
+
+ svtmopa_lane_za16_f16_f16 (1, f16x2, f16, u8, -1); /* { dg-error {passing -1
to argument 5 of 'svtmopa_lane_za16_f16_f16', which expects a value in the
range \[0, 3\]} } */
+ svtmopa_lane_za16_f16_f16 (1, f16x2, f16, u8, 4); /* { dg-error {passing 4
to argument 5 of 'svtmopa_lane_za16_f16_f16', which expects a value in the
range \[0, 3\]} } */
+}
+
diff --git a/gcc/testsuite/lib/target-supports.exp
b/gcc/testsuite/lib/target-supports.exp
index 2b450669c3d..066ef42d440 100644
--- a/gcc/testsuite/lib/target-supports.exp
+++ b/gcc/testsuite/lib/target-supports.exp
@@ -12683,6 +12683,7 @@ set exts_sve2 {
"sme-f8f16" "sme-f8f32"
"sme-b16b16" "sme-f16f16" "sme-i16i64" "sme" "sme2" "sme2p1"
"ssve-fp8dot2" "ssve-fp8dot4" "ssve-fp8fma"
+ "sme-tmop"
}
foreach { aarch64_ext } $exts {
--
2.51.0
