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/vwmacc_vx-i16-m1.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m4.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf4.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m1.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m4.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-mf2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m1.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m4.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf2.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf4.c: New
test.
* gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf8.c: New
test.
Signed-off-by: Pan Li <[email protected]>
---
.../rvv/autovec/group_overlap/group_overlap.h | 199 ++++++++++++++++++
.../autovec/group_overlap/vwmacc_vx-i16-m1.c | 58 +++++
.../autovec/group_overlap/vwmacc_vx-i16-m2.c | 50 +++++
.../autovec/group_overlap/vwmacc_vx-i16-m4.c | 46 ++++
.../autovec/group_overlap/vwmacc_vx-i16-mf2.c | 21 ++
.../autovec/group_overlap/vwmacc_vx-i16-mf4.c | 21 ++
.../autovec/group_overlap/vwmacc_vx-i32-m1.c | 58 +++++
.../autovec/group_overlap/vwmacc_vx-i32-m2.c | 50 +++++
.../autovec/group_overlap/vwmacc_vx-i32-m4.c | 46 ++++
.../autovec/group_overlap/vwmacc_vx-i32-mf2.c | 21 ++
.../autovec/group_overlap/vwmacc_vx-i8-m1.c | 58 +++++
.../autovec/group_overlap/vwmacc_vx-i8-m2.c | 50 +++++
.../autovec/group_overlap/vwmacc_vx-i8-m4.c | 46 ++++
.../autovec/group_overlap/vwmacc_vx-i8-mf2.c | 21 ++
.../autovec/group_overlap/vwmacc_vx-i8-mf4.c | 21 ++
.../autovec/group_overlap/vwmacc_vx-i8-mf8.c | 21 ++
16 files changed, 787 insertions(+)
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m1.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m4.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf4.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m1.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m4.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-mf2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m1.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m4.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf2.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf4.c
create mode 100644
gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-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 cfccb0fff91..dcf4747fed6 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
@@ -811,6 +811,145 @@
ST_F ((void *)out, vd14, VL); OUT += VL; \
ST_F ((void *)out, vd15, VL); OUT += VL; \
+#define LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X4(NT, WT, LD_NF, LD_WF, OUT_F, \
+ ST_F, OUT, START, X, 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; \
+ 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, X, vs0, VL); \
+ WT vd1 = OUT_F (vw1, X, vs1, VL); \
+ WT vd2 = OUT_F (vw2, X, vs2, VL); \
+ WT vd3 = OUT_F (vw3, X, vs3, 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_VX_BODY_X8(NT, WT, LD_NF, LD_WF, OUT_F, \
+ ST_F, OUT, START, X, 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; \
+ 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, X, vs0, VL); \
+ WT vd1 = OUT_F (vw1, X, vs1, VL); \
+ WT vd2 = OUT_F (vw2, X, vs2, VL); \
+ WT vd3 = OUT_F (vw3, X, vs3, VL); \
+ WT vd4 = OUT_F (vw4, X, vs4, VL); \
+ WT vd5 = OUT_F (vw5, X, vs5, VL); \
+ WT vd6 = OUT_F (vw6, X, vs6, VL); \
+ WT vd7 = OUT_F (vw7, X, vs7, 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_VX_BODY_X16(NT, WT, LD_NF, LD_WF, OUT_F, \
+ ST_F, OUT, START, X, 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; \
+ 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, X, vs0, VL); \
+ WT vd1 = OUT_F (vw1, X, vs1, VL); \
+ WT vd2 = OUT_F (vw2, X, vs2, VL); \
+ WT vd3 = OUT_F (vw3, X, vs3, VL); \
+ WT vd4 = OUT_F (vw4, X, vs4, VL); \
+ WT vd5 = OUT_F (vw5, X, vs5, VL); \
+ WT vd6 = OUT_F (vw6, X, vs6, VL); \
+ WT vd7 = OUT_F (vw7, X, vs7, VL); \
+ WT vd8 = OUT_F (vw8, X, vs8, VL); \
+ WT vd9 = OUT_F (vw9, X, vs9, VL); \
+ WT vd10 = OUT_F (vw10, X, vs10, VL); \
+ WT vd11 = OUT_F (vw11, X, vs11, VL); \
+ WT vd12 = OUT_F (vw12, X, vs12, VL); \
+ WT vd13 = OUT_F (vw13, X, vs13, VL); \
+ WT vd14 = OUT_F (vw14, X, vs14, VL); \
+ WT vd15 = OUT_F (vw15, X, vs15, 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; \
+
#define LOOP_DUAL_WIDEN_TERNARY_BODY_SU_X4(NT, NUT, WT, LD_NF, \
LD_NUF, LD_WF, OUT_F, ST_F, \
OUT, START, VL) \
@@ -1013,6 +1152,33 @@
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 the multiplier comes from
+ a scalar register, thus the narrowed source is the only vector operand that
+ can overlap the destination register group. RI_F reinterprets the widened
+ accumulator as the narrowed element type and GET_F extracts the
+ highest-numbered half of it. */
+#define LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2(NT, WT, WNT, LD_WF, \
+ RI_F, GET_F, OUT_F, \
+ ST_F, OUT, START, X, 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); \
+ WNT vr1 = RI_F (vw1); \
+ \
+ NT vs0 = GET_F (vr0, 1); \
+ NT vs1 = GET_F (vr1, 1); \
+ \
+ WT vd0 = OUT_F (vw0, X, vs0, VL); \
+ WT vd1 = OUT_F (vw1, X, vs1, VL); \
+ \
+ asm volatile("nop" ::: "memory"); \
+ \
+ 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
@@ -1188,4 +1354,37 @@
} \
}
+#define DEF_GROUP_OVERLAP_TERNARY_4(VL_F, NT, WT, ST, LD_NF, LD_WF, OUT_F, \
+ ST_F, NAME, LOOP_BODY) \
+ void test_group_overlap_##NAME##_##NT##_ternary_4(uint8_t *data, \
+ uint8_t *out, \
+ ST x, \
+ size_t limit) \
+ { \
+ uint8_t *start = data; \
+ uint8_t *end = data + limit; \
+ size_t vl = VL_F (); \
+ \
+ while (start < end) { \
+ LOOP_BODY (NT, WT, LD_NF, LD_WF, OUT_F, ST_F, out, start, x, vl); \
+ } \
+ }
+
+#define DEF_GROUP_OVERLAP_TERNARY_5(VL_F, NT, WT, WNT, ST, LD_WF, RI_F, \
+ GET_F, OUT_F, ST_F, NAME, LOOP_BODY) \
+ void test_group_overlap_##NAME##_##NT##_ternary_5(uint8_t *data, \
+ uint8_t *out, \
+ ST x, \
+ size_t limit) \
+ { \
+ uint8_t *start = data; \
+ uint8_t *end = data + limit; \
+ size_t vl = VL_F (); \
+ \
+ while (start < end) { \
+ LOOP_BODY (NT, WT, WNT, LD_WF, RI_F, GET_F, OUT_F, ST_F, out, start, \
+ x, vl); \
+ } \
+ }
+
#endif
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m1.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m1.c
new file mode 100644
index 00000000000..0bf156a496f
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m1.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e16m1,
+ vint16m1_t,
+ vint32m2_t,
+ int16_t,
+ __riscv_vle16_v_i16m1,
+ __riscv_vle32_v_i32m2,
+ __riscv_vwmacc_vx_i32m2,
+ __riscv_vse32_v_i32m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e16m1,
+ vint16m1_t,
+ vint32m2_t,
+ vint16m2_t,
+ int16_t,
+ __riscv_vle32_v_i32m2,
+ __riscv_vreinterpret_v_i32m2_i16m2,
+ __riscv_vget_v_i16m2_i16m1,
+ __riscv_vwmacc_vx_i32m2,
+ __riscv_vse32_v_i32m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v6,s0,v27([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,s0,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v14,s0,v31([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,s0,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,s0,v25([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v10,s0,v29([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v30,s0,v23([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,s0,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,s0,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,s0,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v22,s0,v21([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,s0,v18([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,s0,v19([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v18,s0,v1([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v5([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,a2,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/vwmacc_vx-i16-m2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m2.c
new file mode 100644
index 00000000000..e6660e7d55d
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m2.c
@@ -0,0 +1,50 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e16m2,
+ vint16m2_t,
+ vint32m4_t,
+ int16_t,
+ __riscv_vle16_v_i16m2,
+ __riscv_vle32_v_i32m4,
+ __riscv_vwmacc_vx_i32m4,
+ __riscv_vse32_v_i32m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e16m2,
+ vint16m2_t,
+ vint32m4_t,
+ vint16m4_t,
+ int16_t,
+ __riscv_vle32_v_i32m4,
+ __riscv_vreinterpret_v_i32m4_i16m4,
+ __riscv_vget_v_i16m4_i16m2,
+ __riscv_vwmacc_vx_i32m4,
+ __riscv_vse32_v_i32m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,a2,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,a2,v2([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v10([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,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/vwmacc_vx-i16-m4.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m4.c
new file mode 100644
index 00000000000..4689c74dc74
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-m4.c
@@ -0,0 +1,46 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e16m4,
+ vint16m4_t,
+ vint32m8_t,
+ int16_t,
+ __riscv_vle16_v_i16m4,
+ __riscv_vle32_v_i32m8,
+ __riscv_vwmacc_vx_i32m8,
+ __riscv_vse32_v_i32m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e16m4,
+ vint16m4_t,
+ vint32m8_t,
+ vint16m8_t,
+ int16_t,
+ __riscv_vle32_v_i32m8,
+ __riscv_vreinterpret_v_i32m8_i16m8,
+ __riscv_vget_v_i16m8_i16m4,
+ __riscv_vwmacc_vx_i32m8,
+ __riscv_vse32_v_i32m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v12([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v8([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v4([^0-9]|$)} 1 } } */
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,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/vwmacc_vx-i16-mf2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf2.c
new file mode 100644
index 00000000000..23c643ac542
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf2.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e16m1,
+ vint16mf2_t,
+ vint32m1_t,
+ int16_t,
+ __riscv_vle16_v_i16mf2,
+ __riscv_vle32_v_i32m1,
+ __riscv_vwmacc_vx_i32m1,
+ __riscv_vse32_v_i32m1,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf4.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf4.c
new file mode 100644
index 00000000000..e0c3a2b1764
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i16-mf4.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e16m1,
+ vint16mf4_t,
+ vint32mf2_t,
+ int16_t,
+ __riscv_vle16_v_i16mf4,
+ __riscv_vle32_v_i32mf2,
+ __riscv_vwmacc_vx_i32mf2,
+ __riscv_vse32_v_i32mf2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m1.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m1.c
new file mode 100644
index 00000000000..2cd7446a94c
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m1.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e32m1,
+ vint32m1_t,
+ vint64m2_t,
+ int32_t,
+ __riscv_vle32_v_i32m1,
+ __riscv_vle64_v_i64m2,
+ __riscv_vwmacc_vx_i64m2,
+ __riscv_vse64_v_i64m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e32m1,
+ vint32m1_t,
+ vint64m2_t,
+ vint32m2_t,
+ int32_t,
+ __riscv_vle64_v_i64m2,
+ __riscv_vreinterpret_v_i64m2_i32m2,
+ __riscv_vget_v_i32m2_i32m1,
+ __riscv_vwmacc_vx_i64m2,
+ __riscv_vse64_v_i64m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v6,s0,v27([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,s0,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v14,s0,v31([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,s0,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,s0,v25([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v10,s0,v29([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v30,s0,v23([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,s0,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,s0,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,s0,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v22,s0,v21([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,s0,v18([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,s0,v19([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v18,s0,v1([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v5([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,a2,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/vwmacc_vx-i32-m2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m2.c
new file mode 100644
index 00000000000..c867c67a4ea
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m2.c
@@ -0,0 +1,50 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e32m2,
+ vint32m2_t,
+ vint64m4_t,
+ int32_t,
+ __riscv_vle32_v_i32m2,
+ __riscv_vle64_v_i64m4,
+ __riscv_vwmacc_vx_i64m4,
+ __riscv_vse64_v_i64m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e32m2,
+ vint32m2_t,
+ vint64m4_t,
+ vint32m4_t,
+ int32_t,
+ __riscv_vle64_v_i64m4,
+ __riscv_vreinterpret_v_i64m4_i32m4,
+ __riscv_vget_v_i32m4_i32m2,
+ __riscv_vwmacc_vx_i64m4,
+ __riscv_vse64_v_i64m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,a2,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,a2,v2([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v10([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,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/vwmacc_vx-i32-m4.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m4.c
new file mode 100644
index 00000000000..96f0dcc6630
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-m4.c
@@ -0,0 +1,46 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e32m4,
+ vint32m4_t,
+ vint64m8_t,
+ int32_t,
+ __riscv_vle32_v_i32m4,
+ __riscv_vle64_v_i64m8,
+ __riscv_vwmacc_vx_i64m8,
+ __riscv_vse64_v_i64m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e32m4,
+ vint32m4_t,
+ vint64m8_t,
+ vint32m8_t,
+ int32_t,
+ __riscv_vle64_v_i64m8,
+ __riscv_vreinterpret_v_i64m8_i32m8,
+ __riscv_vget_v_i32m8_i32m4,
+ __riscv_vwmacc_vx_i64m8,
+ __riscv_vse64_v_i64m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v12([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v8([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v4([^0-9]|$)} 1 } } */
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,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/vwmacc_vx-i32-mf2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-mf2.c
new file mode 100644
index 00000000000..3d28f4d3510
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i32-mf2.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e32m1,
+ vint32mf2_t,
+ vint64m1_t,
+ int32_t,
+ __riscv_vle32_v_i32mf2,
+ __riscv_vle64_v_i64m1,
+ __riscv_vwmacc_vx_i64m1,
+ __riscv_vse64_v_i64m1,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m1.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m1.c
new file mode 100644
index 00000000000..d7007c46e9b
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m1.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m1,
+ vint8m1_t,
+ vint16m2_t,
+ int8_t,
+ __riscv_vle8_v_i8m1,
+ __riscv_vle16_v_i16m2,
+ __riscv_vwmacc_vx_i16m2,
+ __riscv_vse16_v_i16m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e8m1,
+ vint8m1_t,
+ vint16m2_t,
+ vint8m2_t,
+ int8_t,
+ __riscv_vle16_v_i16m2,
+ __riscv_vreinterpret_v_i16m2_i8m2,
+ __riscv_vget_v_i8m2_i8m1,
+ __riscv_vwmacc_vx_i16m2,
+ __riscv_vse16_v_i16m2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v6,s0,v27([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,s0,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v14,s0,v31([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,s0,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,s0,v25([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v10,s0,v29([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v30,s0,v23([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,s0,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,s0,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v26,s0,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,s0,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v22,s0,v21([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,s0,v18([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,s0,v19([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v18,s0,v1([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v5([^0-9]|$)} 1 } } */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v2,a2,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/vwmacc_vx-i8-m2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m2.c
new file mode 100644
index 00000000000..4507afe275d
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m2.c
@@ -0,0 +1,50 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m2,
+ vint8m2_t,
+ vint16m4_t,
+ int8_t,
+ __riscv_vle8_v_i8m2,
+ __riscv_vle16_v_i16m4,
+ __riscv_vwmacc_vx_i16m4,
+ __riscv_vse16_v_i16m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X8)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e8m2,
+ vint8m2_t,
+ vint16m4_t,
+ vint8m4_t,
+ int8_t,
+ __riscv_vle16_v_i16m4,
+ __riscv_vreinterpret_v_i16m4_i8m4,
+ __riscv_vget_v_i8m4_i8m2,
+ __riscv_vwmacc_vx_i16m4,
+ __riscv_vse16_v_i16m4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v12,a2,v30([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v28([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,v26([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v24([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v22([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v28,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v20,a2,v2([^0-9]|$)} 1 } }
*/
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v10([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v4,a2,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/vwmacc_vx-i8-m4.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m4.c
new file mode 100644
index 00000000000..265bd590d81
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-m4.c
@@ -0,0 +1,46 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m4,
+ vint8m4_t,
+ vint16m8_t,
+ int8_t,
+ __riscv_vle8_v_i8m4,
+ __riscv_vle16_v_i16m8,
+ __riscv_vwmacc_vx_i16m8,
+ __riscv_vse16_v_i16m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X4)
+
+DEF_GROUP_OVERLAP_TERNARY_5(
+ __riscv_vsetvlmax_e8m4,
+ vint8m4_t,
+ vint16m8_t,
+ vint8m8_t,
+ int8_t,
+ __riscv_vle16_v_i16m8,
+ __riscv_vreinterpret_v_i16m8_i8m8,
+ __riscv_vget_v_i8m8_i8m4,
+ __riscv_vwmacc_vx_i16m8,
+ __riscv_vse16_v_i16m8,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_OVERLAP_X2)
+
+/* ternary_4: the accumulator occupies the whole destination register group and
+ is live when the narrowed source is read, so the source cannot be allocated
+ inside the destination register group. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v24,a2,v12([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v0([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v8([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,v4([^0-9]|$)} 1 } } */
+
+/* ternary_5: the narrowed source is the highest-numbered half of its own
+ accumulator, thus it overlaps the highest-numbered part of the destination
+ register group. Without the group overlap the source would have to be
+ copied out to a disjoint register group first. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v16,a2,v20([^0-9]|$)} 1 } }
*/
+/* { dg-final { scan-assembler-times {vwmacc\.vx\s+v8,a2,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/vwmacc_vx-i8-mf2.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf2.c
new file mode 100644
index 00000000000..9d5cbe9bc4b
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf2.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf2_t,
+ vint16m1_t,
+ int8_t,
+ __riscv_vle8_v_i8mf2,
+ __riscv_vle16_v_i16m1,
+ __riscv_vwmacc_vx_i16m1,
+ __riscv_vse16_v_i16m1,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf4.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf4.c
new file mode 100644
index 00000000000..be663f220e0
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf4.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf4_t,
+ vint16mf2_t,
+ int8_t,
+ __riscv_vle8_v_i8mf4,
+ __riscv_vle16_v_i16mf2,
+ __riscv_vwmacc_vx_i16mf2,
+ __riscv_vse16_v_i16mf2,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
diff --git
a/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf8.c
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf8.c
new file mode 100644
index 00000000000..65f9f275dbe
--- /dev/null
+++
b/gcc/testsuite/gcc.target/riscv/rvv/autovec/group_overlap/vwmacc_vx-i8-mf8.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d" } */
+
+#include "group_overlap.h"
+
+DEF_GROUP_OVERLAP_TERNARY_4(
+ __riscv_vsetvlmax_e8m1,
+ vint8mf8_t,
+ vint16mf4_t,
+ int8_t,
+ __riscv_vle8_v_i8mf8,
+ __riscv_vle16_v_i16mf4,
+ __riscv_vwmacc_vx_i16mf4,
+ __riscv_vse16_v_i16mf4,
+ vwmacc_vx,
+ LOOP_DUAL_WIDEN_TERNARY_VX_BODY_X16)
+
+/* The fractional LMUL source has EMUL < 1, thus the widened destination
+ register group must not overlap the source at all. */
+/* { dg-final { scan-assembler-times {vwmacc\.vx} 16 } } */
+/* { dg-final { scan-assembler-not
{vwmacc\.vx\s+(v[0-9]+),[a-z0-9]+,\1([^0-9]|$)} } } */
--
2.43.0
