From: Pan Li <[email protected]>
Add test cases for register group overlap, please
note it is not overlap as much as possible.
gcc/testsuite/ChangeLog:
* gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h:
Add test helper macros.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m1.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m4.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf4.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m1.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m4.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-mf2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m1.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m4.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf2.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf4.c: New test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf8.c: New test.
Signed-off-by: Pan Li <[email protected]>
---
.../rvv/autovec/group_overlap/group_overlap.h | 237 ++++++++++++++++++
.../group_overlap/vwmaccsu_vv-i16-m1.c | 66 +++++
.../group_overlap/vwmaccsu_vv-i16-m2.c | 58 +++++
.../group_overlap/vwmaccsu_vv-i16-m4.c | 54 ++++
.../group_overlap/vwmaccsu_vv-i16-mf2.c | 23 ++
.../group_overlap/vwmaccsu_vv-i16-mf4.c | 23 ++
.../group_overlap/vwmaccsu_vv-i32-m1.c | 66 +++++
.../group_overlap/vwmaccsu_vv-i32-m2.c | 58 +++++
.../group_overlap/vwmaccsu_vv-i32-m4.c | 54 ++++
.../group_overlap/vwmaccsu_vv-i32-mf2.c | 23 ++
.../autovec/group_overlap/vwmaccsu_vv-i8-m1.c | 66 +++++
.../autovec/group_overlap/vwmaccsu_vv-i8-m2.c | 58 +++++
.../autovec/group_overlap/vwmaccsu_vv-i8-m4.c | 54 ++++
.../group_overlap/vwmaccsu_vv-i8-mf2.c | 23 ++
.../group_overlap/vwmaccsu_vv-i8-mf4.c | 23 ++
.../group_overlap/vwmaccsu_vv-i8-mf8.c | 23 ++
16 files changed, 909 insertions(+)
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m1.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m4.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf4.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m1.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m4.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-mf2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m1.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m4.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf2.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf4.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf8.c
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
index ac2abed02ba..19d546d930d 100644
--- a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/group_overlap.h
@@ -700,6 +700,176 @@
ST_F ((void *)out, vd14, VL); OUT += VL; \
ST_F ((void *)out, vd15, VL); OUT += VL; \
+#define LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X4(NT, NUT, WT, LD_NF, \
+ LD_NUF, LD_WF, OUT_F, ST_F, \
+ OUT, START, VL) \
+ NT vs0 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs1 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs2 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs3 = LD_NF ((void *)START, VL); START += VL; \
+ NUT vt0 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt1 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt2 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt3 = LD_NUF ((void *)START, VL); START += VL; \
+ WT vw0 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw1 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw2 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw3 = LD_WF ((void *)START, VL); START += VL; \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ WT vd0 = OUT_F (vw0, vs0, vt0, VL); \
+ WT vd1 = OUT_F (vw1, vs1, vt1, VL); \
+ WT vd2 = OUT_F (vw2, vs2, vt2, VL); \
+ WT vd3 = OUT_F (vw3, vs3, vt3, VL); \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ ST_F ((void *)out, vd0, VL); OUT += VL; \
+ ST_F ((void *)out, vd1, VL); OUT += VL; \
+ ST_F ((void *)out, vd2, VL); OUT += VL; \
+ ST_F ((void *)out, vd3, VL); OUT += VL; \
+
+#define LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X8(NT, NUT, WT, LD_NF, \
+ LD_NUF, LD_WF, OUT_F, ST_F, \
+ OUT, START, VL) \
+ NT vs0 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs1 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs2 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs3 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs4 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs5 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs6 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs7 = LD_NF ((void *)START, VL); START += VL; \
+ NUT vt0 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt1 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt2 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt3 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt4 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt5 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt6 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt7 = LD_NUF ((void *)START, VL); START += VL; \
+ WT vw0 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw1 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw2 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw3 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw4 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw5 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw6 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw7 = LD_WF ((void *)START, VL); START += VL; \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ WT vd0 = OUT_F (vw0, vs0, vt0, VL); \
+ WT vd1 = OUT_F (vw1, vs1, vt1, VL); \
+ WT vd2 = OUT_F (vw2, vs2, vt2, VL); \
+ WT vd3 = OUT_F (vw3, vs3, vt3, VL); \
+ WT vd4 = OUT_F (vw4, vs4, vt4, VL); \
+ WT vd5 = OUT_F (vw5, vs5, vt5, VL); \
+ WT vd6 = OUT_F (vw6, vs6, vt6, VL); \
+ WT vd7 = OUT_F (vw7, vs7, vt7, VL); \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ ST_F ((void *)out, vd0, VL); OUT += VL; \
+ ST_F ((void *)out, vd1, VL); OUT += VL; \
+ ST_F ((void *)out, vd2, VL); OUT += VL; \
+ ST_F ((void *)out, vd3, VL); OUT += VL; \
+ ST_F ((void *)out, vd4, VL); OUT += VL; \
+ ST_F ((void *)out, vd5, VL); OUT += VL; \
+ ST_F ((void *)out, vd6, VL); OUT += VL; \
+ ST_F ((void *)out, vd7, VL); OUT += VL; \
+
+#define LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16(NT, NUT, WT, LD_NF, \
+ LD_NUF, LD_WF, OUT_F, ST_F, \
+ OUT, START, VL) \
+ NT vs0 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs1 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs2 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs3 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs4 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs5 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs6 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs7 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs8 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs9 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs10 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs11 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs12 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs13 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs14 = LD_NF ((void *)START, VL); START += VL; \
+ NT vs15 = LD_NF ((void *)START, VL); START += VL; \
+ NUT vt0 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt1 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt2 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt3 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt4 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt5 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt6 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt7 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt8 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt9 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt10 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt11 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt12 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt13 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt14 = LD_NUF ((void *)START, VL); START += VL; \
+ NUT vt15 = LD_NUF ((void *)START, VL); START += VL; \
+ WT vw0 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw1 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw2 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw3 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw4 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw5 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw6 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw7 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw8 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw9 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw10 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw11 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw12 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw13 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw14 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw15 = LD_WF ((void *)START, VL); START += VL; \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ WT vd0 = OUT_F (vw0, vs0, vt0, VL); \
+ WT vd1 = OUT_F (vw1, vs1, vt1, VL); \
+ WT vd2 = OUT_F (vw2, vs2, vt2, VL); \
+ WT vd3 = OUT_F (vw3, vs3, vt3, VL); \
+ WT vd4 = OUT_F (vw4, vs4, vt4, VL); \
+ WT vd5 = OUT_F (vw5, vs5, vt5, VL); \
+ WT vd6 = OUT_F (vw6, vs6, vt6, VL); \
+ WT vd7 = OUT_F (vw7, vs7, vt7, VL); \
+ WT vd8 = OUT_F (vw8, vs8, vt8, VL); \
+ WT vd9 = OUT_F (vw9, vs9, vt9, VL); \
+ WT vd10 = OUT_F (vw10, vs10, vt10, VL); \
+ WT vd11 = OUT_F (vw11, vs11, vt11, VL); \
+ WT vd12 = OUT_F (vw12, vs12, vt12, VL); \
+ WT vd13 = OUT_F (vw13, vs13, vt13, VL); \
+ WT vd14 = OUT_F (vw14, vs14, vt14, VL); \
+ WT vd15 = OUT_F (vw15, vs15, vt15, VL); \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ ST_F ((void *)out, vd0, VL); OUT += VL; \
+ ST_F ((void *)out, vd1, VL); OUT += VL; \
+ ST_F ((void *)out, vd2, VL); OUT += VL; \
+ ST_F ((void *)out, vd3, VL); OUT += VL; \
+ ST_F ((void *)out, vd4, VL); OUT += VL; \
+ ST_F ((void *)out, vd5, VL); OUT += VL; \
+ ST_F ((void *)out, vd6, VL); OUT += VL; \
+ ST_F ((void *)out, vd7, VL); OUT += VL; \
+ ST_F ((void *)out, vd8, VL); OUT += VL; \
+ ST_F ((void *)out, vd9, VL); OUT += VL; \
+ ST_F ((void *)out, vd10, VL); OUT += VL; \
+ ST_F ((void *)out, vd11, VL); OUT += VL; \
+ ST_F ((void *)out, vd12, VL); OUT += VL; \
+ ST_F ((void *)out, vd13, VL); OUT += VL; \
+ ST_F ((void *)out, vd14, VL); OUT += VL; \
+ ST_F ((void *)out, vd15, VL); OUT += VL; \
+
/* The widened destination register group of a dual widen ternary insn is tied
to the accumulator, which is still live when the narrowed sources are read.
Feeding one narrowed source from the highest-numbered half of the
@@ -732,6 +902,39 @@
ST_F ((void *)out, vd0, VL); OUT += VL; \
ST_F ((void *)out, vd1, VL); OUT += VL; \
+/* Like LOOP_DUAL_WIDEN_TERNARY_BODY_OVERLAP_X2 but for the mixed signed and
+ unsigned narrowed sources. The first insn overlaps the destination register
+ group with the signed source, the second one with the unsigned source, thus
+ both narrowed operands are covered. RI_F reinterprets the widened
+ accumulator as the signed narrowed element type, RI_UF and RI_NUF do the
+ same for the unsigned narrowed element type, GET_F and GET_UF extract the
+ highest-numbered half of it. */
+#define LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2(NT, NUT, WT, WNT, \
+ WNUT, LD_NF, LD_NUF, LD_WF, \
+ RI_F, RI_UF, RI_NUF, GET_F, \
+ GET_UF, OUT_F, ST_F, OUT, \
+ START, VL) \
+ NUT vt0 = LD_NUF ((void *)START, VL); START += VL; \
+ NT vt1 = LD_NF ((void *)START, VL); START += VL; \
+ WT vw0 = LD_WF ((void *)START, VL); START += VL; \
+ WT vw1 = LD_WF ((void *)START, VL); START += VL; \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ WNT vr0 = RI_F (vw0); \
+ WNUT vr1 = RI_NUF (RI_UF (vw1)); \
+ \
+ NT vs0 = GET_F (vr0, 1); \
+ NUT vs1 = GET_UF (vr1, 1); \
+ \
+ WT vd0 = OUT_F (vw0, vs0, vt0, VL); \
+ WT vd1 = OUT_F (vw1, vt1, vs1, VL); \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ ST_F ((void *)out, vd0, VL); OUT += VL; \
+ ST_F ((void *)out, vd1, VL); OUT += VL; \
+
#define DEF_GROUP_OVERLAP_UNARY_0(VL_F, NT, WT, LD_F, OUT_F, ST_F, NAME, \
LOOP_BODY) \
void test_group_overlap_##NAME##_##NT##_unary_0(uint8_t *data, \
@@ -824,4 +1027,38 @@
} \
}
+#define DEF_GROUP_OVERLAP_TERNARY_2(VL_F, NT, NUT, WT, LD_NF, LD_NUF, \
+ LD_WF, OUT_F, ST_F, NAME, LOOP_BODY) \
+ void test_group_overlap_##NAME##_##NT##_ternary_2(uint8_t *data, \
+ uint8_t *out, \
+ size_t limit) \
+ { \
+ uint8_t *start = data; \
+ uint8_t *end = data + limit; \
+ size_t vl = VL_F (); \
+ \
+ while (start < end) { \
+ LOOP_BODY (NT, NUT, WT, LD_NF, LD_NUF, LD_WF, OUT_F, ST_F, out, \
+ start, vl); \
+ } \
+ }
+
+#define DEF_GROUP_OVERLAP_TERNARY_3(VL_F, NT, NUT, WT, WNT, WNUT, LD_NF, \
+ LD_NUF, LD_WF, RI_F, RI_UF, RI_NUF, \
+ GET_F, GET_UF, OUT_F, ST_F, NAME, \
+ LOOP_BODY) \
+ void test_group_overlap_##NAME##_##NT##_ternary_3(uint8_t *data, \
+ uint8_t *out, \
+ size_t limit) \
+ { \
+ uint8_t *start = data; \
+ uint8_t *end = data + limit; \
+ size_t vl = VL_F (); \
+ \
+ while (start < end) { \
+ LOOP_BODY (NT, NUT, WT, WNT, WNUT, LD_NF, LD_NUF, LD_WF, RI_F, \
+ RI_UF, RI_NUF, GET_F, GET_UF, OUT_F, ST_F, out, start, vl);\
+ } \
+ }
+
#endif
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m1.c
new file mode 100644
index 00000000000..851bbeebf0b
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m1.c
@@ -0,0 +1,66 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e16m1,
+ vint16m1_t,
+ vuint16m1_t,
+ vint32m2_t,
+ __riscv_vle16_v_i16m1,
+ __riscv_vle16_v_u16m1,
+ __riscv_vle32_v_i32m2,
+ __riscv_vwmaccsu_vv_i32m2,
+ __riscv_vse32_v_i32m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e16m1,
+ vint16m1_t,
+ vuint16m1_t,
+ vint32m2_t,
+ vint16m2_t,
+ vuint16m2_t,
+ __riscv_vle16_v_i16m1,
+ __riscv_vle16_v_u16m1,
+ __riscv_vle32_v_i32m2,
+ __riscv_vreinterpret_v_i32m2_i16m2,
+ __riscv_vreinterpret_v_i32m2_u32m2,
+ __riscv_vreinterpret_v_u32m2_u16m2,
+ __riscv_vget_v_i16m2_i16m1,
+ __riscv_vget_v_u16m2_u16m1,
+ __riscv_vwmaccsu_vv_i32m2,
+ __riscv_vse32_v_i32m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v30,v29([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v16,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v0,v31([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v14,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v30,v28,v27([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v26,v25([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v6,v20,v19([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v26,v24,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v18,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v10,v1,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v22,v21([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v14,v19,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v21,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v18,v20,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v0,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v22,v0,v1([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v5,v8([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m2.c
new file mode 100644
index 00000000000..de37d84cc70
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m2.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e16m2,
+ vint16m2_t,
+ vuint16m2_t,
+ vint32m4_t,
+ __riscv_vle16_v_i16m2,
+ __riscv_vle16_v_u16m2,
+ __riscv_vle32_v_i32m4,
+ __riscv_vwmaccsu_vv_i32m4,
+ __riscv_vse32_v_i32m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e16m2,
+ vint16m2_t,
+ vuint16m2_t,
+ vint32m4_t,
+ vint16m4_t,
+ vuint16m4_t,
+ __riscv_vle16_v_i16m2,
+ __riscv_vle16_v_u16m2,
+ __riscv_vle32_v_i32m4,
+ __riscv_vreinterpret_v_i32m4_i16m4,
+ __riscv_vreinterpret_v_i32m4_u32m4,
+ __riscv_vreinterpret_v_u32m4_u16m4,
+ __riscv_vget_v_i16m4_i16m2,
+ __riscv_vget_v_u16m4_u16m2,
+ __riscv_vwmaccsu_vv_i32m4,
+ __riscv_vse32_v_i32m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v0,v30([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v26([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v24,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v20,v18([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v16,v14([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v12,v10([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v8,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v2([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v10,v16([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v2,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m4.c
new file mode 100644
index 00000000000..75f2482a4e8
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-m4.c
@@ -0,0 +1,54 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e16m4,
+ vint16m4_t,
+ vuint16m4_t,
+ vint32m8_t,
+ __riscv_vle16_v_i16m4,
+ __riscv_vle16_v_u16m4,
+ __riscv_vle32_v_i32m8,
+ __riscv_vwmaccsu_vv_i32m8,
+ __riscv_vse32_v_i32m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e16m4,
+ vint16m4_t,
+ vuint16m4_t,
+ vint32m8_t,
+ vint16m8_t,
+ vuint16m8_t,
+ __riscv_vle16_v_i16m4,
+ __riscv_vle16_v_u16m4,
+ __riscv_vle32_v_i32m8,
+ __riscv_vreinterpret_v_i32m8_i16m8,
+ __riscv_vreinterpret_v_i32m8_u32m8,
+ __riscv_vreinterpret_v_u32m8_u16m8,
+ __riscv_vget_v_i16m8_i16m4,
+ __riscv_vget_v_u16m8_u16m4,
+ __riscv_vwmaccsu_vv_i32m8,
+ __riscv_vse32_v_i32m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v0,v4([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v20([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v16,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v8,v12([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v20,v0([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf2.c
new file mode 100644
index 00000000000..7bad3831168
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf2.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e16m1,
+ vint16mf2_t,
+ vuint16mf2_t,
+ vint32m1_t,
+ __riscv_vle16_v_i16mf2,
+ __riscv_vle16_v_u16mf2,
+ __riscv_vle32_v_i32m1,
+ __riscv_vwmaccsu_vv_i32m1,
+ __riscv_vse32_v_i32m1,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf4.c
new file mode 100644
index 00000000000..02ab01cee9a
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i16-mf4.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e16m1,
+ vint16mf4_t,
+ vuint16mf4_t,
+ vint32mf2_t,
+ __riscv_vle16_v_i16mf4,
+ __riscv_vle16_v_u16mf4,
+ __riscv_vle32_v_i32mf2,
+ __riscv_vwmaccsu_vv_i32mf2,
+ __riscv_vse32_v_i32mf2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m1.c
new file mode 100644
index 00000000000..b8e1089f0f8
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m1.c
@@ -0,0 +1,66 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e32m1,
+ vint32m1_t,
+ vuint32m1_t,
+ vint64m2_t,
+ __riscv_vle32_v_i32m1,
+ __riscv_vle32_v_u32m1,
+ __riscv_vle64_v_i64m2,
+ __riscv_vwmaccsu_vv_i64m2,
+ __riscv_vse64_v_i64m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e32m1,
+ vint32m1_t,
+ vuint32m1_t,
+ vint64m2_t,
+ vint32m2_t,
+ vuint32m2_t,
+ __riscv_vle32_v_i32m1,
+ __riscv_vle32_v_u32m1,
+ __riscv_vle64_v_i64m2,
+ __riscv_vreinterpret_v_i64m2_i32m2,
+ __riscv_vreinterpret_v_i64m2_u64m2,
+ __riscv_vreinterpret_v_u64m2_u32m2,
+ __riscv_vget_v_i32m2_i32m1,
+ __riscv_vget_v_u32m2_u32m1,
+ __riscv_vwmaccsu_vv_i64m2,
+ __riscv_vse64_v_i64m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v30,v29([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v16,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v0,v31([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v14,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v30,v28,v27([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v26,v25([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v6,v20,v19([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v26,v24,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v18,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v10,v1,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v22,v21([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v14,v19,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v21,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v18,v20,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v0,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v22,v0,v1([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v5,v8([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m2.c
new file mode 100644
index 00000000000..6feaf3e45ff
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m2.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e32m2,
+ vint32m2_t,
+ vuint32m2_t,
+ vint64m4_t,
+ __riscv_vle32_v_i32m2,
+ __riscv_vle32_v_u32m2,
+ __riscv_vle64_v_i64m4,
+ __riscv_vwmaccsu_vv_i64m4,
+ __riscv_vse64_v_i64m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e32m2,
+ vint32m2_t,
+ vuint32m2_t,
+ vint64m4_t,
+ vint32m4_t,
+ vuint32m4_t,
+ __riscv_vle32_v_i32m2,
+ __riscv_vle32_v_u32m2,
+ __riscv_vle64_v_i64m4,
+ __riscv_vreinterpret_v_i64m4_i32m4,
+ __riscv_vreinterpret_v_i64m4_u64m4,
+ __riscv_vreinterpret_v_u64m4_u32m4,
+ __riscv_vget_v_i32m4_i32m2,
+ __riscv_vget_v_u32m4_u32m2,
+ __riscv_vwmaccsu_vv_i64m4,
+ __riscv_vse64_v_i64m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v0,v30([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v26([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v24,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v20,v18([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v16,v14([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v12,v10([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v8,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v2([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v10,v16([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v2,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m4.c
new file mode 100644
index 00000000000..037bd27aa39
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-m4.c
@@ -0,0 +1,54 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e32m4,
+ vint32m4_t,
+ vuint32m4_t,
+ vint64m8_t,
+ __riscv_vle32_v_i32m4,
+ __riscv_vle32_v_u32m4,
+ __riscv_vle64_v_i64m8,
+ __riscv_vwmaccsu_vv_i64m8,
+ __riscv_vse64_v_i64m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e32m4,
+ vint32m4_t,
+ vuint32m4_t,
+ vint64m8_t,
+ vint32m8_t,
+ vuint32m8_t,
+ __riscv_vle32_v_i32m4,
+ __riscv_vle32_v_u32m4,
+ __riscv_vle64_v_i64m8,
+ __riscv_vreinterpret_v_i64m8_i32m8,
+ __riscv_vreinterpret_v_i64m8_u64m8,
+ __riscv_vreinterpret_v_u64m8_u32m8,
+ __riscv_vget_v_i32m8_i32m4,
+ __riscv_vget_v_u32m8_u32m4,
+ __riscv_vwmaccsu_vv_i64m8,
+ __riscv_vse64_v_i64m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v0,v4([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v20([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v16,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v8,v12([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v20,v0([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-mf2.c
new file mode 100644
index 00000000000..a9f2c4a7532
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i32-mf2.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e32m1,
+ vint32mf2_t,
+ vuint32mf2_t,
+ vint64m1_t,
+ __riscv_vle32_v_i32mf2,
+ __riscv_vle32_v_u32mf2,
+ __riscv_vle64_v_i64m1,
+ __riscv_vwmaccsu_vv_i64m1,
+ __riscv_vse64_v_i64m1,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m1.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m1.c
new file mode 100644
index 00000000000..5ac45092fb6
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m1.c
@@ -0,0 +1,66 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m1,
+ vint8m1_t,
+ vuint8m1_t,
+ vint16m2_t,
+ __riscv_vle8_v_i8m1,
+ __riscv_vle8_v_u8m1,
+ __riscv_vle16_v_i16m2,
+ __riscv_vwmaccsu_vv_i16m2,
+ __riscv_vse16_v_i16m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e8m1,
+ vint8m1_t,
+ vuint8m1_t,
+ vint16m2_t,
+ vint8m2_t,
+ vuint8m2_t,
+ __riscv_vle8_v_i8m1,
+ __riscv_vle8_v_u8m1,
+ __riscv_vle16_v_i16m2,
+ __riscv_vreinterpret_v_i16m2_i8m2,
+ __riscv_vreinterpret_v_i16m2_u16m2,
+ __riscv_vreinterpret_v_u16m2_u8m2,
+ __riscv_vget_v_i8m2_i8m1,
+ __riscv_vget_v_u8m2_u8m1,
+ __riscv_vwmaccsu_vv_i16m2,
+ __riscv_vse16_v_i16m2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v30,v29([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v16,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v0,v31([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v14,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v30,v28,v27([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v26,v25([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v6,v20,v19([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v26,v24,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v18,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v10,v1,v15([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v22,v21([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v14,v19,v17([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v21,v23([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v18,v20,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v0,v1([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v22,v0,v1([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v5,v8([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v2,v1,v3([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m2.c
new file mode 100644
index 00000000000..0a13ef31762
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m2.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m2,
+ vint8m2_t,
+ vuint8m2_t,
+ vint16m4_t,
+ __riscv_vle8_v_i8m2,
+ __riscv_vle8_v_u8m2,
+ __riscv_vle16_v_i16m4,
+ __riscv_vwmaccsu_vv_i16m4,
+ __riscv_vse16_v_i16m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e8m2,
+ vint8m2_t,
+ vuint8m2_t,
+ vint16m4_t,
+ vint8m4_t,
+ vuint8m4_t,
+ __riscv_vle8_v_i8m2,
+ __riscv_vle8_v_u8m2,
+ __riscv_vle16_v_i16m4,
+ __riscv_vreinterpret_v_i16m4_i8m4,
+ __riscv_vreinterpret_v_i16m4_u16m4,
+ __riscv_vreinterpret_v_u16m4_u8m4,
+ __riscv_vget_v_i8m4_i8m2,
+ __riscv_vget_v_u8m4_u8m2,
+ __riscv_vwmaccsu_vv_i16m4,
+ __riscv_vse16_v_i16m4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v0,v30([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v26([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v28,v24,v22([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v20,v18([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v20,v16,v14([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v12,v10([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v12,v8,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v2([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v10,v16([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v4,v2,v6([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m4.c
new file mode 100644
index 00000000000..a465500f782
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-m4.c
@@ -0,0 +1,54 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m4,
+ vint8m4_t,
+ vuint8m4_t,
+ vint16m8_t,
+ __riscv_vle8_v_i8m4,
+ __riscv_vle8_v_u8m4,
+ __riscv_vle16_v_i16m8,
+ __riscv_vwmaccsu_vv_i16m8,
+ __riscv_vse16_v_i16m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_3(
+ __riscv_vsetvlmax_e8m4,
+ vint8m4_t,
+ vuint8m4_t,
+ vint16m8_t,
+ vint8m8_t,
+ vuint8m8_t,
+ __riscv_vle8_v_i8m4,
+ __riscv_vle8_v_u8m4,
+ __riscv_vle16_v_i16m8,
+ __riscv_vreinterpret_v_i16m8_i8m8,
+ __riscv_vreinterpret_v_i16m8_u16m8,
+ __riscv_vreinterpret_v_u16m8_u8m8,
+ __riscv_vget_v_i8m8_i8m4,
+ __riscv_vget_v_u8m8_u8m4,
+ __riscv_vwmaccsu_vv_i16m8,
+ __riscv_vse16_v_i16m8,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_OVERLAP_X2)
+
+/* ternary_2: the accumulator occupies the whole destination register group and
+ is live when the narrowed sources are read, so no source can be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v0,v4([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v0,v28,v20([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v24,v16,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v8,v12([^0-9]|$)} 1 } } */
+
+/* ternary_3: one narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. The signed source is the overlapping one in the first insn,
+ the unsigned source in the second one. Without the group overlap the
+ sources would have to be copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v16,v20,v0([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv\s+v8,v4,v12([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-not {vmv[0-9]+r\.v} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf2.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf2.c
new file mode 100644
index 00000000000..993468887a1
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf2.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf2_t,
+ vuint8mf2_t,
+ vint16m1_t,
+ __riscv_vle8_v_i8mf2,
+ __riscv_vle8_v_u8mf2,
+ __riscv_vle16_v_i16m1,
+ __riscv_vwmaccsu_vv_i16m1,
+ __riscv_vse16_v_i16m1,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf4.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf4.c
new file mode 100644
index 00000000000..3a94240d07d
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf4.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf4_t,
+ vuint8mf4_t,
+ vint16mf2_t,
+ __riscv_vle8_v_i8mf4,
+ __riscv_vle8_v_u8mf4,
+ __riscv_vle16_v_i16mf2,
+ __riscv_vwmaccsu_vv_i16mf2,
+ __riscv_vse16_v_i16mf2,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf8.c b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf8.c
new file mode 100644
index 00000000000..a8747f5c265
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmaccsu_vv-i8-mf8.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_2(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf8_t,
+ vuint8mf8_t,
+ vint16mf4_t,
+ __riscv_vle8_v_i8mf8,
+ __riscv_vle8_v_u8mf8,
+ __riscv_vle16_v_i16mf4,
+ __riscv_vwmaccsu_vv_i16mf4,
+ __riscv_vse16_v_i16mf4,
+ vwmaccsu_vv,
+ LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap either source at all. */
+/* { dg-final { scan-assembler-times {vwmaccsu\.vv} 16 } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),\1,} } } */
+/* { dg-final { scan-assembler-not {vwmaccsu\.vv\s+(v[0-9]+),v[0-9]+,\1([^0-9]|$)} } } */
--
2.43.0
lmpx.com only provides a reader for public news (NNTP) servers. It is not
affiliated with the servers or forums shown here and is not responsible for
the content of articles, which is written by their respective authors.