[PATCH v3 2/2] aarch64: Add SME2.1 structured sparsity outer product intrinsics
Claudio Bantaloukas <[email protected]>
| Newsgroups | gmane.comp.gcc.patches |
|---|---|
| Message-ID | <[email protected]> |
This patch adds support for the intrinsics implementing FEAT_SME_TMOP listed
below. All of these require the +sme-tmop arch option.
A new intrinsic shape and a new register constraint is required. This patch
adds these, along with tests. Also added a pair of predicates to tighten the
immediates selecting ZA tiles based on tile width.
Specifically:
- svtmopa_lane_za16[_mf8_mf8] (requires +sme-f8f16)
- svtmopa_lane_za16[_bf16_bf16] (requires +sme-b16b16)
- svtmopa_lane_za16[_f16_f16] (requires +sme-f16f16)
- svtmopa_lane_za32[_mf8_mf8] (requires +sme-f8f32)
- svtmopa_lane_za32[_bf16_bf16] (requires +sme-b16b16)
- svtmopa_lane_za32[_f16_f16]
- svtmopa_lane_za32[_f32_f32]
- svtmopa_lane_za32[_s16_s16]
- svtmopa_lane_za32[_u16_u16]
- svtmopa_lane_za32[_s8_s8]
- svtmopa_lane_za32[_u8_u8]
- svtmopa_lane_za32[_s8_u8]
- svtmopa_lane_za32[_u8_s8]
gcc/
* config/aarch64/aarch64-acle-builtins.cc:
(function_resolver::resolve_to): Fix documentation comment.
* config/aarch64/aarch64-acle-builtins.h: (TYPES_tmop_base,
TYPES_tmop_h_bfloat, TYPES_tmop_h_float, TYPES_tmop_h_mf8,
TYPES_tmop_s_mf8): Add new DEF_SVE_TYPES_ARRAYs.
* config/aarch64/aarch64-c.cc: (__ARM_FEATURE_SME_TMOP): Add define.
* config/aarch64/aarch64-sme.md:
(@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><SVE_FULL_BHI:mode>): Add
new insn.
(@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><VNx16QI_ONLY:mode>):
Likewise.
(@aarch64_sme_lane_<optab><mode><mode>):
Likewise.
(@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><SVE_FULL_HF:mode>):
Likewise.
(@aarch64_sme_lane_<optab><SME_TMOPA_FP8:mode><VNx16QI_ONLY:mode>):
Likewise.
* config/aarch64/aarch64-sve-builtins-shapes.cc:
(ternary_za_uint_dual_single): Add new shape.
* config/aarch64/aarch64-sve-builtins-shapes.h: Likewise.
* config/aarch64/aarch64-sve-builtins-sme.cc: (svtmopa_lane_za_impl):
Add new function impl class.
(svtmopa_lane_za): Add new FUNCTION.
* config/aarch64/aarch64-sve-builtins-sme.def: (svtmopa_lane): Add new
DEF_SME_ZA_FUNCTION_GS entries.
* config/aarch64/aarch64-sve-builtins-sme.h: (svtmopa_lane_za): Add new.
* config/aarch64/aarch64.h: (TARGET_STREAMING_SME_TMOP): Add new define.
* config/aarch64/constraints.md (Uwo): Add new constraint for Control
Vector Register in TMOP operations.
* config/aarch64/iterators.md: (SME_TMOPA_BHSF, SME_TMOPA_FP8): Add new mode
iterators.
(UNSPEC_SME_FTMOPA, UNSPEC_SME_FTMOPA_FP8, UNSPEC_SME_STMOPA,
UNSPEC_SME_SUTMOPA, UNSPEC_SME_USTMOPA, UNSPEC_SME_UTMOPA): Add new
unspecs.
(SME_TMOP_INT, SME_TMOP_INT_CROSS, SME_TMOP_FP, SME_TMOP_FP8): Add new
iterators.
(za_imm_bits): Add new mode_attr.
(optab): Update with new unspecs.
* config/aarch64/predicates.md: (aarch64_imm1, aarch64_imm4): Add new
predicates.
gcc/testsuite/
* gcc.target/aarch64/pragma_cpp_predefs_6.c: Add tests checking that the
sme-tmop prefef is off by default, and checks for feature dependencies.
* gcc.target/aarch64/sme2/acle-asm/test_sme2_acle.h: (TEST_ZA_TMOP) Add
testing macro.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c:
New test.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_f16_f16.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f16_f16.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f32_f32.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_s8.c: Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_u8.c: Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s16_s16.c:
Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_s8.c: Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_u8.c: Likewise.
* gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u16_u16.c:
Likewise.
* gcc.target/aarch64/sve/acle/general-c/ternary_za_uint_dual_single_1.c:
Likewise.
* lib/target-supports.exp: Add check_effective_target support for
sme-tmop.
---
gcc/config/aarch64/aarch64-acle-builtins.cc | 2 +-
gcc/config/aarch64/aarch64-acle-builtins.h | 29 +++++
gcc/config/aarch64/aarch64-c.cc | 2 +
gcc/config/aarch64/aarch64-sme.md | 112 ++++++++++++++++++
.../aarch64/aarch64-sve-builtins-shapes.cc | 39 ++++++
.../aarch64/aarch64-sve-builtins-shapes.h | 1 +
.../aarch64/aarch64-sve-builtins-sme.cc | 36 ++++++
.../aarch64/aarch64-sve-builtins-sme.def | 24 ++++
gcc/config/aarch64/aarch64-sve-builtins-sme.h | 1 +
gcc/config/aarch64/aarch64.h | 3 +
gcc/config/aarch64/constraints.md | 4 +
gcc/config/aarch64/iterators.md | 30 +++++
gcc/config/aarch64/predicates.md | 11 ++
.../gcc.target/aarch64/pragma_cpp_predefs_6.c | 17 +++
.../aarch64/sme2/acle-asm/test_sme2_acle.h | 18 +++
.../sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za16_f16_f16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c | 83 +++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_f16_f16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_f32_f32.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c | 83 +++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_s16_s16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_s8_s8.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_s8_u8.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_u16_u16.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_u8_s8.c | 76 ++++++++++++
.../sme2/acle-asm/tmopa_lane_za32_u8_u8.c | 76 ++++++++++++
.../general-c/ternary_za_uint_dual_single_1.c | 87 ++++++++++++++
gcc/testsuite/lib/target-supports.exp | 1 +
30 files changed, 1418 insertions(+), 1 deletion(-)
create mode 100644 gcc/testsuite/gcc.target/aarch64/pragma_cpp_predefs_6.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_f16_f16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_mf8_mf8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_bf16_bf16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f16_f16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_f32_f32.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_mf8_mf8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s16_s16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_s8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_s8_u8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u16_u16.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_s8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za32_u8_u8.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/acle/general-c/ternary_za_uint_dual_single_1.c
diff --git a/gcc/config/aarch64/aarch64-acle-builtins.cc b/gcc/config/aarch64/aarch64-acle-builtins.cc
index 98ba6c8d347..98b549305a9 100644
--- a/gcc/config/aarch64/aarch64-acle-builtins.cc
+++ b/gcc/config/aarch64/aarch64-acle-builtins.cc
@@ -1099,7 +1099,7 @@ function_resolver::lookup_form (mode_suffix_index mode, sve_type type)
}
/* Resolve the function to one with the mode suffix given by MODE, the
- type suffixes given by TYPE0 and TYPE1, and group suffix given by
+ type suffixes given by TYPE0, TYPE1 and TYPE2, and group suffix given by
GROUP. Return its function decl on success, otherwise report an
error and return error_mark_node. */
tree
diff --git a/gcc/config/aarch64/aarch64-acle-builtins.h b/gcc/config/aarch64/aarch64-acle-builtins.h
index 8dd5e138719..bc646df3b36 100644
--- a/gcc/config/aarch64/aarch64-acle-builtins.h
+++ b/gcc/config/aarch64/aarch64-acle-builtins.h
@@ -1815,6 +1815,30 @@ function_expander::result_mode () const
#define TYPES_mop_i16i64_unsigned(S, D, T) \
D (za64, u16)
+/* _za32 x { _s8_s8 _u8_u8
+ _s8_u8 _u8_s8
+ _s16_s16 _u16_u16
+ _bf16_bf16 _f16_f16
+ _f32_f32 }. */
+#define TYPES_tmop_base(S, D, T) \
+ T (za32, s8, s8), T (za32, u8, u8), \
+ T (za32, s8, u8), T (za32, u8, s8), \
+ T (za32, s16, s16), T (za32, u16, u16), \
+ T (za32, bf16, bf16), T (za32, f16, f16), \
+ T (za32, f32, f32)
+
+#define TYPES_tmop_h_bfloat(S, D, T) \
+ T (za16, bf16, bf16)
+
+#define TYPES_tmop_h_float(S, D, T) \
+ T (za16, f16, f16)
+
+#define TYPES_tmop_h_mf8(S, D, T) \
+ T (za16, mf8, mf8)
+
+#define TYPES_tmop_s_mf8(S, D, T) \
+ T (za32, mf8, mf8)
+
/* _za. */
#define TYPES_za(S, D, T) \
S (za)
@@ -2107,6 +2131,11 @@ DEF_SVE_TYPES_ARRAY (mop_base_unsigned);
DEF_SVE_TYPES_ARRAY (mop_i16i64);
DEF_SVE_TYPES_ARRAY (mop_i16i64_signed);
DEF_SVE_TYPES_ARRAY (mop_i16i64_unsigned);
+DEF_SVE_TYPES_ARRAY (tmop_base);
+DEF_SVE_TYPES_ARRAY (tmop_h_float);
+DEF_SVE_TYPES_ARRAY (tmop_h_bfloat);
+DEF_SVE_TYPES_ARRAY (tmop_h_mf8);
+DEF_SVE_TYPES_ARRAY (tmop_s_mf8);
DEF_SVE_TYPES_ARRAY (za);
DEF_SVE_TYPES_ARRAY (b_float);
diff --git a/gcc/config/aarch64/aarch64-c.cc b/gcc/config/aarch64/aarch64-c.cc
index 482c796d639..43e78468d9a 100644
--- a/gcc/config/aarch64/aarch64-c.cc
+++ b/gcc/config/aarch64/aarch64-c.cc
@@ -322,6 +322,8 @@ aarch64_update_cpp_builtins (cpp_reader *pfile)
aarch64_def_or_undef (TARGET_SME2p2, "__ARM_FEATURE_SME2p2", pfile);
aarch64_def_or_undef (TARGET_FAMINMAX, "__ARM_FEATURE_FAMINMAX", pfile);
aarch64_def_or_undef (TARGET_PCDPHINT, "__ARM_FEATURE_PCDPHINT", pfile);
+ aarch64_def_or_undef (AARCH64_HAVE_ISA (SME_TMOP),
+ "__ARM_FEATURE_SME_TMOP", pfile);
aarch64_def_or_undef (AARCH64_HAVE_ISA (SSVE_FEXPA),
"__ARM_FEATURE_SSVE_FEXPA", pfile);
builtin_define ("__ARM_PREFETCH_RANGE");
diff --git a/gcc/config/aarch64/aarch64-sme.md b/gcc/config/aarch64/aarch64-sme.md
index 7091f566ba8..fa4c3016e1d 100644
--- a/gcc/config/aarch64/aarch64-sme.md
+++ b/gcc/config/aarch64/aarch64-sme.md
@@ -46,11 +46,13 @@
;; ---- [INT] Dot product
;; ---- [INT] Ternary widening arithmetic on ZA slice
;; ---- [INT] Sum of outer products
+;; ---- [INT] Sparse outer product
;; ---- [FP] Dot product
;; ---- [FP8] Dot product
;; ---- [FP] Ternary arithmetic on ZA slice
;; ---- [FP] Ternary widening arithmetic on ZA slice
;; ---- [FP] Sum of outer products
+;; ---- [FP] Sparse outer product
;;
;; == Table lookup
;; ---- Table lookup
@@ -1814,6 +1816,53 @@ (define_insn "@aarch64_sme_<optab><VNx4SI_ONLY:mode><VNx4SI_ONLY:mode>"
"<optab>\tza%0.s, %1/m, %2/m, %3.s, %4.s"
)
+;; -------------------------------------------------------------------------
+;; ---- [INT] Sparse outer product
+;; -------------------------------------------------------------------------
+;; Includes:
+;; - STMOPA
+;; - UTMOPA
+;; - SUTMOPA
+;; - USTMOPA
+;; -------------------------------------------------------------------------
+;; svtmopa_lane_za32[_s16_s16]
+;; svtmopa_lane_za32[_u16_u16]
+;; svtmopa_lane_za32[_s8_s8]
+;; svtmopa_lane_za32[_u8_u8]
+(define_insn "@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><SVE_FULL_BHI:mode>"
+ [(set (reg:VNx4SI_ONLY ZA_REGNUM)
+ (unspec:VNx4SI_ONLY
+ [(reg:VNx4SI_ONLY ZA_REGNUM)
+ (reg:DI SME_STATE_REGNUM)
+ (match_operand:DI 0 "aarch64_imm2")
+ (match_operand:<SVE_FULL_BHI:VDOUBLE> 1 "aligned_register_operand" "Uw2")
+ (match_operand:SVE_FULL_BHI 2 "register_operand" "w")
+ (match_operand:VNx16QI 3 "register_operand" "Uwo")
+ (match_operand:DI 4 "aarch64_imm2")
+ ]
+ SME_TMOP_INT))]
+ "TARGET_STREAMING_SME_TMOP"
+ "<optab>\tza%0.s, %1, %2.<SVE_FULL_BHI:Vetype>, %3[%4]"
+)
+
+;; svtmopa_lane_za32[_s8_u8]
+;; svtmopa_lane_za32[_u8_s8]
+(define_insn "@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><VNx16QI_ONLY:mode>"
+ [(set (reg:VNx4SI_ONLY ZA_REGNUM)
+ (unspec:VNx4SI_ONLY
+ [(reg:VNx4SI_ONLY ZA_REGNUM)
+ (reg:DI SME_STATE_REGNUM)
+ (match_operand:DI 0 "aarch64_imm2")
+ (match_operand:<VNx16QI_ONLY:VDOUBLE> 1 "aligned_register_operand" "Uw2")
+ (match_operand:VNx16QI_ONLY 2 "register_operand" "w")
+ (match_operand:VNx16QI 3 "register_operand" "Uwo")
+ (match_operand:DI 4 "aarch64_imm2")
+ ]
+ SME_TMOP_INT_CROSS))]
+ "TARGET_STREAMING_SME_TMOP"
+ "<optab>\tza%0.s, %1, %2.<VNx16QI_ONLY:Vetype>, %3[%4]"
+)
+
;; -------------------------------------------------------------------------
;; ---- [FP] Dot product
;; -------------------------------------------------------------------------
@@ -2737,6 +2786,69 @@ (define_insn "@aarch64_sme_<optab><SME_ZA_F8F16_32:mode><VNx16QI_ONLY:mode>"
"<optab>\tza%0.<SME_ZA_F8F16_32:Vetype>, %1/m, %2/m, %3.b, %4.b"
)
+;; -------------------------------------------------------------------------
+;; ---- [FP] Sparse outer product
+;; -------------------------------------------------------------------------
+;; Includes:
+;; - BFTMOPA (SME_TMOP)
+;; - FTMOPA (SME_TMOP)
+;; -------------------------------------------------------------------------
+;; svtmopa_lane_za16[_bf16_bf16]
+;; svtmopa_lane_za16[_f16_f16]
+;; svtmopa_lane_za32[_f32_f32]
+(define_insn "@aarch64_sme_lane_<optab><mode><mode>"
+ [(set (reg:SME_TMOPA_BHSF ZA_REGNUM)
+ (unspec:SME_TMOPA_BHSF
+ [(reg:SME_TMOPA_BHSF ZA_REGNUM)
+ (reg:DI SME_STATE_REGNUM)
+ (match_operand:DI 0 "aarch64_imm<SME_TMOPA_BHSF:za_imm_bits>")
+ (match_operand:<VDOUBLE> 1 "aligned_register_operand" "Uw2")
+ (match_operand:SME_TMOPA_BHSF 2 "register_operand" "w")
+ (match_operand:VNx16QI 3 "register_operand" "Uwo")
+ (match_operand:DI 4 "aarch64_imm2")
+ ]
+ SME_TMOP_FP))]
+ "TARGET_STREAMING_SME_TMOP"
+ "<b><optab>\tza%0.<Vetype>, %1, %2.<Vetype>, %3[%4]"
+)
+
+;; svtmopa_lane_za32[_bf16_bf16]
+;; svtmopa_lane_za32[_f16_f16]
+(define_insn "@aarch64_sme_lane_<optab><VNx4SI_ONLY:mode><SVE_FULL_HF:mode>"
+ [(set (reg:SVE_FULL_HF ZA_REGNUM)
+ (unspec:SVE_FULL_HF
+ [(reg:VNx4SI_ONLY ZA_REGNUM)
+ (reg:DI SME_STATE_REGNUM)
+ (match_operand:DI 0 "aarch64_imm2")
+ (match_operand:<SVE_FULL_HF:VDOUBLE> 1 "aligned_register_operand" "Uw2")
+ (match_operand:SVE_FULL_HF 2 "register_operand" "w")
+ (match_operand:VNx16QI 3 "register_operand" "Uwo")
+ (match_operand:DI 4 "aarch64_imm2")
+ ]
+ SME_TMOP_FP))]
+ "TARGET_STREAMING_SME_TMOP"
+ "<SVE_FULL_HF:b><optab>\tza%0.s, %1, %2.<SVE_FULL_HF:Vetype>, %3[%4]"
+)
+
+;; svtmopa_lane_za16[_mf8_mf8]_fpm
+;; svtmopa_lane_za32[_mf8_mf8]_fpm
+(define_insn "@aarch64_sme_lane_<optab><SME_TMOPA_FP8:mode><VNx16QI_ONLY:mode>"
+ [(set (reg:SME_TMOPA_FP8 ZA_REGNUM)
+ (unspec:SME_TMOPA_FP8
+ [(reg:SME_TMOPA_FP8 ZA_REGNUM)
+ (reg:DI SME_STATE_REGNUM)
+ (match_operand:DI 0 "aarch64_imm<SME_TMOPA_FP8:za_imm_bits>")
+ (match_operand:<VNx16QI_ONLY:VDOUBLE> 1 "aligned_register_operand" "Uw2")
+ (match_operand:VNx16QI_ONLY 2 "register_operand" "w")
+ (match_operand:VNx16QI 3 "register_operand" "Uwo")
+ (match_operand:DI 4 "aarch64_imm2")
+ (reg:DI FPM_REGNUM)
+ ]
+ SME_TMOP_FP8))]
+ "TARGET_STREAMING_SME_TMOP"
+ "<optab>\tza%0.<SME_TMOPA_FP8:Vetype>, %1, %2.b, %3[%4]"
+)
+
;; =========================================================================
;; == Table lookup
;; =========================================================================
diff --git a/gcc/config/aarch64/aarch64-sve-builtins-shapes.cc b/gcc/config/aarch64/aarch64-sve-builtins-shapes.cc
index cf128bc9464..66c75f80111 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins-shapes.cc
+++ b/gcc/config/aarch64/aarch64-sve-builtins-shapes.cc
@@ -4940,6 +4940,45 @@ struct ternary_uintq_intq_opt_n_def
};
SHAPE (ternary_uintq_intq_opt_n);
+/* void svfoo_t0[_t1_t2](uint64_t, sv<t1>x2_t, sv<t2>_t, svuint8_t, uint64_t)
+ where the first argument is a ZA tile.
+ and the fifth argument is a control index (0-3) */
+struct ternary_za_uint_dual_single_def : public overloaded_base<1>
+{
+ void
+ build (function_builder &b, const function_group_info &group) const override
+ {
+ b.add_overloaded_functions (group, MODE_none);
+ build_all (b, "_,su64,u1,v2,vu8,su64", group, MODE_none);
+ }
+
+ bool
+ check (function_checker &c) const override
+ {
+ return (c.require_immediate_range (0, 0, c.num_za_tiles () - 1)
+ && c.require_immediate_range (4, 0, 3));
+ }
+
+ tree
+ resolve (function_resolver &r) const override
+ {
+ sve_type type1;
+ type_suffix_index type2;
+ if (!r.check_num_arguments (r.fpm_mode == FPM_set ? 6: 5)
+ || !r.require_integer_immediate (0)
+ || (type1 = r.infer_vector_or_tuple_type (1, 2)) == NUM_TYPE_SUFFIXES
+ || (type2 = r.infer_vector_type (2)) == NUM_TYPE_SUFFIXES
+ || !r.require_vector_type (3, VECTOR_TYPE_svuint8_t)
+ || !r.require_integer_immediate (4)
+ || (r.fpm_mode == FPM_set && !r.require_scalar_type (5, "uint64_t")))
+ return error_mark_node;
+
+ return r.resolve_to (r.mode_suffix_id, r.type_suffix_ids[0],
+ type1.type, type2);
+ }
+};
+SHAPE (ternary_za_uint_dual_single);
+
/* svbool_t svfoo[_<t0>](sv<t0>_t, sv<t0>_t, uint64_t)
where the final argument is an integer constant expression in the
diff --git a/gcc/config/aarch64/aarch64-sve-builtins-shapes.h b/gcc/config/aarch64/aarch64-sve-builtins-shapes.h
index 0fca00b36d1..ac3cf09c5bd 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins-shapes.h
+++ b/gcc/config/aarch64/aarch64-sve-builtins-shapes.h
@@ -240,6 +240,7 @@ namespace aarch64_acle
extern const function_shape *const ternary_uintq_intq;
extern const function_shape *const ternary_uintq_intq_lane;
extern const function_shape *const ternary_uintq_intq_opt_n;
+ extern const function_shape *const ternary_za_uint_dual_single;
extern const function_shape *const tmad;
extern const function_shape *const unary;
extern const function_shape *const unary_convert;
diff --git a/gcc/config/aarch64/aarch64-sve-builtins-sme.cc b/gcc/config/aarch64/aarch64-sve-builtins-sme.cc
index f2def0fb837..7a3a3230abd 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins-sme.cc
+++ b/gcc/config/aarch64/aarch64-sve-builtins-sme.cc
@@ -461,6 +461,41 @@ public:
}
};
+class svtmopa_lane_za_impl : public read_write_za<function_base>
+{
+public:
+ int
+ unspec_for (const function_instance &instance) const
+ {
+ if (instance.fpm_mode == FPM_set)
+ return UNSPEC_SME_FTMOPA_FP8;
+ const auto &suffix1 = instance.type_suffix (1);
+ if (!suffix1.integer_p)
+ return UNSPEC_SME_FTMOPA;
+ const auto &suffix2 = instance.type_suffix (2);
+ if (suffix1.unsigned_p && suffix2.unsigned_p)
+ return UNSPEC_SME_UTMOPA;
+ else if (!suffix1.unsigned_p && !suffix2.unsigned_p)
+ return UNSPEC_SME_STMOPA;
+ else if (suffix1.unsigned_p && !suffix2.unsigned_p)
+ return UNSPEC_SME_USTMOPA;
+ else
+ return UNSPEC_SME_SUTMOPA;
+ }
+
+ rtx
+ expand (function_expander &e) const override
+ {
+ machine_mode za_mode = e.vector_mode (0);
+ machine_mode v_mode = e.tuple_mode (1);
+ if (GET_MODE_UNIT_BITSIZE (za_mode) == GET_MODE_UNIT_BITSIZE (v_mode))
+ za_mode = v_mode;
+ insn_code icode
+ = code_for_aarch64_sme_lane (unspec_for (e), za_mode, v_mode);
+ return e.use_exact_insn (icode);
+ }
+};
+
class svundef_za_impl : public write_za<function_base>
{
public:
@@ -675,6 +710,7 @@ FUNCTION (svsudot_lane_za, sme_2mode_lane_function, (UNSPEC_SME_SUDOT))
FUNCTION (svsuvdot_lane_za, sme_2mode_lane_function, (UNSPEC_SME_SUVDOT))
FUNCTION (svsumopa_za, sme_2mode_function, (UNSPEC_SME_SUMOPA))
FUNCTION (svsumops_za, sme_2mode_function, (UNSPEC_SME_SUMOPS))
+FUNCTION (svtmopa_lane_za, svtmopa_lane_za_impl,)
FUNCTION (svundef_za, svundef_za_impl, )
FUNCTION (svusdot_za, sme_2mode_function, (UNSPEC_NONE, UNSPEC_SME_USDOT))
FUNCTION (svusdot_lane_za, sme_2mode_lane_function, (UNSPEC_NONE, UNSPEC_SME_USDOT))
diff --git a/gcc/config/aarch64/aarch64-sve-builtins-sme.def b/gcc/config/aarch64/aarch64-sve-builtins-sme.def
index fa2de697e0b..24d8411bb8b 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins-sme.def
+++ b/gcc/config/aarch64/aarch64-sve-builtins-sme.def
@@ -313,6 +313,30 @@ DEF_SME_ZA_FUNCTION_GS_FPM (svmla, binary_za_slice_opt_single, za_s_mf8, vg1x24,
DEF_SME_ZA_FUNCTION_GS_FPM (svmopa, binary_za_m, za_s_mf8, none, za_m, set)
#undef REQUIRED_EXTENSIONS
+#define REQUIRED_EXTENSIONS streaming_only (AARCH64_FL_SME_TMOP)
+DEF_SME_ZA_FUNCTION_GS (svtmopa_lane, ternary_za_uint_dual_single, tmop_base, none, none)
+#undef REQUIRED_EXTENSIONS
+
+#define REQUIRED_EXTENSIONS streaming_only (AARCH64_FL_SME_TMOP \
+ | AARCH64_FL_SME_F16F16)
+DEF_SME_ZA_FUNCTION_GS (svtmopa_lane, ternary_za_uint_dual_single, tmop_h_float, none, none)
+#undef REQUIRED_EXTENSIONS
+
+#define REQUIRED_EXTENSIONS streaming_only (AARCH64_FL_SME_TMOP \
+ | AARCH64_FL_SME_B16B16)
+DEF_SME_ZA_FUNCTION_GS (svtmopa_lane, ternary_za_uint_dual_single, tmop_h_bfloat, none, none)
+#undef REQUIRED_EXTENSIONS
+
+#define REQUIRED_EXTENSIONS streaming_only (AARCH64_FL_SME_TMOP \
+ | AARCH64_FL_SME_F8F16)
+DEF_SME_ZA_FUNCTION_GS_FPM (svtmopa_lane, ternary_za_uint_dual_single, tmop_h_mf8, none, none, set)
+#undef REQUIRED_EXTENSIONS
+
+#define REQUIRED_EXTENSIONS streaming_only (AARCH64_FL_SME_TMOP \
+ | AARCH64_FL_SME_F8F32)
+DEF_SME_ZA_FUNCTION_GS_FPM (svtmopa_lane, ternary_za_uint_dual_single, tmop_s_mf8, none, none, set)
+#undef REQUIRED_EXTENSIONS
+
#undef DEF_SME_ZA_FUNCTION
#undef DEF_SME_ZA_FUNCTION_GS
#undef DEF_SME_ZA_FUNCTION_GS_FPM
diff --git a/gcc/config/aarch64/aarch64-sve-builtins-sme.h b/gcc/config/aarch64/aarch64-sve-builtins-sme.h
index cc264ea8486..32d98d6f23a 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins-sme.h
+++ b/gcc/config/aarch64/aarch64-sve-builtins-sme.h
@@ -70,6 +70,7 @@ namespace aarch64_acle
extern const function_base *const svsuvdot_lane_za;
extern const function_base *const svsumopa_za;
extern const function_base *const svsumops_za;
+ extern const function_base *const svtmopa_lane_za;
extern const function_base *const svusdot_za;
extern const function_base *const svusdot_lane_za;
extern const function_base *const svusvdot_lane_za;
diff --git a/gcc/config/aarch64/aarch64.h b/gcc/config/aarch64/aarch64.h
index 11fb7a9481e..28c690db9ac 100644
--- a/gcc/config/aarch64/aarch64.h
+++ b/gcc/config/aarch64/aarch64.h
@@ -368,6 +368,9 @@ constexpr auto AARCH64_FL_DEFAULT_ISA_MODE ATTRIBUTE_UNUSED
#define TARGET_STREAMING_SME2p3 (TARGET_STREAMING && AARCH64_HAVE_ISA (SME2p3))
+#define TARGET_STREAMING_SME_TMOP \
+ (AARCH64_HAVE_ISA (SME_TMOP) && TARGET_STREAMING)
+
#define TARGET_SME_B16B16 AARCH64_HAVE_ISA (SME_B16B16)
/* ARMv8.3-A features. */
diff --git a/gcc/config/aarch64/constraints.md b/gcc/config/aarch64/constraints.md
index 847739e3270..ed3d1a82c31 100644
--- a/gcc/config/aarch64/constraints.md
+++ b/gcc/config/aarch64/constraints.md
@@ -64,6 +64,10 @@ (define_register_constraint "Uwt" "FP_REGS"
"@internal The first register in a tuple of 4 strided FPRs."
"(regno & 0xc) == 0")
+(define_register_constraint "Uwo" "FP_REGS"
+ "@internal Control Vector Register (One of Z20-Z23 or Z28-Z31)."
+ "(regno & 0x14) == 0x14")
+
(define_register_constraint "Upa" "PR_REGS"
"SVE predicate registers p0 - p15.")
diff --git a/gcc/config/aarch64/iterators.md b/gcc/config/aarch64/iterators.md
index 8bfa6429026..9ca59218a36 100644
--- a/gcc/config/aarch64/iterators.md
+++ b/gcc/config/aarch64/iterators.md
@@ -859,6 +859,13 @@ (define_mode_iterator SME_MOP_HSDF [VNx4SF
(VNx8HF "TARGET_STREAMING_SME_F16F16")
(VNx8BF "TARGET_STREAMING_SME_B16B16")])
+(define_mode_iterator SME_TMOPA_BHSF [(VNx8HF "TARGET_STREAMING_SME_F16F16")
+ (VNx8BF "TARGET_STREAMING_SME_B16B16")
+ VNx4SF])
+
+(define_mode_iterator SME_TMOPA_FP8 [(VNx8HI "TARGET_STREAMING_SME_F8F16")
+ (VNx4SI "TARGET_STREAMING_SME_F8F32")])
+
;; ------------------------------------------------------------------
;; Unspec enumerations for Advance SIMD. These could well go into
;; aarch64.md but for their use in int_iterators here.
@@ -1367,6 +1374,8 @@ (define_c_enum "unspec"
UNSPEC_SME_FMOPA
UNSPEC_SME_FMOPS
UNSPEC_SME_FSUB
+ UNSPEC_SME_FTMOPA
+ UNSPEC_SME_FTMOPA_FP8
UNSPEC_SME_LD1_HOR
UNSPEC_SME_LD1_VER
UNSPEC_SME_READ
@@ -1383,12 +1392,14 @@ (define_c_enum "unspec"
UNSPEC_SME_SMOPS
UNSPEC_SME_ST1_HOR
UNSPEC_SME_ST1_VER
+ UNSPEC_SME_STMOPA
UNSPEC_SME_SUB
UNSPEC_SME_SUB_WRITE
UNSPEC_SME_SUDOT
UNSPEC_SME_SUVDOT
UNSPEC_SME_SUMOPA
UNSPEC_SME_SUMOPS
+ UNSPEC_SME_SUTMOPA
UNSPEC_SME_UDOT
UNSPEC_SME_UVDOT
UNSPEC_SME_UMLA
@@ -1399,6 +1410,8 @@ (define_c_enum "unspec"
UNSPEC_SME_USVDOT
UNSPEC_SME_USMOPA
UNSPEC_SME_USMOPS
+ UNSPEC_SME_USTMOPA
+ UNSPEC_SME_UTMOPA
UNSPEC_SME_WRITE
UNSPEC_SME_WRITE_HOR
UNSPEC_SME_WRITE_VER
@@ -2970,6 +2983,12 @@ (define_mode_attr za16_32_last_offset [(VNx4SI "3")(VNx8HI "1")])
(define_mode_attr za32_last_offset [(VNx16QI "3") (VNx32QI "3") (VNx64QI "3")
(VNx8HI "1") (VNx16HI "1") (VNx32HI "1")])
+;; The number of bits required to specify a ZA tile of a certain element size.
+(define_mode_attr za_imm_bits [(VNx8HI "1") (VNx4SI "2") (VNx2DI "3")
+ (VNx1TI "4")
+ (VNx8BF "1") (VNx4BF "2") (VNx2BF "3")
+ (VNx8HF "1") (VNx4SF "2") (VNx2DF "3")])
+
(define_mode_attr vg_modifier [(VNx16QI "")
(VNx32QI ", vgx2")
(VNx64QI ", vgx4")
@@ -4344,6 +4363,11 @@ (define_int_iterator SME_FP8_FVDOT_HALF [
UNSPEC_SME_FVDOTT_FP8
])
+(define_int_iterator SME_TMOP_INT [UNSPEC_SME_STMOPA UNSPEC_SME_UTMOPA])
+(define_int_iterator SME_TMOP_INT_CROSS [UNSPEC_SME_SUTMOPA UNSPEC_SME_USTMOPA])
+(define_int_iterator SME_TMOP_FP [UNSPEC_SME_FTMOPA])
+(define_int_iterator SME_TMOP_FP8 [UNSPEC_SME_FTMOPA_FP8])
+
;; Iterators for atomic operations.
(define_int_iterator ATOMIC_LDOP
@@ -4504,6 +4528,8 @@ (define_int_attr optab [(UNSPEC_ANDF "and")
(UNSPEC_SME_FMOPA "fmopa")
(UNSPEC_SME_FMOPS "fmops")
(UNSPEC_SME_FSUB "fsub")
+ (UNSPEC_SME_FTMOPA "ftmopa")
+ (UNSPEC_SME_FTMOPA_FP8 "ftmopa")
(UNSPEC_SME_LD1_HOR "ld1_hor")
(UNSPEC_SME_LD1_VER "ld1_ver")
(UNSPEC_SME_READ_HOR "read_hor")
@@ -4518,6 +4544,8 @@ (define_int_attr optab [(UNSPEC_ANDF "and")
(UNSPEC_SME_SMOPS "smops")
(UNSPEC_SME_ST1_HOR "st1_hor")
(UNSPEC_SME_ST1_VER "st1_ver")
+ (UNSPEC_SME_STMOPA "stmopa")
+ (UNSPEC_SME_SUTMOPA "sutmopa")
(UNSPEC_SME_SUB "sub")
(UNSPEC_SME_SUB_WRITE "sub_write")
(UNSPEC_SME_SUDOT "sudot")
@@ -4531,9 +4559,11 @@ (define_int_attr optab [(UNSPEC_ANDF "and")
(UNSPEC_SME_UMOPA "umopa")
(UNSPEC_SME_UMOPS "umops")
(UNSPEC_SME_USDOT "usdot")
+ (UNSPEC_SME_USTMOPA "ustmopa")
(UNSPEC_SME_USVDOT "usvdot")
(UNSPEC_SME_USMOPA "usmopa")
(UNSPEC_SME_USMOPS "usmops")
+ (UNSPEC_SME_UTMOPA "utmopa")
(UNSPEC_SME_WRITE_HOR "write_hor")
(UNSPEC_SME_WRITE_VER "write_ver")
(UNSPEC_SQCADD90 "sqcadd90")
diff --git a/gcc/config/aarch64/predicates.md b/gcc/config/aarch64/predicates.md
index 6cd809cd615..15d362f04ab 100644
--- a/gcc/config/aarch64/predicates.md
+++ b/gcc/config/aarch64/predicates.md
@@ -286,6 +286,11 @@ (define_predicate "aarch64_shift_imm3"
(and (match_code "const_int")
(match_test "UINTVAL (op) <= 4")))
+;; The imm1 field is a 1-bit field that only accepts immediates 0 and 1.
+(define_predicate "aarch64_imm1"
+ (and (match_code "const_int")
+ (match_test "UINTVAL (op) <= 1")))
+
;; The imm2 field is a 2-bit field that only accepts immediates in the
;; range 0..3.
(define_predicate "aarch64_imm2"
@@ -298,6 +303,12 @@ (define_predicate "aarch64_imm3"
(and (match_code "const_int")
(match_test "UINTVAL (op) <= 7")))
+;; The imm4 field is a 4-bit field that only accepts immediates in the
+;; range 0..15.
+(define_predicate "aarch64_imm4"
+ (and (match_code "const_int")
+ (match_test "UINTVAL (op) <= 15")))
+
;; An immediate that fits into 24 bits, but needs splitting.
(define_predicate "aarch64_split_imm24"
(match_code "const_int")
diff --git a/gcc/testsuite/gcc.target/aarch64/pragma_cpp_predefs_6.c b/gcc/testsuite/gcc.target/aarch64/pragma_cpp_predefs_6.c
new file mode 100644
index 00000000000..8cd6cdc066f
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/pragma_cpp_predefs_6.c
@@ -0,0 +1,17 @@
+#pragma GCC target "arch=armv9.5-a+sme"
+#ifdef __ARM_FEATURE_SME_TMOP
+#error Foo
+#endif
+
+#pragma GCC target "arch=armv9-a+sme-tmop"
+#ifndef __ARM_FEATURE_SME_TMOP
+#error Foo
+#endif
+
+#ifndef __ARM_FEATURE_SME
+#error Foo
+#endif
+
+#ifndef __ARM_FEATURE_SME2
+#error Foo
+#endif
diff --git a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/test_sme2_acle.h b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/test_sme2_acle.h
index ff237983ad9..82e837cd7c7 100644
--- a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/test_sme2_acle.h
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/test_sme2_acle.h
@@ -121,4 +121,22 @@
INVOKE (CODE1, CODE2); \
}
+#define TEST_ZA_TMOP(NAME, TYPE1, TYPE2, CODE1, CODE2) \
+ PROTO (NAME, void, (fpm_t fpm0)) \
+ { \
+ register TYPE1 z0 __asm ("z0"); \
+ register TYPE1 z1 __asm ("z1"); \
+ register TYPE1 z2 __asm ("z2"); \
+ register TYPE2 z4 __asm ("z4"); \
+ register svuint8_t z19 __asm ("z19"); \
+ register svuint8_t z20 __asm ("z20"); \
+ register svuint8_t z24 __asm ("z24"); \
+ register svuint8_t z27 __asm ("z27"); \
+ register svuint8_t z28 __asm ("z28"); \
+ __asm volatile ("" : "=w" (z0), "=w" (z1), "=w" (z2), \
+ "=w" (z4), "=w" (z19), "=w" (z20), \
+ "=w" (z24), "=w" (z27), "=w" (z28)); \
+ INVOKE (CODE1, CODE2); \
+ }
+
#endif
diff --git a/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c
new file mode 100644
index 00000000000..49fe1af0e93
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sme2/acle-asm/tmopa_lane_za16_bf16_bf16.c
@@ -0,0 +1,76 @@
+/* { dg-do assemble { target { aarch64_asm_sme-b16b16_ok && aarch64_asm_sme-tmop_ok} } } */
+/* { dg-do compile { target { ! { aarch64_asm_sme-b16b16_ok && aarch64_asm_sme-tmop_ok } } } } */
+/* { dg-final { check-function-bodies "**" "" "-DCHECK_ASM" } } */
+
+#include "test_sme2_acle.h"
+
+#pragma GCC target "+sme-tmop+sme-b16b16"
+
+/*
+** tmopa_lane_za16_bf16_bf16_0_z0_z4_z20_0:
+** bftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z4_z20_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za16 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with different values.
+** tmopa_lane_za16_bf16_bf16_1_z2_z4_z20_3:
+** bftmopa za1\.h, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_1_z2_z4_z20_3, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (1, z2, z4, z20, 3),
+ svtmopa_lane_za16 (1, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za16_bf16_bf16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** bftmopa za0\.h, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z1_z4_z20_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za16 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** bftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z4_z19_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za16 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** bftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z4_z24_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za16 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** bftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z4_z27_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za16 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_bf16_bf16_0_z0_z4_z28_0:
+** bftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_bf16_bf16_0_z0_z4_z28_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za16_bf16_bf16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za16 (0, z0, z4, 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..2c746f287a2
--- /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_z4_z20_0:
+** ftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z4_z20_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za16 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with different values.
+** tmopa_lane_za16_f16_f16_1_z2_z4_z20_3:
+** ftmopa za1\.h, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_1_z2_z4_z20_3, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (1, z2, z4, z20, 3),
+ svtmopa_lane_za16 (1, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za16_f16_f16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.h, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z1_z4_z20_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za16 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_f16_f16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z4_z19_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za16 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_f16_f16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z4_z24_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za16 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_f16_f16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** ftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z4_z27_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za16 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_f16_f16_0_z0_z4_z28_0:
+** ftmopa za0\.h, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_f16_f16_0_z0_z4_z28_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za16_f16_f16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za16 (0, z0, z4, 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..6adca3a08a4
--- /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_z4_z20_0:
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z4_z20_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z4, z20, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z4, z20, 0, fpm0))
+
+/* ZA slice and offset with different values.
+** tmopa_lane_za16_mf8_mf8_1_z2_z4_z20_3:
+** msr fpmr, x0
+** ftmopa za1\.h, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_1_z2_z4_z20_3, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (1, z2, z4, z20, 3, fpm0),
+ svtmopa_lane_za16_fpm (1, z2, z4, z20, 3, fpm0))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za16_mf8_mf8_0_z1_z4_z20_0:
+** msr fpmr, x0
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.h, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z1_z4_z20_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z1, z4, z20, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z1, z4, z20, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z4_z19_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z4, z19, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z4, z19, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z4_z24_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z4, z24, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z4, z24, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z4_z27_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z4, z27, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z4, z27, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za16_mf8_mf8_0_z0_z4_z28_0:
+** msr fpmr, x0
+** ftmopa za0\.h, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za16_mf8_mf8_0_z0_z4_z28_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za16_mf8_mf8_fpm (0, z0, z4, z28, 0, fpm0),
+ svtmopa_lane_za16_fpm (0, z0, z4, 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..0aca9883a26
--- /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_z4_z20_0:
+** bftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z4_z20_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_bf16_bf16_3_z2_z4_z20_3:
+** bftmopa za3\.s, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_3_z2_z4_z20_3, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_bf16_bf16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** bftmopa za0\.s, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z1_z4_z20_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z4_z19_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z4_z24_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** bftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z4_z27_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_bf16_bf16_0_z0_z4_z28_0:
+** bftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_bf16_bf16_0_z0_z4_z28_0, svbfloat16x2_t, svbfloat16_t,
+ svtmopa_lane_za32_bf16_bf16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..c34fd1e8f72
--- /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_z4_z20_0:
+** ftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z4_z20_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_f16_f16_3_z2_z4_z20_3:
+** ftmopa za3\.s, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_3_z2_z4_z20_3, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_f16_f16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z1_z4_z20_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f16_f16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z4_z19_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f16_f16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z4_z24_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f16_f16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** ftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z4_z27_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f16_f16_0_z0_z4_z28_0:
+** ftmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f16_f16_0_z0_z4_z28_0, svfloat16x2_t, svfloat16_t,
+ svtmopa_lane_za32_f16_f16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..12398c7aeaf
--- /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_z4_z20_0:
+** ftmopa za0\.s, {z0\.s - z1\.s}, z4\.s, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z4_z20_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_f32_f32_3_z2_z4_z20_3:
+** ftmopa za3\.s, {z2\.s - z3\.s}, z4\.s, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_3_z2_z4_z20_3, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_f32_f32_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.s - \2\.s}, z4\.s, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z1_z4_z20_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f32_f32_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z4\.s, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z4_z19_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f32_f32_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z4\.s, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z4_z24_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f32_f32_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** ftmopa za0\.s, {z0\.s - z1\.s}, z4\.s, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z4_z27_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_f32_f32_0_z0_z4_z28_0:
+** ftmopa za0\.s, {z0\.s - z1\.s}, z4\.s, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_f32_f32_0_z0_z4_z28_0, svfloat32x2_t, svfloat32_t,
+ svtmopa_lane_za32_f32_f32 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..e73053180a7
--- /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_z4_z20_0:
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z4_z20_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z4, z20, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z4, z20, 0, fpm0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_mf8_mf8_3_z2_z4_z20_3:
+** msr fpmr, x0
+** ftmopa za3\.s, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_3_z2_z4_z20_3, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (3, z2, z4, z20, 3, fpm0),
+ svtmopa_lane_za32_fpm (3, z2, z4, z20, 3, fpm0))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_mf8_mf8_0_z1_z4_z20_0:
+** msr fpmr, x0
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ftmopa za0\.s, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z1_z4_z20_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z1, z4, z20, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z1, z4, z20, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z4_z19_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z4, z19, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z4, z19, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z4_z24_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z4, z24, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z4, z24, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z4_z27_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z4, z27, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z4, z27, 0, fpm0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_mf8_mf8_0_z0_z4_z28_0:
+** msr fpmr, x0
+** ftmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_mf8_mf8_0_z0_z4_z28_0, svmfloat8x2_t, svmfloat8_t,
+ svtmopa_lane_za32_mf8_mf8_fpm (0, z0, z4, z28, 0, fpm0),
+ svtmopa_lane_za32_fpm (0, z0, z4, 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..42b47fdfd4a
--- /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_z4_z20_0:
+** stmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z4_z20_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s16_s16_3_z2_z4_z20_3:
+** stmopa za3\.s, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_3_z2_z4_z20_3, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s16_s16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** stmopa za0\.s, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z1_z4_z20_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s16_s16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z4_z19_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s16_s16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z4_z24_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s16_s16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** stmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z4_z27_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s16_s16_0_z0_z4_z28_0:
+** stmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s16_s16_0_z0_z4_z28_0, svint16x2_t, svint16_t,
+ svtmopa_lane_za32_s16_s16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..6eb02f8a5cd
--- /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_z4_z20_0:
+** stmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z4_z20_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s8_s8_3_z2_z4_z20_3:
+** stmopa za3\.s, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_3_z2_z4_z20_3, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s8_s8_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** stmopa za0\.s, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z1_z4_z20_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_s8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z4_z19_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_s8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z4_z24_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_s8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** stmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z4_z27_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_s8_0_z0_z4_z28_0:
+** stmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_s8_0_z0_z4_z28_0, svint8x2_t, svint8_t,
+ svtmopa_lane_za32_s8_s8 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..9057aa660ba
--- /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_z4_z20_0:
+** sutmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z4_z20_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_s8_u8_3_z2_z4_z20_3:
+** sutmopa za3\.s, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_3_z2_z4_z20_3, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_s8_u8_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** sutmopa za0\.s, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z1_z4_z20_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_u8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z4_z19_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_u8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z4_z24_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_u8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** sutmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z4_z27_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_s8_u8_0_z0_z4_z28_0:
+** sutmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_s8_u8_0_z0_z4_z28_0, svint8x2_t, svuint8_t,
+ svtmopa_lane_za32_s8_u8 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..5aa23f77a43
--- /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_z4_z20_0:
+** utmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z4_z20_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u16_u16_3_z2_z4_z20_3:
+** utmopa za3\.s, {z2\.h - z3\.h}, z4\.h, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_3_z2_z4_z20_3, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u16_u16_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** utmopa za0\.s, {\1\.h - \2\.h}, z4\.h, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z1_z4_z20_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u16_u16_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z4_z19_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u16_u16_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z4_z24_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u16_u16_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** utmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z4_z27_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u16_u16_0_z0_z4_z28_0:
+** utmopa za0\.s, {z0\.h - z1\.h}, z4\.h, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u16_u16_0_z0_z4_z28_0, svuint16x2_t, svuint16_t,
+ svtmopa_lane_za32_u16_u16 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..76ed6019793
--- /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_z4_z20_0:
+** ustmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z4_z20_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u8_s8_3_z2_z4_z20_3:
+** ustmopa za3\.s, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_3_z2_z4_z20_3, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u8_s8_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** ustmopa za0\.s, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z1_z4_z20_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_s8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z4_z19_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_s8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z4_z24_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_s8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** ustmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z4_z27_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_s8_0_z0_z4_z28_0:
+** ustmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_s8_0_z0_z4_z28_0, svuint8x2_t, svint8_t,
+ svtmopa_lane_za32_u8_s8 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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..075b5da689b
--- /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_z4_z20_0:
+** utmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z4_z20_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z4, z20, 0),
+ svtmopa_lane_za32 (0, z0, z4, z20, 0))
+
+/* ZA slice and offset with maximum values.
+** tmopa_lane_za32_u8_u8_3_z2_z4_z20_3:
+** utmopa za3\.s, {z2\.b - z3\.b}, z4\.b, z20\[3\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_3_z2_z4_z20_3, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (3, z2, z4, z20, 3),
+ svtmopa_lane_za32 (3, z2, z4, z20, 3))
+
+/* The first register on the second argument must be even.
+** tmopa_lane_za32_u8_u8_0_z1_z4_z20_0:
+** mov (z\d+)\.d, z1\.d
+** mov (z\d+)\.d, z2\.d
+** utmopa za0\.s, {\1\.b - \2\.b}, z4\.b, z20\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z1_z4_z20_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z1, z4, z20, 0),
+ svtmopa_lane_za32 (0, z1, z4, z20, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_u8_0_z0_z4_z19_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z19.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z4_z19_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z4, z19, 0),
+ svtmopa_lane_za32 (0, z0, z4, z19, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_u8_0_z0_z4_z24_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z24.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z4_z24_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z4, z24, 0),
+ svtmopa_lane_za32 (0, z0, z4, z24, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_u8_0_z0_z4_z27_0:
+** mov z(2[0-3]|2[89]|3[01]).d, z27.d
+** utmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z\1\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z4_z27_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z4, z27, 0),
+ svtmopa_lane_za32 (0, z0, z4, z27, 0))
+
+/* zk register must be one of Z20-Z23 or Z28-z31.
+** tmopa_lane_za32_u8_u8_0_z0_z4_z28_0:
+** utmopa za0\.s, {z0\.b - z1\.b}, z4\.b, z28\[0\]
+** ret
+*/
+TEST_ZA_TMOP (tmopa_lane_za32_u8_u8_0_z0_z4_z28_0, svuint8x2_t, svuint8_t,
+ svtmopa_lane_za32_u8_u8 (0, z0, z4, z28, 0),
+ svtmopa_lane_za32 (0, z0, z4, 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 987e57dfac7..1e7c133971d 100644
--- a/gcc/testsuite/lib/target-supports.exp
+++ b/gcc/testsuite/lib/target-supports.exp
@@ -13078,6 +13078,7 @@ set exts {
"sme-f8f16" "sme-f8f32"
"sme-b16b16" "sme-f16f16" "sme-i16i64" "sme" "sme2" "sme2p1" "sme2p2"
"sme2p3"
+ "sme-tmop"
"ssve-fp8dot2" "ssve-fp8dot4" "ssve-fp8fma" "sve-bfscale" "sme-lutv2"
"ssve-fexpa" "ssve-bitperm" "ssve-aes"
}
--
2.47.3