Add RVV intrinsics for the non-batched Zvdota family[1]: Zvqwdota8i, Zvqwdota16i, Zvfwdota16bf and Zvfqwdota8f.
The public intrinsic names follow the riscv-rvv-intrinsic-doc naming scheme and encode VS2, VS1 and VD types in the suffix. The instruction mnemonic is selected from the VS2 type while vtype.altfmt is selected from the VS1 type. This commit also introduces the vfloat8e4m3 and vfloat8e5m2 vector types needed for the Zvfqwdota8f intrinsics. [1] https://github.com/riscv-non-isa/riscv-rvv-intrinsic-doc/commit/c50f171d24a4186a86855b7750ff0bf697c486bb gcc/ChangeLog: * config/riscv/genrvv-type-indexer.cc (float8type): New helper for vfloat8e4m3/vfloat8e5m2 type index generation. (emit_zvdota_type_indexes): New index. (emit_float8_type_index): Ditto. (main): Emit Zvdota accumulator and FP8 type indexes. * config/riscv/riscv-vector-builtins-bases.cc (zvdota_vs2_argno): New helper to locate VS2 in operand list. (zvdota_int_altfmt): Select ALTFMT from VS1 signedness. (zvdota_fp8_altfmt): Select ALTFMT from VS1 FP8 format. (zvdota_unsigned_vs2_p): New type predicate. (zvdota_e5m2_vs2_p): Ditto. (class vqwdota): New expander. (class vfwdota): Ditto. (class vfqwdota): Ditto. (BASE): Register vqwdota, vfwdota and vfqwdota. * config/riscv/riscv-vector-builtins-bases.h: Declare new bases. * config/riscv/riscv-vector-builtins-functions.def: New definitions. * config/riscv/riscv-vector-builtins-shapes.cc (supports_vectype_p): Allow zvdota shape. (struct zvdota_def): New shape with three type suffixes. (SHAPE): Register it. * config/riscv/riscv-vector-builtins-shapes.h: Declare it. * config/riscv/riscv-vector-builtins.cc (get_builtin_partition): Handle new Zvdota partitions. (register_builtin_types_on_null): Register float8e4m3/float8e5m2 scalar types. (get_zvdota_acc_type_index): New index. (get_float8_type_index): Ditto. (rvv_arg_type_info::get_function_type_index): Handle Zvdota accumulator and FP8 vector base types. (function_expander::use_zvdota_insn): New. (registered_function::overloaded_hash): Use ABI-mangled names. * config/riscv/riscv-vector-builtins.def: Add vfloat8e4m3 and vfloat8e5m2 types and Zvdota accumulator base types. * config/riscv/riscv-vector-builtins.h (RVV_NO_ALTFMT): New definition. (enum required_ext): Add Zvdota extensions. (enum rvv_builtin_partition): Ditto. (required_ext_to_isa_name): Handle them. (required_extensions_specified): Ditto. gcc/testsuite/ChangeLog: * gcc.target/riscv/rvv/base/zvdota-intrinsic.c: New test. * gcc.target/riscv/rvv/base/zvdota-no-rm-fail.c: New test. --- gcc/config/riscv/genrvv-type-indexer.cc | 114 ++++++ .../riscv/riscv-vector-builtins-bases.cc | 115 ++++++ .../riscv/riscv-vector-builtins-bases.h | 4 + .../riscv/riscv-vector-builtins-functions.def | 32 ++ .../riscv/riscv-vector-builtins-shapes.cc | 41 ++- .../riscv/riscv-vector-builtins-shapes.h | 1 + gcc/config/riscv/riscv-vector-builtins.cc | 344 +++++++++++++++++- gcc/config/riscv/riscv-vector-builtins.def | 36 +- gcc/config/riscv/riscv-vector-builtins.h | 29 ++ .../riscv/rvv/base/zvdota-intrinsic.c | 106 ++++++ .../riscv/rvv/base/zvdota-no-rm-fail.c | 25 ++ 11 files changed, 843 insertions(+), 4 deletions(-) create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-intrinsic.c create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-no-rm-fail.c diff --git a/gcc/config/riscv/genrvv-type-indexer.cc b/gcc/config/riscv/genrvv-type-indexer.cc index ddc0d9863fc..8ba0e0c42b5 100644 --- a/gcc/config/riscv/genrvv-type-indexer.cc +++ b/gcc/config/riscv/genrvv-type-indexer.cc @@ -133,6 +133,17 @@ bfloat16_type (int lmul_log2) return mode.str (); } +std::string +float8type (const char *format, int lmul_log2) +{ + if (!valid_type (8, lmul_log2, /*float_t*/ false)) + return "INVALID"; + + std::stringstream mode; + mode << "vfloat8" << format << to_lmul (lmul_log2) << "_t"; + return mode.str (); +} + std::string bfloat16_wide_type (int lmul_log2) { @@ -244,6 +255,90 @@ same_ratio_eew_bf16_type (unsigned sew, int lmul_log2) return bfloat16_type (elmul_log2); } +static void +emit_zvdota_type_indexes (FILE *fp, const char *acc = "INVALID", + const char *uacc = "INVALID", + const char *f32_acc = "INVALID", + const char *fp8e4m3 = "INVALID", + const char *fp8e5m2 = "INVALID") +{ + fprintf (fp, " /*ZVDOTA_ACC*/ %s,\n", acc); + fprintf (fp, " /*ZVDOTA_UACC*/ %s,\n", uacc); + fprintf (fp, " /*ZVDOTA_F32_ACC*/ %s,\n", f32_acc); + fprintf (fp, " /*FP8E4M3*/ %s,\n", fp8e4m3); + fprintf (fp, " /*FP8E5M2*/ %s,\n", fp8e5m2); +} + +static void +emit_float8_type_index (FILE *fp, const char *format, int lmul_log2) +{ + fprintf (fp, "DEF_RVV_TYPE_INDEX (\n"); + fprintf (fp, " /*VECTOR*/ %s,\n", + float8type (format, lmul_log2).c_str ()); + fprintf (fp, " /*MASK*/ %s,\n", maskmode (8, lmul_log2).c_str ()); + fprintf (fp, " /*SIGNED*/ %s,\n", + inttype (8, lmul_log2, /*unsigned_p*/ false).c_str ()); + fprintf (fp, " /*UNSIGNED*/ %s,\n", + inttype (8, lmul_log2, /*unsigned_p*/ true).c_str ()); + fprintf (fp, " /*SIGNED_EEW8_INDEX*/ %s,\n", + inttype (8, lmul_log2, /*unsigned_p*/ false).c_str ()); + for (unsigned eew : {8, 16, 32, 64}) + fprintf (fp, " /*EEW%d_INDEX*/ %s,\n", eew, + same_ratio_eew_type (8, lmul_log2, eew, + /*unsigned_p*/ true, false) + .c_str ()); + fprintf (fp, " /*SHIFT*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC*/ INVALID,\n"); + fprintf (fp, " /*QUAD_TRUNC*/ INVALID,\n"); + fprintf (fp, " /*QUAD_EMUL*/ INVALID,\n"); + fprintf (fp, " /*QUAD_EMUL_SIGNED*/ INVALID,\n"); + fprintf (fp, " /*QUAD_EMUL_UNSIGNED*/ INVALID,\n"); + fprintf (fp, " /*QUAD_FIX*/ INVALID,\n"); + fprintf (fp, " /*QUAD_FIX_SIGNED*/ INVALID,\n"); + fprintf (fp, " /*QUAD_FIX_UNSIGNED*/ INVALID,\n"); + fprintf (fp, " /*OCT_TRUNC*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_SCALAR*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_SIGNED*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_UNSIGNED*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_UNSIGNED_SCALAR*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_BFLOAT_SCALAR*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_BFLOAT*/ INVALID,\n"); + fprintf (fp, " /*DOUBLE_TRUNC_FLOAT*/ INVALID,\n"); + fprintf (fp, " /*FLOAT*/ INVALID,\n"); + fprintf (fp, " /*LMUL1*/ %s,\n", float8type (format, 0).c_str ()); + fprintf (fp, " /*WLMUL1*/ INVALID,\n"); + fprintf (fp, " /*QLMUL1*/ INVALID,\n"); + fprintf (fp, " /*QLMUL1_SIGNED*/ INVALID,\n"); + fprintf (fp, " /*QLMUL1_UNSIGNED*/ INVALID,\n"); + fprintf (fp, " /*XFQF*/ INVALID,\n"); + for (unsigned eew : {8, 16, 32, 64}) + fprintf (fp, " /*EEW%d_INTERPRET*/ INVALID,\n", eew); + + for (unsigned boolsize : BOOL_SIZE_LIST) + fprintf (fp, " /*BOOL%d_INTERPRET*/ INVALID,\n", boolsize); + + for (unsigned eew : EEW_SIZE_LIST) + fprintf (fp, " /*SIGNED_EEW%d_LMUL1_INTERPRET*/ INVALID,\n", eew); + + for (unsigned eew : EEW_SIZE_LIST) + fprintf (fp, " /*UNSIGNED_EEW%d_LMUL1_INTERPRET*/ INVALID,\n", eew); + + fprintf (fp, " /*X2*/ INVALID,\n"); + + for (unsigned lmul_log2_offset : {1, 2, 3, 4, 5, 6}) + { + unsigned multiple_of_lmul = 1 << lmul_log2_offset; + fprintf (fp, " /*X%d_VLMUL_EXT*/ INVALID,\n", multiple_of_lmul); + } + + emit_zvdota_type_indexes (fp, "INVALID", "INVALID", "vfloat32m1_t", + float8type ("e4m3", lmul_log2).c_str (), + float8type ("e5m2", lmul_log2).c_str ()); + fprintf (fp, " /*TUPLE_SUBPART*/ %s\n", + float8type (format, lmul_log2).c_str ()); + fprintf (fp, ")\n"); +} + int main (int argc, const char **argv) { @@ -315,6 +410,7 @@ main (int argc, const char **argv) unsigned multiple_of_lmul = 1 << lmul_log2_offset; fprintf (fp, " /*X%d_INTERPRET*/ INVALID,\n", multiple_of_lmul); } + emit_zvdota_type_indexes (fp); fprintf (fp, " /*TUPLE_SUBPART*/ INVALID\n"); fprintf (fp, ")\n"); } @@ -444,9 +540,24 @@ main (int argc, const char **argv) inttype (sew, lmul_log2 + lmul_log2_offset, unsigned_p) .c_str ()); } + if (nf == 1 && sew == 8) + emit_zvdota_type_indexes ( + fp, inttype (32, LMUL1_LOG2, /*unsigned_p*/ false).c_str (), + inttype (32, LMUL1_LOG2, /*unsigned_p*/ true).c_str ()); + else if (nf == 1 && sew == 16) + emit_zvdota_type_indexes ( + fp, inttype (64, LMUL1_LOG2, /*unsigned_p*/ false).c_str (), + inttype (64, LMUL1_LOG2, /*unsigned_p*/ true).c_str ()); + else + emit_zvdota_type_indexes (fp); fprintf (fp, " /*TUPLE_SUBPART*/ %s\n", inttype (sew, lmul_log2, 1, unsigned_p).c_str ()); fprintf (fp, ")\n"); + if (sew == 8 && nf == 1 && unsigned_p) + { + emit_float8_type_index (fp, "e4m3", lmul_log2); + emit_float8_type_index (fp, "e5m2", lmul_log2); + } } // Build for vbfloat16 for (int lmul_log2 : {-2, -1, 0, 1, 2, 3}) @@ -520,6 +631,8 @@ main (int argc, const char **argv) fprintf (fp, " /*X%d_VLMUL_EXT*/ %s,\n", multiple_of_lmul, bfloat16_type (lmul_log2 + lmul_log2_offset).c_str ()); } + emit_zvdota_type_indexes (fp, "INVALID", "INVALID", + nf == 1 ? "vfloat32m1_t" : "INVALID"); fprintf (fp, " /*TUPLE_SUBPART*/ %s\n", bfloat16_type (lmul_log2, 1U).c_str ()); fprintf (fp, ")\n"); @@ -609,6 +722,7 @@ main (int argc, const char **argv) fprintf (fp, " /*X%d_VLMUL_EXT*/ %s,\n", multiple_of_lmul, floattype (sew, lmul_log2 + lmul_log2_offset).c_str ()); } + emit_zvdota_type_indexes (fp); fprintf (fp, " /*TUPLE_SUBPART*/ %s\n", floattype (sew, lmul_log2, 1).c_str ()); fprintf (fp, ")\n"); diff --git a/gcc/config/riscv/riscv-vector-builtins-bases.cc b/gcc/config/riscv/riscv-vector-builtins-bases.cc index 0ae9b34fa2b..a41d0753422 100644 --- a/gcc/config/riscv/riscv-vector-builtins-bases.cc +++ b/gcc/config/riscv/riscv-vector-builtins-bases.cc @@ -2462,6 +2462,113 @@ public: } }; +static unsigned +zvdota_vs2_argno (const function_expander &e) +{ + return e.pred == PRED_TYPE_m || e.pred == PRED_TYPE_tum ? 2 : 1; +} + +static unsigned +zvdota_int_altfmt (const function_expander &e) +{ + return e.op_info->args[2].base_type == RVV_BASE_unsigned_vector + ? ALTFMT_NONE + : ALTFMT_ALT; +} + +static unsigned +zvdota_fp8_altfmt (const function_expander &e) +{ + return e.op_info->args[2].base_type == RVV_BASE_fp8e5m2_vector + ? ALTFMT_ALT + : ALTFMT_NONE; +} + +static bool +zvdota_unsigned_vs2_p (vector_type_index type_idx) +{ + switch (type_idx) + { + case VECTOR_TYPE_vuint8mf8_t: + case VECTOR_TYPE_vuint8mf4_t: + case VECTOR_TYPE_vuint8mf2_t: + case VECTOR_TYPE_vuint8m1_t: + case VECTOR_TYPE_vuint8m2_t: + case VECTOR_TYPE_vuint8m4_t: + case VECTOR_TYPE_vuint8m8_t: + case VECTOR_TYPE_vuint16mf4_t: + case VECTOR_TYPE_vuint16mf2_t: + case VECTOR_TYPE_vuint16m1_t: + case VECTOR_TYPE_vuint16m2_t: + case VECTOR_TYPE_vuint16m4_t: + case VECTOR_TYPE_vuint16m8_t: + return true; + default: + return false; + } +} + +static bool +zvdota_e5m2_vs2_p (vector_type_index type_idx) +{ + switch (type_idx) + { + case VECTOR_TYPE_vfloat8e5m2mf8_t: + case VECTOR_TYPE_vfloat8e5m2mf4_t: + case VECTOR_TYPE_vfloat8e5m2mf2_t: + case VECTOR_TYPE_vfloat8e5m2m1_t: + case VECTOR_TYPE_vfloat8e5m2m2_t: + case VECTOR_TYPE_vfloat8e5m2m4_t: + case VECTOR_TYPE_vfloat8e5m2m8_t: + return true; + default: + return false; + } +} + +class vqwdota : public function_base +{ +public: + bool has_merge_operand_p () const override { return false; } + + rtx expand (function_expander &e) const override + { + machine_mode mode = e.arg_mode (zvdota_vs2_argno (e)); + int unspec = zvdota_unsigned_vs2_p (e.type.index) ? UNSPEC_VQWDOTAU + : UNSPEC_VQWDOTAS; + return e.use_zvdota_insn ( + code_for_pred_vv_zvdota (unspec, mode), zvdota_int_altfmt (e)); + } +}; + +class vfwdota : public function_base +{ +public: + bool has_merge_operand_p () const override { return false; } + + rtx expand (function_expander &e) const override + { + return e.use_zvdota_insn ( + code_for_pred_vfwdota_vv_zvdota (e.arg_mode (zvdota_vs2_argno (e))), + ALTFMT_ALT); + } +}; + +class vfqwdota : public function_base +{ +public: + bool has_merge_operand_p () const override { return false; } + + rtx expand (function_expander &e) const override + { + machine_mode mode = e.arg_mode (zvdota_vs2_argno (e)); + int unspec = zvdota_e5m2_vs2_p (e.type.index) ? UNSPEC_VFQWDOTA_ALT + : UNSPEC_VFQWDOTA; + return e.use_zvdota_insn ( + code_for_pred_vv_zvdota (unspec, mode), zvdota_fp8_altfmt (e)); + } +}; + static CONSTEXPR const vsetvl<false> vsetvl_obj; static CONSTEXPR const vsetvl<true> vsetvlmax_obj; static CONSTEXPR const loadstore<false, LST_UNIT_STRIDE, false> vle_obj; @@ -2786,6 +2893,10 @@ static CONSTEXPR const vfwcvtbf16_f vfwcvtbf16_f_obj; /* Zvfbfwma; */ static CONSTEXPR const vfwmaccbf16<NO_FRM> vfwmaccbf16_obj; static CONSTEXPR const vfwmaccbf16<HAS_FRM> vfwmaccbf16_frm_obj; +/* Zvdota. */ +static CONSTEXPR const vqwdota vqwdota_obj; +static CONSTEXPR const vfwdota vfwdota_obj; +static CONSTEXPR const vfqwdota vfqwdota_obj; /* Declare the function base NAME, pointing it to an instance of class <NAME>_obj. */ @@ -3114,4 +3225,8 @@ BASE (vfwcvtbf16_f) /* Zvfbfwma */ BASE (vfwmaccbf16) BASE (vfwmaccbf16_frm) +/* Zvdota */ +BASE (vqwdota) +BASE (vfwdota) +BASE (vfqwdota) } // end namespace riscv_vector diff --git a/gcc/config/riscv/riscv-vector-builtins-bases.h b/gcc/config/riscv/riscv-vector-builtins-bases.h index 9261d353e22..f561c632a9c 100644 --- a/gcc/config/riscv/riscv-vector-builtins-bases.h +++ b/gcc/config/riscv/riscv-vector-builtins-bases.h @@ -352,6 +352,10 @@ extern const function_base *const vfwcvtbf16_f; /* Zvfbfwma */ extern const function_base *const vfwmaccbf16; extern const function_base *const vfwmaccbf16_frm; +/* Zvdota */ +extern const function_base *const vqwdota; +extern const function_base *const vfwdota; +extern const function_base *const vfqwdota; } } // end namespace riscv_vector diff --git a/gcc/config/riscv/riscv-vector-builtins-functions.def b/gcc/config/riscv/riscv-vector-builtins-functions.def index f64d42d1823..0d037937cad 100644 --- a/gcc/config/riscv/riscv-vector-builtins-functions.def +++ b/gcc/config/riscv/riscv-vector-builtins-functions.def @@ -431,6 +431,38 @@ DEF_RVV_FUNCTION (vfwmsac_frm, alu_frm, full_preds, f_wwfv_ops) DEF_RVV_FUNCTION (vfwnmsac_frm, alu_frm, full_preds, f_wwvv_ops) DEF_RVV_FUNCTION (vfwnmsac_frm, alu_frm, full_preds, f_wwfv_ops) +#undef REQUIRED_EXTENSIONS +#define REQUIRED_EXTENSIONS ZVQWDOTA8I_EXT +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota8i_ss_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota8i_su_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota8i_us_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota8i_uu_vv_ops) + +#undef REQUIRED_EXTENSIONS +#define REQUIRED_EXTENSIONS ZVQWDOTA16I_EXT +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota16i_ss_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota16i_su_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota16i_us_vv_ops) +DEF_RVV_FUNCTION (vqwdota, zvdota, no_mu_preds, zvqwdota16i_uu_vv_ops) + +#undef REQUIRED_EXTENSIONS +#define REQUIRED_EXTENSIONS ZVFWDOTA16BF_EXT +DEF_RVV_FUNCTION (vfwdota, zvdota, no_mu_preds, zvfwdota16bf_vv_ops) + +#undef REQUIRED_EXTENSIONS +#define REQUIRED_EXTENSIONS ZVFQWDOTA8F_EXT +DEF_RVV_FUNCTION (vfqwdota, zvdota, no_mu_preds, + zvfqwdota8f_e4m3_e4m3_vv_ops) +DEF_RVV_FUNCTION (vfqwdota, zvdota, no_mu_preds, + zvfqwdota8f_e4m3_e5m2_vv_ops) +DEF_RVV_FUNCTION (vfqwdota, zvdota, no_mu_preds, + zvfqwdota8f_e5m2_e4m3_vv_ops) +DEF_RVV_FUNCTION (vfqwdota, zvdota, no_mu_preds, + zvfqwdota8f_e5m2_e5m2_vv_ops) + +#undef REQUIRED_EXTENSIONS +#define REQUIRED_EXTENSIONS VECTOR_EXT + // 13.8. Vector Floating-Point Square-Root Instruction DEF_RVV_FUNCTION (vfsqrt, alu, full_preds, f_v_ops) diff --git a/gcc/config/riscv/riscv-vector-builtins-shapes.cc b/gcc/config/riscv/riscv-vector-builtins-shapes.cc index a113bae56e1..dce48bef42e 100644 --- a/gcc/config/riscv/riscv-vector-builtins-shapes.cc +++ b/gcc/config/riscv/riscv-vector-builtins-shapes.cc @@ -97,7 +97,8 @@ supports_vectype_p (const function_group_info &group, unsigned int vec_type_idx) || *group.shape == shapes::fault_load || *group.shape == shapes::seg_loadstore || *group.shape == shapes::seg_indexed_loadstore - || *group.shape == shapes::seg_fault_load) + || *group.shape == shapes::seg_fault_load + || *group.shape == shapes::zvdota) return true; return false; } @@ -424,6 +425,43 @@ struct alu_def : public build_base } }; +/* zvdota_def class. */ +struct zvdota_def : public build_base +{ + char *get_name (function_builder &b, const function_instance &instance, + bool overloaded_p) const override + { + /* Return nullptr if it can not be overloaded. */ + if (overloaded_p && !instance.base->can_be_overloaded_p (instance.pred)) + return nullptr; + + b.append_base_name (instance.base_name); + + if (!overloaded_p) + { + vector_type_index vs2_type_idx + = instance.op_info->args[1].get_function_type_index ( + instance.type.index); + vector_type_index vs1_type_idx + = instance.op_info->args[2].get_function_type_index ( + instance.type.index); + vector_type_index ret_type_idx + = instance.op_info->ret.get_function_type_index ( + instance.type.index); + + b.append_name (operand_suffixes[instance.op_info->op]); + b.append_name (type_suffixes[vs2_type_idx].vector); + b.append_name (type_suffixes[vs1_type_idx].vector); + b.append_name (type_suffixes[ret_type_idx].vector); + } + + if (overloaded_p && instance.pred == PRED_TYPE_m) + return b.finish_name (); + b.append_name (predication_suffixes[instance.pred]); + return b.finish_name (); + } +}; + /* The base class for frm build. */ struct build_frm_base : public build_base { @@ -1427,6 +1465,7 @@ SHAPE(seg_fault_load, seg_fault_load) SHAPE(crypto_vv, crypto_vv) SHAPE(crypto_vi, crypto_vi) SHAPE(crypto_vv_no_op_type, crypto_vv_no_op_type) +SHAPE(zvdota, zvdota) SHAPE (sf_vqmacc, sf_vqmacc) SHAPE (sf_vfnrclip, sf_vfnrclip) SHAPE(sf_vcix_se, sf_vcix_se) diff --git a/gcc/config/riscv/riscv-vector-builtins-shapes.h b/gcc/config/riscv/riscv-vector-builtins-shapes.h index d9502d4c5b4..db3c645fb37 100644 --- a/gcc/config/riscv/riscv-vector-builtins-shapes.h +++ b/gcc/config/riscv/riscv-vector-builtins-shapes.h @@ -32,6 +32,7 @@ extern const function_shape *const th_loadstore_width; extern const function_shape *const th_indexed_loadstore_width; extern const function_shape *const alu; extern const function_shape *const alu_frm; +extern const function_shape *const zvdota; extern const function_shape *const widen_alu; extern const function_shape *const widen_alu_frm; extern const function_shape *const no_mask_policy; diff --git a/gcc/config/riscv/riscv-vector-builtins.cc b/gcc/config/riscv/riscv-vector-builtins.cc index 0fcc60b5942..4894c0183b3 100644 --- a/gcc/config/riscv/riscv-vector-builtins.cc +++ b/gcc/config/riscv/riscv-vector-builtins.cc @@ -590,6 +590,74 @@ static const rvv_type_info f16_ops[] = { #include "riscv-vector-builtins-types.def" {NUM_VECTOR_TYPES, 0}}; +static const rvv_type_info zvqwdota8i_i_ops[] = { + {VECTOR_TYPE_vint8mf8_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vint8mf4_t, 0}, + {VECTOR_TYPE_vint8mf2_t, 0}, + {VECTOR_TYPE_vint8m1_t, 0}, + {VECTOR_TYPE_vint8m2_t, 0}, + {VECTOR_TYPE_vint8m4_t, 0}, + {VECTOR_TYPE_vint8m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvqwdota8i_u_ops[] = { + {VECTOR_TYPE_vuint8mf8_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vuint8mf4_t, 0}, + {VECTOR_TYPE_vuint8mf2_t, 0}, + {VECTOR_TYPE_vuint8m1_t, 0}, + {VECTOR_TYPE_vuint8m2_t, 0}, + {VECTOR_TYPE_vuint8m4_t, 0}, + {VECTOR_TYPE_vuint8m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvqwdota16i_i_ops[] = { + {VECTOR_TYPE_vint16mf4_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vint16mf2_t, 0}, + {VECTOR_TYPE_vint16m1_t, 0}, + {VECTOR_TYPE_vint16m2_t, 0}, + {VECTOR_TYPE_vint16m4_t, 0}, + {VECTOR_TYPE_vint16m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvqwdota16i_u_ops[] = { + {VECTOR_TYPE_vuint16mf4_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vuint16mf2_t, 0}, + {VECTOR_TYPE_vuint16m1_t, 0}, + {VECTOR_TYPE_vuint16m2_t, 0}, + {VECTOR_TYPE_vuint16m4_t, 0}, + {VECTOR_TYPE_vuint16m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvfwdota16bf_ops[] = { + {VECTOR_TYPE_vbfloat16mf4_t, + RVV_REQUIRE_ELEN_BF_16 | RVV_REQUIRE_MIN_VLEN_64}, + {VECTOR_TYPE_vbfloat16mf2_t, RVV_REQUIRE_ELEN_BF_16}, + {VECTOR_TYPE_vbfloat16m1_t, RVV_REQUIRE_ELEN_BF_16}, + {VECTOR_TYPE_vbfloat16m2_t, RVV_REQUIRE_ELEN_BF_16}, + {VECTOR_TYPE_vbfloat16m4_t, RVV_REQUIRE_ELEN_BF_16}, + {VECTOR_TYPE_vbfloat16m8_t, RVV_REQUIRE_ELEN_BF_16}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvfqwdota8f_e4m3_ops[] = { + {VECTOR_TYPE_vfloat8e4m3mf8_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vfloat8e4m3mf4_t, 0}, + {VECTOR_TYPE_vfloat8e4m3mf2_t, 0}, + {VECTOR_TYPE_vfloat8e4m3m1_t, 0}, + {VECTOR_TYPE_vfloat8e4m3m2_t, 0}, + {VECTOR_TYPE_vfloat8e4m3m4_t, 0}, + {VECTOR_TYPE_vfloat8e4m3m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + +static const rvv_type_info zvfqwdota8f_e5m2_ops[] = { + {VECTOR_TYPE_vfloat8e5m2mf8_t, RVV_REQUIRE_ELEN_64}, + {VECTOR_TYPE_vfloat8e5m2mf4_t, 0}, + {VECTOR_TYPE_vfloat8e5m2mf2_t, 0}, + {VECTOR_TYPE_vfloat8e5m2m1_t, 0}, + {VECTOR_TYPE_vfloat8e5m2m2_t, 0}, + {VECTOR_TYPE_vfloat8e5m2m4_t, 0}, + {VECTOR_TYPE_vfloat8e5m2m8_t, 0}, + {NUM_VECTOR_TYPES, 0}}; + static CONSTEXPR const rvv_arg_type_info rvv_arg_type_info_end = rvv_arg_type_info (NUM_BASE_TYPES); @@ -689,6 +757,41 @@ static CONSTEXPR const rvv_arg_type_info vv_args[] = {rvv_arg_type_info (RVV_BASE_vector), rvv_arg_type_info (RVV_BASE_vector), rvv_arg_type_info_end}; +static CONSTEXPR const rvv_arg_type_info zvdota_ss_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_signed_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_su_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_unsigned_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_us_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_signed_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_uu_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_uacc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_unsigned_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_f_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_fp8_e4m3_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_fp8e4m3_vector), rvv_arg_type_info_end}; + +static CONSTEXPR const rvv_arg_type_info zvdota_fp8_e5m2_args[] + = {rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), + rvv_arg_type_info (RVV_BASE_vector), + rvv_arg_type_info (RVV_BASE_fp8e5m2_vector), rvv_arg_type_info_end}; + /* A list of args for vector_type func (vector_type, vector_type, vector_type) * function. */ static CONSTEXPR const rvv_arg_type_info vvv_args[] @@ -1530,6 +1633,84 @@ static CONSTEXPR const rvv_op_info f_vvvv_ops rvv_arg_type_info (RVV_BASE_vector), /* Return type */ vvv_args /* Args */}; +static CONSTEXPR const rvv_op_info zvqwdota8i_ss_vv_ops + = {zvqwdota8i_i_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_ss_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota8i_su_vv_ops + = {zvqwdota8i_i_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_su_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota8i_us_vv_ops + = {zvqwdota8i_u_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_us_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota8i_uu_vv_ops + = {zvqwdota8i_u_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_uacc), /* Return type */ + zvdota_uu_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota16i_ss_vv_ops + = {zvqwdota16i_i_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_ss_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota16i_su_vv_ops + = {zvqwdota16i_i_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_su_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota16i_us_vv_ops + = {zvqwdota16i_u_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_acc), /* Return type */ + zvdota_us_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvqwdota16i_uu_vv_ops + = {zvqwdota16i_u_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_uacc), /* Return type */ + zvdota_uu_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvfwdota16bf_vv_ops + = {zvfwdota16bf_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), /* Return type */ + zvdota_f_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvfqwdota8f_e4m3_e4m3_vv_ops + = {zvfqwdota8f_e4m3_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), /* Return type */ + zvdota_fp8_e4m3_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvfqwdota8f_e4m3_e5m2_vv_ops + = {zvfqwdota8f_e4m3_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), /* Return type */ + zvdota_fp8_e5m2_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvfqwdota8f_e5m2_e4m3_vv_ops + = {zvfqwdota8f_e5m2_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), /* Return type */ + zvdota_fp8_e4m3_args /* Args */}; + +static CONSTEXPR const rvv_op_info zvfqwdota8f_e5m2_e5m2_vv_ops + = {zvfqwdota8f_e5m2_ops, /* Types */ + OP_TYPE_vv, /* Suffix */ + rvv_arg_type_info (RVV_BASE_zvdota_f32_acc), /* Return type */ + zvdota_fp8_e5m2_args /* Args */}; + /* A static operand information for vector_type func (vector_type, scalar_type, * vector_type) function registration. */ static CONSTEXPR const rvv_op_info f_vvfv_ops @@ -3458,7 +3639,8 @@ static CONSTEXPR const function_type_info function_types[] = { SIGNED_EEW64_LMUL1_INTERPRET, UNSIGNED_EEW8_LMUL1_INTERPRET, \ UNSIGNED_EEW16_LMUL1_INTERPRET, UNSIGNED_EEW32_LMUL1_INTERPRET, \ UNSIGNED_EEW64_LMUL1_INTERPRET, X2, X2_VLMUL_EXT, X4_VLMUL_EXT, X8_VLMUL_EXT,\ - X16_VLMUL_EXT, X32_VLMUL_EXT, X64_VLMUL_EXT, TUPLE_SUBPART) \ + X16_VLMUL_EXT, X32_VLMUL_EXT, X64_VLMUL_EXT, ZVDOTA_ACC, ZVDOTA_UACC, \ + ZVDOTA_F32_ACC, FP8E4M3, FP8E5M2, TUPLE_SUBPART) \ { \ VECTOR_TYPE_##VECTOR, \ VECTOR_TYPE_INVALID, \ @@ -3530,7 +3712,11 @@ static CONSTEXPR const function_type_info function_types[] = { VECTOR_TYPE_##X16_VLMUL_EXT, \ VECTOR_TYPE_##X32_VLMUL_EXT, \ VECTOR_TYPE_##X64_VLMUL_EXT, \ - VECTOR_TYPE_INVALID, \ + VECTOR_TYPE_##ZVDOTA_ACC, \ + VECTOR_TYPE_##ZVDOTA_UACC, \ + VECTOR_TYPE_##ZVDOTA_F32_ACC, \ + VECTOR_TYPE_##FP8E4M3, \ + VECTOR_TYPE_##FP8E5M2, \ VECTOR_TYPE_##TUPLE_SUBPART, \ }, #include "riscv-vector-builtins.def" @@ -3656,6 +3842,14 @@ get_builtin_partition (required_ext ext, const function_instance &instance) return RVV_PARTITION_ZVFBFWMA; case ZVFOFP8MIN_EXT: return RVV_PARTITION_ZVFOFP8MIN; + case ZVQWDOTA8I_EXT: + return RVV_PARTITION_ZVQWDOTA8I; + case ZVQWDOTA16I_EXT: + return RVV_PARTITION_ZVQWDOTA16I; + case ZVFWDOTA16BF_EXT: + return RVV_PARTITION_ZVFWDOTA16BF; + case ZVFQWDOTA8F_EXT: + return RVV_PARTITION_ZVFQWDOTA8F; case XSFVQMACCQOQ_EXT: return RVV_PARTITION_XSFVQMACCQOQ; case XSFVQMACCDOD_EXT: @@ -4000,6 +4194,8 @@ register_builtin_types_on_null () tree int16_type_node = get_typenode_from_name (INT16_TYPE); tree int32_type_node = get_typenode_from_name (INT32_TYPE); tree int64_type_node = get_typenode_from_name (INT64_TYPE); + tree float8e4m3_type_node = build_distinct_type_copy (uint8_type_node); + tree float8e5m2_type_node = build_distinct_type_copy (uint8_type_node); machine_mode mode; #define DEF_RVV_TYPE(NAME, NCHARS, ABI_NAME, SCALAR_TYPE, VECTOR_MODE, \ @@ -4216,9 +4412,113 @@ rvv_arg_type_info::get_scalar_float_type (vector_type_index type_idx) const return NULL_TREE; } +static vector_type_index +get_zvdota_acc_type_index (vector_type_index type_idx, bool unsigned_p) +{ + switch (type_idx) + { + case VECTOR_TYPE_vint8mf8_t: + case VECTOR_TYPE_vuint8mf8_t: + case VECTOR_TYPE_vint8mf4_t: + case VECTOR_TYPE_vuint8mf4_t: + case VECTOR_TYPE_vint8mf2_t: + case VECTOR_TYPE_vuint8mf2_t: + case VECTOR_TYPE_vint8m1_t: + case VECTOR_TYPE_vuint8m1_t: + case VECTOR_TYPE_vint8m2_t: + case VECTOR_TYPE_vuint8m2_t: + case VECTOR_TYPE_vint8m4_t: + case VECTOR_TYPE_vuint8m4_t: + case VECTOR_TYPE_vint8m8_t: + case VECTOR_TYPE_vuint8m8_t: + return unsigned_p ? VECTOR_TYPE_vuint32m1_t : VECTOR_TYPE_vint32m1_t; + + case VECTOR_TYPE_vint16mf4_t: + case VECTOR_TYPE_vuint16mf4_t: + case VECTOR_TYPE_vint16mf2_t: + case VECTOR_TYPE_vuint16mf2_t: + case VECTOR_TYPE_vint16m1_t: + case VECTOR_TYPE_vuint16m1_t: + case VECTOR_TYPE_vint16m2_t: + case VECTOR_TYPE_vuint16m2_t: + case VECTOR_TYPE_vint16m4_t: + case VECTOR_TYPE_vuint16m4_t: + case VECTOR_TYPE_vint16m8_t: + case VECTOR_TYPE_vuint16m8_t: + return unsigned_p ? VECTOR_TYPE_vuint64m1_t : VECTOR_TYPE_vint64m1_t; + + default: + return VECTOR_TYPE_INVALID; + } +} + +static vector_type_index +get_float8_type_index (vector_type_index type_idx, bool e5m2_p) +{ + switch (type_idx) + { + case VECTOR_TYPE_vfloat8e4m3mf8_t: + case VECTOR_TYPE_vfloat8e5m2mf8_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2mf8_t + : VECTOR_TYPE_vfloat8e4m3mf8_t; + case VECTOR_TYPE_vfloat8e4m3mf4_t: + case VECTOR_TYPE_vfloat8e5m2mf4_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2mf4_t + : VECTOR_TYPE_vfloat8e4m3mf4_t; + case VECTOR_TYPE_vfloat8e4m3mf2_t: + case VECTOR_TYPE_vfloat8e5m2mf2_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2mf2_t + : VECTOR_TYPE_vfloat8e4m3mf2_t; + case VECTOR_TYPE_vfloat8e4m3m1_t: + case VECTOR_TYPE_vfloat8e5m2m1_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2m1_t + : VECTOR_TYPE_vfloat8e4m3m1_t; + case VECTOR_TYPE_vfloat8e4m3m2_t: + case VECTOR_TYPE_vfloat8e5m2m2_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2m2_t + : VECTOR_TYPE_vfloat8e4m3m2_t; + case VECTOR_TYPE_vfloat8e4m3m4_t: + case VECTOR_TYPE_vfloat8e5m2m4_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2m4_t + : VECTOR_TYPE_vfloat8e4m3m4_t; + case VECTOR_TYPE_vfloat8e4m3m8_t: + case VECTOR_TYPE_vfloat8e5m2m8_t: + return e5m2_p ? VECTOR_TYPE_vfloat8e5m2m8_t + : VECTOR_TYPE_vfloat8e4m3m8_t; + default: + return VECTOR_TYPE_INVALID; + } +} + vector_type_index rvv_arg_type_info::get_function_type_index (vector_type_index type_idx) const { + vector_type_index function_type_idx = VECTOR_TYPE_INVALID; + switch (base_type) + { + case RVV_BASE_zvdota_acc: + function_type_idx = get_zvdota_acc_type_index (type_idx, false); + break; + case RVV_BASE_zvdota_uacc: + function_type_idx = get_zvdota_acc_type_index (type_idx, true); + break; + case RVV_BASE_zvdota_f32_acc: + function_type_idx = VECTOR_TYPE_vfloat32m1_t; + break; + case RVV_BASE_fp8e4m3_vector: + function_type_idx = get_float8_type_index (type_idx, false); + break; + case RVV_BASE_fp8e5m2_vector: + function_type_idx = get_float8_type_index (type_idx, true); + break; + default: + break; + } + + if (function_type_idx != VECTOR_TYPE_INVALID) + return builtin_types[function_type_idx].vector ? function_type_idx + : NUM_VECTOR_TYPES; + tree type = builtin_types[function_types[type_idx].type_indexes[base_type]].vector; return type ? function_types[type_idx].type_indexes[base_type] @@ -5090,6 +5390,37 @@ function_expander::use_ternop_insn (bool vd_accum_p, insn_code icode) return generate_insn (icode); } +/* Implement Zvdota destructive-accumulator calls. The machine patterns use + the source vector mode to control vsetvl and tie the accumulator input to + the output register. */ +rtx +function_expander::use_zvdota_insn (insn_code icode, unsigned altfmt) +{ + int arg_offset = 0; + + if (use_real_mask_p (pred)) + add_input_operand (arg_offset++); + else + add_all_one_mask_operand (mask_mode ()); + + add_input_operand (arg_offset++); + add_input_operand (arg_offset++); + add_input_operand (arg_offset++); + + for (int argno = arg_offset; argno < call_expr_nargs (exp); argno++) + add_input_operand (argno); + + add_input_operand (Pmode, get_tail_policy_for_pred (pred)); + add_input_operand (Pmode, get_mask_policy_for_pred (pred)); + add_input_operand (Pmode, get_avl_type_rtx (avl_type::NONVLMAX)); + + if (altfmt != RVV_NO_ALTFMT + && opno < insn_data[icode].n_generator_args) + add_input_operand (Pmode, gen_int_mode (altfmt, Pmode)); + + return generate_insn (icode); +} + /* Implement the call using instruction ICODE, with a 1:1 mapping between arguments and input operands. There are operands that cannot be broadcast using v[f]mv. In that case @@ -5363,6 +5694,15 @@ registered_function::overloaded_hash () const : TYPE_UNSIGNED (type); mode_p = POINTER_TYPE_P (type) ? TYPE_MODE (TREE_TYPE (type)) : TYPE_MODE (type); + if (tree rvv_type = POINTER_TYPE_P (type) ? TREE_TYPE (type) : type) + if (tree attr = lookup_vector_type_attribute (rvv_type)) + if (tree id = TREE_VALUE (chain_index (0, TREE_VALUE (attr)))) + { + const char *mangled_name = IDENTIFIER_POINTER (id); + h.add (mangled_name, strlen (mangled_name)); + continue; + } + if (POINTER_TYPE_P (type) || lookup_vector_type_attribute (type)) { h.add_int (unsigned_p); diff --git a/gcc/config/riscv/riscv-vector-builtins.def b/gcc/config/riscv/riscv-vector-builtins.def index e4e23ce32b2..eeade4e8e66 100644 --- a/gcc/config/riscv/riscv-vector-builtins.def +++ b/gcc/config/riscv/riscv-vector-builtins.def @@ -83,7 +83,8 @@ along with GCC; see the file COPYING3. If not see SIGNED_EEW64_LMUL1_INTERPRET, UNSIGNED_EEW8_LMUL1_INTERPRET, \ UNSIGNED_EEW16_LMUL1_INTERPRET, UNSIGNED_EEW32_LMUL1_INTERPRET, \ UNSIGNED_EEW64_LMUL1_INTERPRET, X2, X2_VLMUL_EXT, X4_VLMUL_EXT, X8_VLMUL_EXT,\ - X16_VLMUL_EXT, X32_VLMUL_EXT, X64_VLMUL_EXT, TUPLE_SUBPART) + X16_VLMUL_EXT, X32_VLMUL_EXT, X64_VLMUL_EXT, ZVDOTA_ACC, ZVDOTA_UACC, \ + ZVDOTA_F32_ACC, FP8E4M3, FP8E5M2, TUPLE_SUBPART) #endif /* Define RVV_VXRM rounding mode enum for fixed-point intrinsics. */ @@ -126,6 +127,10 @@ DEF_RVV_TYPE (vint8mf8_t, 15, __rvv_int8mf8_t, int8, RVVMF8QI, _i8mf8, _i8, _e8mf8) DEF_RVV_TYPE (vuint8mf8_t, 16, __rvv_uint8mf8_t, uint8, RVVMF8QI, _u8mf8, _u8, _e8mf8) +DEF_RVV_TYPE (vfloat8e4m3mf8_t, 21, __rvv_float8e4m3mf8_t, float8e4m3, + RVVMF8QI, _f8e4m3mf8, _f8e4m3, _e8mf8) +DEF_RVV_TYPE (vfloat8e5m2mf8_t, 21, __rvv_float8e5m2mf8_t, float8e5m2, + RVVMF8QI, _f8e5m2mf8, _f8e5m2, _e8mf8) /* Define tuple types for SEW = 8, LMUL = MF8. */ DEF_RVV_TUPLE_TYPE (vint8mf8x2_t, 17, __rvv_int8mf8x2_t, vint8mf8_t, int8, 2, _i8mf8x2) DEF_RVV_TUPLE_TYPE (vuint8mf8x2_t, 18, __rvv_uint8mf8x2_t, vuint8mf8_t, uint8, 2, _u8mf8x2) @@ -147,6 +152,10 @@ DEF_RVV_TYPE (vint8mf4_t, 15, __rvv_int8mf4_t, int8, RVVMF4QI, _i8mf4, _i8, _e8mf4) DEF_RVV_TYPE (vuint8mf4_t, 16, __rvv_uint8mf4_t, uint8, RVVMF4QI, _u8mf4, _u8, _e8mf4) +DEF_RVV_TYPE (vfloat8e4m3mf4_t, 21, __rvv_float8e4m3mf4_t, float8e4m3, + RVVMF4QI, _f8e4m3mf4, _f8e4m3, _e8mf4) +DEF_RVV_TYPE (vfloat8e5m2mf4_t, 21, __rvv_float8e5m2mf4_t, float8e5m2, + RVVMF4QI, _f8e5m2mf4, _f8e5m2, _e8mf4) /* Define tuple types for SEW = 8, LMUL = MF4. */ DEF_RVV_TUPLE_TYPE (vint8mf4x2_t, 17, __rvv_int8mf4x2_t, vint8mf4_t, int8, 2, _i8mf4x2) DEF_RVV_TUPLE_TYPE (vuint8mf4x2_t, 18, __rvv_uint8mf4x2_t, vuint8mf4_t, uint8, 2, _u8mf4x2) @@ -168,6 +177,10 @@ DEF_RVV_TYPE (vint8mf2_t, 15, __rvv_int8mf2_t, int8, RVVMF2QI, _i8mf2, _i8, _e8mf2) DEF_RVV_TYPE (vuint8mf2_t, 16, __rvv_uint8mf2_t, uint8, RVVMF2QI, _u8mf2, _u8, _e8mf2) +DEF_RVV_TYPE (vfloat8e4m3mf2_t, 21, __rvv_float8e4m3mf2_t, float8e4m3, + RVVMF2QI, _f8e4m3mf2, _f8e4m3, _e8mf2) +DEF_RVV_TYPE (vfloat8e5m2mf2_t, 21, __rvv_float8e5m2mf2_t, float8e5m2, + RVVMF2QI, _f8e5m2mf2, _f8e5m2, _e8mf2) /* Define tuple types for SEW = 8, LMUL = MF2. */ DEF_RVV_TUPLE_TYPE (vint8mf2x2_t, 17, __rvv_int8mf2x2_t, vint8mf2_t, int8, 2, _i8mf2x2) DEF_RVV_TUPLE_TYPE (vuint8mf2x2_t, 18, __rvv_uint8mf2x2_t, vuint8mf2_t, uint8, 2, _u8mf2x2) @@ -188,6 +201,10 @@ DEF_RVV_TUPLE_TYPE (vuint8mf2x8_t, 18, __rvv_uint8mf2x8_t, vuint8mf2_t, uint8, 8 DEF_RVV_TYPE (vint8m1_t, 14, __rvv_int8m1_t, int8, RVVM1QI, _i8m1, _i8, _e8m1) DEF_RVV_TYPE (vuint8m1_t, 15, __rvv_uint8m1_t, uint8, RVVM1QI, _u8m1, _u8, _e8m1) +DEF_RVV_TYPE (vfloat8e4m3m1_t, 20, __rvv_float8e4m3m1_t, float8e4m3, + RVVM1QI, _f8e4m3m1, _f8e4m3, _e8m1) +DEF_RVV_TYPE (vfloat8e5m2m1_t, 20, __rvv_float8e5m2m1_t, float8e5m2, + RVVM1QI, _f8e5m2m1, _f8e5m2, _e8m1) /* Define tuple types for SEW = 8, LMUL = M1. */ DEF_RVV_TUPLE_TYPE (vint8m1x2_t, 16, __rvv_int8m1x2_t, vint8m1_t, int8, 2, _i8m1x2) DEF_RVV_TUPLE_TYPE (vuint8m1x2_t, 17, __rvv_uint8m1x2_t, vuint8m1_t, uint8, 2, _u8m1x2) @@ -208,6 +225,10 @@ DEF_RVV_TUPLE_TYPE (vuint8m1x8_t, 17, __rvv_uint8m1x8_t, vuint8m1_t, uint8, 8, _ DEF_RVV_TYPE (vint8m2_t, 14, __rvv_int8m2_t, int8, RVVM2QI, _i8m2, _i8, _e8m2) DEF_RVV_TYPE (vuint8m2_t, 15, __rvv_uint8m2_t, uint8, RVVM2QI, _u8m2, _u8, _e8m2) +DEF_RVV_TYPE (vfloat8e4m3m2_t, 20, __rvv_float8e4m3m2_t, float8e4m3, + RVVM2QI, _f8e4m3m2, _f8e4m3, _e8m2) +DEF_RVV_TYPE (vfloat8e5m2m2_t, 20, __rvv_float8e5m2m2_t, float8e5m2, + RVVM2QI, _f8e5m2m2, _f8e5m2, _e8m2) /* Define tuple types for SEW = 8, LMUL = M2. */ DEF_RVV_TUPLE_TYPE (vint8m2x2_t, 16, __rvv_int8m2x2_t, vint8m2_t, int8, 2, _i8m2x2) DEF_RVV_TUPLE_TYPE (vuint8m2x2_t, 17, __rvv_uint8m2x2_t, vuint8m2_t, uint8, 2, _u8m2x2) @@ -220,6 +241,10 @@ DEF_RVV_TUPLE_TYPE (vuint8m2x4_t, 17, __rvv_uint8m2x4_t, vuint8m2_t, uint8, 4, _ DEF_RVV_TYPE (vint8m4_t, 14, __rvv_int8m4_t, int8, RVVM4QI, _i8m4, _i8, _e8m4) DEF_RVV_TYPE (vuint8m4_t, 15, __rvv_uint8m4_t, uint8, RVVM4QI, _u8m4, _u8, _e8m4) +DEF_RVV_TYPE (vfloat8e4m3m4_t, 20, __rvv_float8e4m3m4_t, float8e4m3, + RVVM4QI, _f8e4m3m4, _f8e4m3, _e8m4) +DEF_RVV_TYPE (vfloat8e5m2m4_t, 20, __rvv_float8e5m2m4_t, float8e5m2, + RVVM4QI, _f8e5m2m4, _f8e5m2, _e8m4) /* Define tuple types for SEW = 8, LMUL = M4. */ DEF_RVV_TUPLE_TYPE (vint8m4x2_t, 16, __rvv_int8m4x2_t, vint8m4_t, int8, 2, _i8m4x2) DEF_RVV_TUPLE_TYPE (vuint8m4x2_t, 17, __rvv_uint8m4x2_t, vuint8m4_t, uint8, 2, _u8m4x2) @@ -228,6 +253,10 @@ DEF_RVV_TUPLE_TYPE (vuint8m4x2_t, 17, __rvv_uint8m4x2_t, vuint8m4_t, uint8, 2, _ DEF_RVV_TYPE (vint8m8_t, 14, __rvv_int8m8_t, int8, RVVM8QI, _i8m8, _i8, _e8m8) DEF_RVV_TYPE (vuint8m8_t, 15, __rvv_uint8m8_t, uint8, RVVM8QI, _u8m8, _u8, _e8m8) +DEF_RVV_TYPE (vfloat8e4m3m8_t, 20, __rvv_float8e4m3m8_t, float8e4m3, + RVVM8QI, _f8e4m3m8, _f8e4m3, _e8m8) +DEF_RVV_TYPE (vfloat8e5m2m8_t, 20, __rvv_float8e5m2m8_t, float8e5m2, + RVVM8QI, _f8e5m2m8, _f8e5m2, _e8m8) /* LMUL = 1/4: Only enable when TARGET_MIN_VLEN > 32. @@ -755,6 +784,11 @@ DEF_RVV_BASE_TYPE (vlmul_ext_x8, get_vector_type (type_idx)) DEF_RVV_BASE_TYPE (vlmul_ext_x16, get_vector_type (type_idx)) DEF_RVV_BASE_TYPE (vlmul_ext_x32, get_vector_type (type_idx)) DEF_RVV_BASE_TYPE (vlmul_ext_x64, get_vector_type (type_idx)) +DEF_RVV_BASE_TYPE (zvdota_acc, get_vector_type (type_idx)) +DEF_RVV_BASE_TYPE (zvdota_uacc, get_vector_type (type_idx)) +DEF_RVV_BASE_TYPE (zvdota_f32_acc, get_vector_type (type_idx)) +DEF_RVV_BASE_TYPE (fp8e4m3_vector, get_vector_type (type_idx)) +DEF_RVV_BASE_TYPE (fp8e5m2_vector, get_vector_type (type_idx)) DEF_RVV_BASE_TYPE (size_ptr, build_pointer_type (size_type_node)) DEF_RVV_BASE_TYPE (tuple_subpart, get_tuple_subpart_type (type_idx)) DEF_RVV_BASE_TYPE (xfqf_float, get_xfqf_float_type (type_idx)) diff --git a/gcc/config/riscv/riscv-vector-builtins.h b/gcc/config/riscv/riscv-vector-builtins.h index aa20cb82c48..10af348c1e6 100644 --- a/gcc/config/riscv/riscv-vector-builtins.h +++ b/gcc/config/riscv/riscv-vector-builtins.h @@ -112,6 +112,10 @@ static const unsigned int CP_WRITE_CSR = 1U << 5; #define RVV_REQUIRE_ELEN_BF_16 (1 << 7) /* Require BF16. */ #define RVV_REQUIRE_ZVFOFP8MIN (1 << 8) /* Require ZVFOFP8MIN extension. */ +/* Sentinel used by builtin expanders for instructions that leave the machine + description "altfmt" attribute at its default. */ +#define RVV_NO_ALTFMT 255U + /* Enumerates the required extensions. */ enum required_ext { @@ -131,6 +135,10 @@ enum required_ext ZVFBFMIN_EXT, /* Zvfbfmin extension */ ZVFBFWMA_EXT, /* Zvfbfwma extension */ ZVFOFP8MIN_EXT, /* Zvfofp8min extension */ + ZVQWDOTA8I_EXT, /* Zvqwdota8i extension */ + ZVQWDOTA16I_EXT, /* Zvqwdota16i extension */ + ZVFWDOTA16BF_EXT, /* Zvfwdota16bf extension */ + ZVFQWDOTA8F_EXT, /* Zvfqwdota8f extension */ XSFVQMACCQOQ_EXT, /* XSFVQMACCQOQ extension */ XSFVQMACCDOD_EXT, /* XSFVQMACCDOD extension */ XSFVFNRCLIPXFQF_EXT, /* XSFVFNRCLIPXFQF extension */ @@ -162,6 +170,10 @@ enum rvv_builtin_partition RVV_PARTITION_ZVFHMIN, RVV_PARTITION_ZVFH, RVV_PARTITION_ZVFOFP8MIN, + RVV_PARTITION_ZVQWDOTA8I, + RVV_PARTITION_ZVQWDOTA16I, + RVV_PARTITION_ZVFWDOTA16BF, + RVV_PARTITION_ZVFQWDOTA8F, RVV_PARTITION_XSFVQMACCQOQ, RVV_PARTITION_XSFVQMACCDOD, RVV_PARTITION_XSFVFNRCLIPXFQF, @@ -216,6 +228,14 @@ static inline const char * required_ext_to_isa_name (enum required_ext required) return "zvfbfwma"; case ZVFOFP8MIN_EXT: return "zvfofp8min"; + case ZVQWDOTA8I_EXT: + return "zvqwdota8i"; + case ZVQWDOTA16I_EXT: + return "zvqwdota16i"; + case ZVFWDOTA16BF_EXT: + return "zvfwdota16bf"; + case ZVFQWDOTA8F_EXT: + return "zvfqwdota8f"; case XSFVQMACCQOQ_EXT: return "xsfvqmaccqoq"; case XSFVQMACCDOD_EXT: @@ -273,6 +293,14 @@ static inline bool required_extensions_specified (enum required_ext required) return TARGET_ZVFBFWMA; case ZVFOFP8MIN_EXT: return TARGET_ZVFOFP8MIN; + case ZVQWDOTA8I_EXT: + return TARGET_ZVQWDOTA8I; + case ZVQWDOTA16I_EXT: + return TARGET_ZVQWDOTA16I; + case ZVFWDOTA16BF_EXT: + return TARGET_ZVFWDOTA16BF; + case ZVFQWDOTA8F_EXT: + return TARGET_ZVFQWDOTA8F; case XSFVQMACCQOQ_EXT: return TARGET_XSFVQMACCQOQ; case XSFVQMACCDOD_EXT: @@ -561,6 +589,7 @@ public: rtx use_contiguous_store_insn (insn_code); rtx use_compare_insn (rtx_code, insn_code); rtx use_ternop_insn (bool, insn_code); + rtx use_zvdota_insn (insn_code, unsigned = RVV_NO_ALTFMT); rtx use_widen_ternop_insn (insn_code); rtx use_scalar_move_insn (insn_code); rtx use_scalar_broadcast_insn (insn_code); diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-intrinsic.c b/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-intrinsic.c new file mode 100644 index 00000000000..0f8c7534755 --- /dev/null +++ b/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-intrinsic.c @@ -0,0 +1,106 @@ +/* { dg-do compile } */ +/* { dg-options "-march=rv64gcv_zvqwdota8i_zvqwdota16i_zvfwdota16bf_zvfqwdota8f_zvfbfmin -mabi=lp64d -O2" } */ + +#include "riscv_vector.h" + +vuint32m1_t +test_vqwdota_uu_8 (vuint32m1_t vd, vuint8mf4_t vs2, + vuint8mf4_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_u8mf4_u8mf4_u32m1 (vd, vs2, vs1, vl); +} + +vint32m1_t +test_vqwdota_us_8 (vint32m1_t vd, vuint8mf4_t vs2, + vint8mf4_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_u8mf4_i8mf4_i32m1 (vd, vs2, vs1, vl); +} + +vint32m1_t +test_vqwdota_su_8 (vint32m1_t vd, vint8mf4_t vs2, + vuint8mf4_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_i8mf4_u8mf4_i32m1 (vd, vs2, vs1, vl); +} + +vint32m1_t +test_vqwdota_ss_8 (vint32m1_t vd, vint8mf4_t vs2, + vint8mf4_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_i8mf4_i8mf4_i32m1 (vd, vs2, vs1, vl); +} + +vuint64m1_t +test_vqwdota_uu_16 (vuint64m1_t vd, vuint16mf2_t vs2, + vuint16mf2_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_u16mf2_u16mf2_u64m1 (vd, vs2, vs1, vl); +} + +vint64m1_t +test_vqwdota_us_16 (vint64m1_t vd, vuint16mf2_t vs2, + vint16mf2_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_u16mf2_i16mf2_i64m1 (vd, vs2, vs1, vl); +} + +vint64m1_t +test_vqwdota_su_16 (vint64m1_t vd, vint16mf2_t vs2, + vuint16mf2_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_i16mf2_u16mf2_i64m1 (vd, vs2, vs1, vl); +} + +vint64m1_t +test_vqwdota_ss_16 (vint64m1_t vd, vint16mf2_t vs2, + vint16mf2_t vs1, size_t vl) +{ + return __riscv_vqwdota_vv_i16mf2_i16mf2_i64m1 (vd, vs2, vs1, vl); +} + +vfloat32m1_t +test_vfqwdota_e4m3_e4m3 (vfloat32m1_t vd, vfloat8e4m3mf4_t vs2, + vfloat8e4m3mf4_t vs1, size_t vl) +{ + return __riscv_vfqwdota_vv_f8e4m3mf4_f8e4m3mf4_f32m1 (vd, vs2, vs1, vl); +} + +vfloat32m1_t +test_vfqwdota_e4m3_e5m2 (vfloat32m1_t vd, vfloat8e4m3mf4_t vs2, + vfloat8e5m2mf4_t vs1, size_t vl) +{ + return __riscv_vfqwdota_vv_f8e4m3mf4_f8e5m2mf4_f32m1 (vd, vs2, vs1, vl); +} + +vfloat32m1_t +test_vfqwdota_e5m2_e4m3 (vfloat32m1_t vd, vfloat8e5m2mf4_t vs2, + vfloat8e4m3mf4_t vs1, size_t vl) +{ + return __riscv_vfqwdota_vv_f8e5m2mf4_f8e4m3mf4_f32m1 (vd, vs2, vs1, vl); +} + +vfloat32m1_t +test_vfqwdota_e5m2_e5m2 (vfloat32m1_t vd, vfloat8e5m2mf4_t vs2, + vfloat8e5m2mf4_t vs1, size_t vl) +{ + return __riscv_vfqwdota_vv_f8e5m2mf4_f8e5m2mf4_f32m1 (vd, vs2, vs1, vl); +} + +vfloat32m1_t +test_vfwdota_bf16 (vfloat32m1_t vd, vbfloat16mf2_t vs2, + vbfloat16mf2_t vs1, size_t vl) +{ + return __riscv_vfwdota_vv_bf16mf2_bf16mf2_f32m1 (vd, vs2, vs1, vl); +} + +/* { dg-final { scan-assembler-times {vqwdotau\.vv} 4 } } */ +/* { dg-final { scan-assembler-times {vqwdotas\.vv} 4 } } */ +/* { dg-final { scan-assembler-times {vfqwdota\.vv} 2 } } */ +/* { dg-final { scan-assembler-times {vfqwdota\.alt\.vv} 2 } } */ +/* { dg-final { scan-assembler-times {vfwdota\.vv} 1 } } */ + +/* { dg-final { scan-assembler-times {vsetvli\s+zero,\s*[a-z0-9]+,\s*e8,\s*mf4,\s*ta,\s*ma} 4 } } */ +/* { dg-final { scan-assembler-times {vsetvli\s+zero,\s*[a-z0-9]+,\s*e8alt,\s*mf4,\s*ta,\s*ma} 4 } } */ +/* { dg-final { scan-assembler-times {vsetvli\s+zero,\s*[a-z0-9]+,\s*e16,\s*mf2,\s*ta,\s*ma} 2 } } */ +/* { dg-final { scan-assembler-times {vsetvli\s+zero,\s*[a-z0-9]+,\s*e16alt,\s*mf2,\s*ta,\s*ma} 3 } } */ diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-no-rm-fail.c b/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-no-rm-fail.c new file mode 100644 index 00000000000..86d834fa754 --- /dev/null +++ b/gcc/testsuite/gcc.target/riscv/rvv/base/zvdota-no-rm-fail.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ +/* { dg-options "-march=rv64gcv_zvfwdota16bf_zvfqwdota8f_zvfbfmin -mabi=lp64d -O2 -Werror=implicit-function-declaration" } */ + +#include "riscv_vector.h" + +void +test_vfwdota_rm (vfloat32m1_t vd, vbfloat16mf2_t vs2, + vbfloat16mf2_t vs1, size_t vl) +{ + (void) __riscv_vfwdota_vv_bf16mf2_bf16mf2_f32m1_rm (vd, vs2, vs1, __RISCV_FRM_RNE, vl); /* { dg-error "implicit declaration of function" } */ +} + +void +test_vfqwdota_rm (vfloat32m1_t vd, vfloat8e4m3mf4_t vs2, + vfloat8e4m3mf4_t vs1, size_t vl) +{ + (void) __riscv_vfqwdota_vv_f8e4m3mf4_f8e4m3mf4_f32m1_rm (vd, vs2, vs1, __RISCV_FRM_RNE, vl); /* { dg-error "implicit declaration of function" } */ +} + +void +test_vfqwdota_e5m2_rm (vfloat32m1_t vd, vfloat8e5m2mf4_t vs2, + vfloat8e4m3mf4_t vs1, size_t vl) +{ + (void) __riscv_vfqwdota_vv_f8e5m2mf4_f8e4m3mf4_f32m1_rm (vd, vs2, vs1, __RISCV_FRM_RNE, vl); /* { dg-error "implicit declaration of function" } */ +} -- 2.43.0
