From: juzhe.zhong@rivai.ai
To: gcc-patches@gcc.gnu.org
Cc: kito.cheng@gmail.com, kito.cheng@sifive.com, palmer@dabbelt.com,
palmer@rivosinc.com, jeffreyalaw@gmail.com, rdapp.gcc@gmail.com,
Juzhe-Zhong <juzhe.zhong@rivai.ai>
Subject: [PATCH] RISC-V: Add mode switching target hook to insert rounding mode config for fixed-point instructions
Date: Wed, 17 May 2023 17:58:18 +0800 [thread overview]
Message-ID: <20230517095818.1285188-1-juzhe.zhong@rivai.ai> (raw)
From: Juzhe-Zhong <juzhe.zhong@rivai.ai>
Hi, this patch support the new coming fixed-point intrinsics:
https://github.com/riscv-non-isa/rvv-intrinsic-doc/pull/222
Insert fixed-point rounding mode configuration by mode switching target hook.
Mode switching target hook is implemented applying LCM (Lazy code Motion).
So the performance && correctness can be well trusted.
Here is the example:
void f (void * in, void *out, int32_t x, int n, int m)
{
for (int i = 0; i < n; i++) {
vint32m1_t v = __riscv_vle32_v_i32m1 (in + i, 4);
vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i, 4);
vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
__riscv_vse32_v_i32m1 (out + 100 + i, v3, 4);
}
for (int i = 0; i < n; i++) {
vint32m1_t v = __riscv_vle32_v_i32m1 (in + i + 1000, 4);
vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i + 1000, 4);
vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
__riscv_vse32_v_i32m1 (out + 100 + i + 1000, v3, 4);
}
}
ASM:
...
csrwi vxrm,2
vsetivli zero,4,e32,m1,tu,ma
...
Loop 1
...
Loop 2
mode switching can global recognize both Loop 1 and Loop 2 are using RDN
rounding mode and hoist such single "csrwi vxrm,2" to dominate both Loop 1
and Loop 2.
Besides, I have add correctness check sanity tests in this patch too.
Ok for trunk ?
gcc/ChangeLog:
* config/riscv/riscv-opts.h (enum riscv_entity): New enum.
* config/riscv/riscv.cc (riscv_emit_mode_set): New function.
(riscv_mode_needed): Ditto.
(riscv_mode_after): Ditto.
(riscv_mode_entry): Ditto.
(riscv_mode_exit): Ditto.
(riscv_mode_priority): Ditto.
(TARGET_MODE_EMIT): New target hook.
(TARGET_MODE_NEEDED): Ditto.
(TARGET_MODE_AFTER): Ditto.
(TARGET_MODE_ENTRY): Ditto.
(TARGET_MODE_EXIT): Ditto.
(TARGET_MODE_PRIORITY): Ditto.
* config/riscv/riscv.h (OPTIMIZE_MODE_SWITCHING): Ditto.
(NUM_MODES_FOR_MODE_SWITCHING): Ditto.
* config/riscv/riscv.md: Add csrwvxrm.
* config/riscv/vector.md (rnu,rne,rdn,rod,none): New attribute.
(vxrmsi): New pattern.
gcc/testsuite/ChangeLog:
* gcc.target/riscv/rvv/base/vxrm-10.c: New test.
* gcc.target/riscv/rvv/base/vxrm-6.c: New test.
* gcc.target/riscv/rvv/base/vxrm-7.c: New test.
* gcc.target/riscv/rvv/base/vxrm-8.c: New test.
* gcc.target/riscv/rvv/base/vxrm-9.c: New test.
---
gcc/config/riscv/riscv-opts.h | 8 ++
gcc/config/riscv/riscv.cc | 104 ++++++++++++++++++
gcc/config/riscv/riscv.h | 6 +-
gcc/config/riscv/riscv.md | 3 +-
gcc/config/riscv/vector.md | 29 +++++
.../gcc.target/riscv/rvv/base/vxrm-10.c | 26 +++++
.../gcc.target/riscv/rvv/base/vxrm-6.c | 15 +++
.../gcc.target/riscv/rvv/base/vxrm-7.c | 16 +++
.../gcc.target/riscv/rvv/base/vxrm-8.c | 18 +++
.../gcc.target/riscv/rvv/base/vxrm-9.c | 26 +++++
10 files changed, 249 insertions(+), 2 deletions(-)
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-10.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-6.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-7.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-8.c
create mode 100644 gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-9.c
diff --git a/gcc/config/riscv/riscv-opts.h b/gcc/config/riscv/riscv-opts.h
index 1b2e6de5e1b..2a16402265a 100644
--- a/gcc/config/riscv/riscv-opts.h
+++ b/gcc/config/riscv/riscv-opts.h
@@ -91,6 +91,14 @@ enum riscv_multilib_select_kind {
select_by_abi,
};
+/* ENTITIES in mode switching. */
+enum riscv_entity
+{
+ RISCV_VXRM = 0,
+ RISCV_FRM,
+ MAX_RISCV_ENTITIES
+};
+
#define MASK_ZICSR (1 << 0)
#define MASK_ZIFENCEI (1 << 1)
diff --git a/gcc/config/riscv/riscv.cc b/gcc/config/riscv/riscv.cc
index de5b87b1a87..0d1b83f4315 100644
--- a/gcc/config/riscv/riscv.cc
+++ b/gcc/config/riscv/riscv.cc
@@ -7513,6 +7513,95 @@ riscv_vectorize_preferred_vector_alignment (const_tree type)
return TYPE_ALIGN (type);
}
+/* Implement Mode switching. */
+
+static void
+riscv_emit_mode_set (int entity, int mode, int prev_mode,
+ HARD_REG_SET regs_live ATTRIBUTE_UNUSED)
+{
+ switch (entity)
+ {
+ case RISCV_VXRM:
+ if (mode != VXRM_MODE_NONE && mode != prev_mode)
+ emit_insn (gen_vxrmsi (gen_int_mode (mode, SImode)));
+ break;
+ default:
+ gcc_unreachable ();
+ }
+}
+
+/* Return mode that entity must be switched into
+ prior to the execution of insn. */
+
+static int
+riscv_mode_needed (int entity, rtx_insn *insn)
+{
+ switch (entity)
+ {
+ case RISCV_VXRM:
+ return recog_memoized (insn) >= 0 ? get_attr_vxrm_mode (insn)
+ : VXRM_MODE_NONE;
+ default:
+ gcc_unreachable ();
+ }
+}
+
+/* Return the mode that an insn results in. */
+
+static int
+riscv_mode_after (int entity, int mode, rtx_insn *insn)
+{
+ switch (entity)
+ {
+ case RISCV_VXRM:
+ if (recog_memoized (insn) >= 0)
+ return reg_mentioned_p (gen_rtx_REG (SImode, VXRM_REGNUM),
+ PATTERN (insn))
+ ? get_attr_vxrm_mode (insn)
+ : mode;
+ else
+ return mode;
+ default:
+ gcc_unreachable ();
+ }
+}
+
+/* Return a mode that ENTITY is assumed to be
+ switched to at function entry. */
+
+static int
+riscv_mode_entry (int entity)
+{
+ switch (entity)
+ {
+ case RISCV_VXRM:
+ return VXRM_MODE_NONE;
+ default:
+ gcc_unreachable ();
+ }
+}
+
+/* Return a mode that ENTITY is assumed to be
+ switched to at function exit. */
+
+static int
+riscv_mode_exit (int entity)
+{
+ switch (entity)
+ {
+ case RISCV_VXRM:
+ return VXRM_MODE_NONE;
+ default:
+ gcc_unreachable ();
+ }
+}
+
+static int
+riscv_mode_priority (int, int n)
+{
+ return n;
+}
+
/* Initialize the GCC target structure. */
#undef TARGET_ASM_ALIGNED_HI_OP
#define TARGET_ASM_ALIGNED_HI_OP "\t.half\t"
@@ -7789,6 +7878,21 @@ riscv_vectorize_preferred_vector_alignment (const_tree type)
#define TARGET_VECTORIZE_PREFERRED_VECTOR_ALIGNMENT \
riscv_vectorize_preferred_vector_alignment
+/* Mode switching hooks. */
+
+#undef TARGET_MODE_EMIT
+#define TARGET_MODE_EMIT riscv_emit_mode_set
+#undef TARGET_MODE_NEEDED
+#define TARGET_MODE_NEEDED riscv_mode_needed
+#undef TARGET_MODE_AFTER
+#define TARGET_MODE_AFTER riscv_mode_after
+#undef TARGET_MODE_ENTRY
+#define TARGET_MODE_ENTRY riscv_mode_entry
+#undef TARGET_MODE_EXIT
+#define TARGET_MODE_EXIT riscv_mode_exit
+#undef TARGET_MODE_PRIORITY
+#define TARGET_MODE_PRIORITY riscv_mode_priority
+
struct gcc_target targetm = TARGET_INITIALIZER;
#include "gt-riscv.h"
diff --git a/gcc/config/riscv/riscv.h b/gcc/config/riscv/riscv.h
index f55bd6112a8..29f2c07ce5d 100644
--- a/gcc/config/riscv/riscv.h
+++ b/gcc/config/riscv/riscv.h
@@ -845,7 +845,7 @@ typedef struct {
"fs0", "fs1", "fa0", "fa1", "fa2", "fa3", "fa4", "fa5", \
"fa6", "fa7", "fs2", "fs3", "fs4", "fs5", "fs6", "fs7", \
"fs8", "fs9", "fs10","fs11","ft8", "ft9", "ft10","ft11", \
- "arg", "frame", "vl", "vtype", "N/A", "N/A", "N/A", "N/A", \
+ "arg", "frame", "vl", "vtype", "vxrm", "N/A", "N/A", "N/A", \
"N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", \
"N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", \
"N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", "N/A", \
@@ -1107,4 +1107,8 @@ extern void riscv_remove_unneeded_save_restore_calls (void);
CONST1_RTX for the simplification. */
#define VECTOR_STORE_FLAG_VALUE(MODE) CONSTM1_RTX (GET_MODE_INNER (MODE))
+/* Mode switching (Lazy code motion) for RVV rounding mode instructions. */
+#define OPTIMIZE_MODE_SWITCHING(ENTITY) (TARGET_VECTOR)
+#define NUM_MODES_FOR_MODE_SWITCHING {VXRM_MODE_NONE}
+
#endif /* ! GCC_RISCV_H */
diff --git a/gcc/config/riscv/riscv.md b/gcc/config/riscv/riscv.md
index 91808d6bd2a..a9179931217 100644
--- a/gcc/config/riscv/riscv.md
+++ b/gcc/config/riscv/riscv.md
@@ -284,6 +284,7 @@
;; Classification of RVV instructions which will be added to each RVV .md pattern and used by scheduler.
;; rdvlenb vector byte length vlenb csrr read
;; rdvl vector length vl csrr read
+;; wrvxrm vector fixed-point rounding mode write
;; vsetvl vector configuration-setting instrucions
;; 7. Vector Loads and Stores
;; vlde vector unit-stride load instructions
@@ -387,7 +388,7 @@
mtc,mfc,const,arith,logical,shift,slt,imul,idiv,move,fmove,fadd,fmul,
fmadd,fdiv,fcmp,fcvt,fsqrt,multi,auipc,sfb_alu,nop,ghost,bitmanip,rotate,
clmul,min,max,minu,maxu,clz,ctz,cpop,
- atomic,condmove,crypto,rdvlenb,rdvl,vsetvl,vlde,vste,vldm,vstm,vlds,vsts,
+ atomic,condmove,crypto,rdvlenb,rdvl,wrvxrm,vsetvl,vlde,vste,vldm,vstm,vlds,vsts,
vldux,vldox,vstux,vstox,vldff,vldr,vstr,
vlsegde,vssegte,vlsegds,vssegts,vlsegdux,vlsegdox,vssegtux,vssegtox,vlsegdff,
vialu,viwalu,vext,vicalu,vshift,vnshift,vicmp,viminmax,
diff --git a/gcc/config/riscv/vector.md b/gcc/config/riscv/vector.md
index a06b84d7473..1c102943622 100644
--- a/gcc/config/riscv/vector.md
+++ b/gcc/config/riscv/vector.md
@@ -407,6 +407,26 @@
(symbol_ref "INTVAL (operands[4])")]
(const_int INVALID_ATTRIBUTE)))
+;; Defines rounding mode of an fixed-point operation.
+
+(define_attr "vxrm_mode" "rnu,rne,rdn,rod,none"
+ (cond [(and (eq_attr "type" "vsalu,vaalu,vsmul,vsshift,vnclip")
+ (match_test "INTVAL(operands[9]) == riscv_vector::VXRM_RNU"))
+ (const_string "rnu")
+
+ (and (eq_attr "type" "vsalu,vaalu,vsmul,vsshift,vnclip")
+ (match_test "INTVAL(operands[9]) == riscv_vector::VXRM_RNE"))
+ (const_string "rne")
+
+ (and (eq_attr "type" "vsalu,vaalu,vsmul,vsshift,vnclip")
+ (match_test "INTVAL(operands[9]) == riscv_vector::VXRM_RDN"))
+ (const_string "rdn")
+
+ (and (eq_attr "type" "vsalu,vaalu,vsmul,vsshift,vnclip")
+ (match_test "INTVAL(operands[9]) == riscv_vector::VXRM_ROD"))
+ (const_string "rod")]
+ (const_string "none")))
+
;; -----------------------------------------------------------------
;; ---- Miscellaneous Operations
;; -----------------------------------------------------------------
@@ -584,6 +604,15 @@
"TARGET_VECTOR"
"")
+;; Set VXRM
+(define_insn "vxrmsi"
+ [(set (reg:SI VXRM_REGNUM)
+ (match_operand 0 "const_int_operand" "i"))]
+ "TARGET_VECTOR"
+ "csrwi\tvxrm,%0"
+ [(set_attr "type" "wrvxrm")
+ (set_attr "mode" "SI")])
+
;; -----------------------------------------------------------------
;; ---- Moves Operations
;; -----------------------------------------------------------------
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-10.c b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-10.c
new file mode 100644
index 00000000000..a707aa1645e
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-10.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d -O3 -fno-schedule-insns -fno-schedule-insns2" } */
+
+#include "riscv_vector.h"
+
+void f (void * in, void *out, int32_t x, int n, int m)
+{
+ for (int i = 0; i < n; i++) {
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in + i, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100 + i, v3, 4);
+ }
+
+ for (int i = 0; i < n; i++) {
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in + i + 1000, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i + 1000, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RNE, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RNE, 4);
+ __riscv_vse32_v_i32m1 (out + 100 + i + 1000, v3, 4);
+ }
+}
+
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*1} 1 } } */
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-6.c b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-6.c
new file mode 100644
index 00000000000..4b346d67c27
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-6.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d -O3" } */
+
+#include "riscv_vector.h"
+
+void f (void * in, void *out, int32_t x, int n, int m)
+{
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100, v3, 4);
+}
+
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-7.c b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-7.c
new file mode 100644
index 00000000000..1ca795ce3f4
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-7.c
@@ -0,0 +1,16 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d -O3" } */
+
+#include "riscv_vector.h"
+
+void f (void * in, void *out, int32_t x, int n, int m)
+{
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RNE, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100, v3, 4);
+}
+
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*1} 1 } } */
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-8.c b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-8.c
new file mode 100644
index 00000000000..5799f731e21
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-8.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d -O3 -fno-schedule-insns -fno-schedule-insns2" } */
+
+#include "riscv_vector.h"
+
+void f (void * in, void *out, int32_t x, int n, int m)
+{
+ for (int i = 0; i < n; i++) {
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in + i, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100 + i, v3, 4);
+ }
+}
+
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2\s+vsetivli\s+zero,\s*4,\s*e32,\s*m1,\s*tu,\s*m[au]} 1 } } */
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2} 1 } } */
diff --git a/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-9.c b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-9.c
new file mode 100644
index 00000000000..13921d4af21
--- /dev/null
+++ b/gcc/testsuite/gcc.target/riscv/rvv/base/vxrm-9.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-options "-march=rv64gcv -mabi=lp64d -O3 -fno-schedule-insns -fno-schedule-insns2" } */
+
+#include "riscv_vector.h"
+
+void f (void * in, void *out, int32_t x, int n, int m)
+{
+ for (int i = 0; i < n; i++) {
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in + i, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100 + i, v3, 4);
+ }
+
+ for (int i = 0; i < n; i++) {
+ vint32m1_t v = __riscv_vle32_v_i32m1 (in + i + 1000, 4);
+ vint32m1_t v2 = __riscv_vle32_v_i32m1_tu (v, in + 100 + i + 1000, 4);
+ vint32m1_t v3 = __riscv_vaadd_vx_i32m1 (v2, 0, VXRM_RDN, 4);
+ v3 = __riscv_vaadd_vx_i32m1 (v3, 3, VXRM_RDN, 4);
+ __riscv_vse32_v_i32m1 (out + 100 + i + 1000, v3, 4);
+ }
+}
+
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2\s+vsetivli\s+zero,\s*4,\s*e32,\s*m1,\s*tu,\s*m[au]} 1 } } */
+/* { dg-final { scan-assembler-times {csrwi\s+vxrm,\s*2} 1 } } */
--
2.36.3
next reply other threads:[~2023-05-17 9:58 UTC|newest]
Thread overview: 8+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-05-17 9:58 juzhe.zhong [this message]
2023-05-17 10:05 ` Kito Cheng
2023-05-17 14:56 ` 钟居哲
2023-05-17 15:18 ` Li, Pan2
2023-05-19 10:41 ` Andreas Schwab
2023-05-19 12:29 ` Li, Pan2
2023-05-19 12:45 ` Li, Pan2
2023-05-19 13:59 ` Andreas Schwab
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20230517095818.1285188-1-juzhe.zhong@rivai.ai \
--to=juzhe.zhong@rivai.ai \
--cc=gcc-patches@gcc.gnu.org \
--cc=jeffreyalaw@gmail.com \
--cc=kito.cheng@gmail.com \
--cc=kito.cheng@sifive.com \
--cc=palmer@dabbelt.com \
--cc=palmer@rivosinc.com \
--cc=rdapp.gcc@gmail.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).