diff --git a/hdl/chisel/src/kelvin/rvv/RvvCore.scala b/hdl/chisel/src/kelvin/rvv/RvvCore.scala
index 75ae92a..2917755 100644
--- a/hdl/chisel/src/kelvin/rvv/RvvCore.scala
+++ b/hdl/chisel/src/kelvin/rvv/RvvCore.scala
@@ -459,16 +459,22 @@
   io.rvv2lsu <> rvvCoreWrapper.io.rvv2lsu
   io.lsu2rvv <> rvvCoreWrapper.io.lsu2rvv
 
-  io.configState.valid        := rvvCoreWrapper.io.configStateValid
-  io.configState.bits.vl      := rvvCoreWrapper.io.configVl
-  io.configState.bits.vstart  := rvvCoreWrapper.io.configVstart
-  io.configState.bits.ma      := rvvCoreWrapper.io.configMa
-  io.configState.bits.ta      := rvvCoreWrapper.io.configTa
-  io.configState.bits.xrm     := rvvCoreWrapper.io.configXrm
-  io.configState.bits.sew     := rvvCoreWrapper.io.configSew
-  io.configState.bits.lmul    := rvvCoreWrapper.io.configLmul
-  io.rvv_idle                 := rvvCoreWrapper.io.rvv_idle
-  io.queue_capacity           := rvvCoreWrapper.io.queue_capacity
+  // Conservatively mark config state as invalid the cycle when CSR instruction
+  // updates vstart, vxrm or vxsat.
+  io.configState.valid := rvvCoreWrapper.io.configStateValid &&
+      !rvvCoreWrapper.io.vcsr_valid &&
+      !io.csr.vstart_write.valid &&
+      !io.csr.vxrm_write.valid &&
+      !io.csr.vxsat_write.valid
+  io.configState.bits.vl     := rvvCoreWrapper.io.configVl
+  io.configState.bits.vstart := rvvCoreWrapper.io.configVstart
+  io.configState.bits.ma     := rvvCoreWrapper.io.configMa
+  io.configState.bits.ta     := rvvCoreWrapper.io.configTa
+  io.configState.bits.xrm    := rvvCoreWrapper.io.configXrm
+  io.configState.bits.sew    := rvvCoreWrapper.io.configSew
+  io.configState.bits.lmul   := rvvCoreWrapper.io.configLmul
+  io.rvv_idle                := rvvCoreWrapper.io.rvv_idle
+  io.queue_capacity          := rvvCoreWrapper.io.queue_capacity
 
   val vstart_wdata = MuxCase(vstart, Seq(
       rvvCoreWrapper.io.vcsr_valid -> rvvCoreWrapper.io.vcsr_vstart,
diff --git a/hdl/chisel/src/kelvin/rvv/RvvDecode.scala b/hdl/chisel/src/kelvin/rvv/RvvDecode.scala
index 03f1b7d..1a85a16 100644
--- a/hdl/chisel/src/kelvin/rvv/RvvDecode.scala
+++ b/hdl/chisel/src/kelvin/rvv/RvvDecode.scala
@@ -45,6 +45,23 @@
     bits(7, 5)
   }
 
+  // If this instruction is a "reduction" instruction, that requires vstart=0
+  def isReduction(): Bool = {
+    (opcode === RvvCompressedOpcode.RVVALU) && (funct3() === "b010".U) &&
+        MuxLookup(funct6(), false.B)(Seq(
+            "b000000".U -> true.B,  // vredsum
+            "b000001".U -> true.B,  // vredand
+            "b000010".U -> true.B,  // vredor
+            "b000011".U -> true.B,  // vredxor
+            "b000100".U -> true.B,  // vredminu
+            "b000101".U -> true.B,  // vredmin
+            "b000110".U -> true.B,  // vredmaxu
+            "b000111".U -> true.B,  // vredmax
+            "b110000".U -> true.B,  // vwredsumu
+            "b110001".U -> true.B,  // vwredsum
+        ))
+  }
+
   // "Addressing Mode" for loads/store (see Section 7.2 of RVV Spec)
   def mop: RvvAddressingMode.Type = {
     RvvAddressingMode(bits(20, 19))
diff --git a/hdl/chisel/src/kelvin/scalar/Decode.scala b/hdl/chisel/src/kelvin/scalar/Decode.scala
index b808286..828666c 100644
--- a/hdl/chisel/src/kelvin/scalar/Decode.scala
+++ b/hdl/chisel/src/kelvin/scalar/Decode.scala
@@ -263,6 +263,8 @@
     val jalrFault = Output(Vec(p.instructionLanes, Bool()))
     val bxxFault = Output(Vec(p.instructionLanes, Bool()))
     val undefFault = Output(Vec(p.instructionLanes, Bool()))
+    val rvvFault = Option.when(p.enableRvv)(
+        Output(Vec(p.instructionLanes, Bool())))
     val bruTarget = Output(Vec(p.instructionLanes, UInt(32.W)))
     val jalrTarget = Input(Vec(p.instructionLanes, new RegfileBranchTargetIO))
 
@@ -420,12 +422,27 @@
   val fence = decodedInsts.map(x => x.isFency() && (io.mactive || io.lsuActive))
 
   // ---------------------------------------------------------------------------
+  // Csr interlock
+  val csrInterlock = (0 until p.instructionLanes).map(i =>
+    if (i == 0) {
+      true.B
+    } else {
+      !decodedInsts(0).isCsr()
+    }
+  )
+
+  // ---------------------------------------------------------------------------
   // Rvv config interlock rules
   // RVV Load store unit requires valid config state on dispatch.
-  val rvvConfigInterlock = if (p.enableRvv) {
+  val configInvalid = if (p.enableRvv) {
     val configChange = decodedInsts.map(
         x => x.rvv.get.valid && x.rvv.get.bits.isVset())
-    val configInvalid = configChange.scan(!io.rvvState.get.valid)(_ || _)
+    configChange.scan(!io.rvvState.get.valid)(_ || _)
+  } else {
+    Seq.fill(p.instructionLanes)(false.B)
+  }
+
+  val rvvConfigInterlock = if (p.enableRvv) {
     val canDispatchRvv = (0 until p.instructionLanes).map(i =>
         !decodedInsts(i).rvv.get.valid || // Don't lock non-rvv
         !decodedInsts(i).rvv.get.bits.isLoadStore() || // Non-LSU can handle change
@@ -437,6 +454,21 @@
   }
 
   // ---------------------------------------------------------------------------
+  // Rvv reduction
+  // Don't allow reduction instructions to execute if vstart != 0
+  val rvvReductionInterlock = if (p.enableRvv) {
+    (0 until p.instructionLanes).map(i => {
+        val invalidReduction =
+            decodedInsts(i).rvv.get.valid &&
+            decodedInsts(i).rvv.get.bits.isReduction() &&
+            (configInvalid(i) || (io.rvvState.get.bits.vstart =/= 0.U))
+        !invalidReduction
+    })
+  } else {
+    Seq.fill(p.instructionLanes)(true.B)
+  }
+
+  // ---------------------------------------------------------------------------
   // Rvv Interlock
   val rvvInterlock = if (p.enableRvv) {
     val isRvv = decodedInsts.map(x => x.rvv.get.valid)
@@ -499,7 +531,9 @@
       !floatWriteAfterWrite(i) && // Avoid WAW hazards
       !branchInterlock(i) && // Only branch/alu can be dispatched after a branch
       !fence(i) &&           // Don't dispatch if fence interlocked
+      csrInterlock(i) &&
       rvvConfigInterlock(i) &&     // Rvv interlock rules
+      rvvReductionInterlock(i) && // Don't dispatch reduction if vstart != 0
       // rvvLsuInterlock(i) &&  // Dispatch only one Rvv LsuOp
       lsuInterlock(i) && // Ensure lsu instructions can be dispatched into queue
       rvvInterlock(i) && // Ensure rvv instructions can be dispatched into queue
@@ -721,6 +755,22 @@
     io.inst(i).ready := lastReady(i + 1)
   }
 
+  // Fault handling for RVV
+  if (p.enableRvv) {
+    for (i <- 0 until p.instructionLanes) {
+      io.rvvFault.get(i) := (if (i == 0) {
+        // Return fault if vstart != 0
+        val isReduction = decodedInsts(i).rvv.get.valid &&
+            decodedInsts(0).rvv.get.bits.isReduction()
+        val vStartNotZero = io.rvvState.get.valid &&
+            (io.rvvState.get.bits.vstart =/= 0.U)
+        io.inst(0).valid && isReduction && vStartNotZero
+      } else {
+        false.B
+      })
+    }
+  }
+
   for (i <- 0 until p.instructionLanes) {
     val d = decodedInsts(i)
     val rs3Addr = io.inst(i).bits.inst(31,27)
@@ -796,6 +846,8 @@
     io.jalFault(i) := decode(i).io.jalFault && !io.branchTaken && decode(i).io.inst.valid
     io.bxxFault(i) := decode(i).io.bxxFault & !io.branchTaken && decode(i).io.inst.valid
     io.undefFault(i) := decode(i).io.undefFault & !io.branchTaken && decode(i).io.inst.valid
+    if (p.enableRvv) { io.rvvFault.get(i) := false.B }
+
     io.bruTarget(i) := decode(i).io.bruTarget
   }
 
diff --git a/hdl/chisel/src/kelvin/scalar/FaultManager.scala b/hdl/chisel/src/kelvin/scalar/FaultManager.scala
index e8ff77b..5470413 100644
--- a/hdl/chisel/src/kelvin/scalar/FaultManager.scala
+++ b/hdl/chisel/src/kelvin/scalar/FaultManager.scala
@@ -32,6 +32,7 @@
         val jalr = Bool()
         val bxx = Bool()
         val undef = Bool()
+        val rvv = if (p.enableRvv) Some(Bool()) else None
       }))
       val pc = Input(Vec(p.instructionLanes, new Bundle {
         val pc = UInt(32.W)
@@ -51,7 +52,13 @@
     val out = Output(Valid(new FaultManagerOutput))
   })
 
-  val faults = VecInit((0 until p.instructionLanes).map(x => (io.in.fault(x).csr | io.in.fault(x).jal | io.in.fault(x).jalr | io.in.fault(x).bxx | io.in.fault(x).undef)))
+  val faults = VecInit((0 until p.instructionLanes).map(x => (
+      io.in.fault(x).csr |
+      io.in.fault(x).jal |
+      io.in.fault(x).jalr |
+      io.in.fault(x).bxx |
+      io.in.fault(x).undef |
+      io.in.fault(x).rvv.getOrElse(false.B))))
   val fault = faults.reduce(_|_)
   val first_fault = PriorityEncoder(faults)
   val undef_fault = io.in.fault.map(_.undef).reduce(_|_)
@@ -64,6 +71,8 @@
   val jalr_fault_idx = PriorityEncoder(io.in.fault.map(_.jalr))
   val bxx_fault = io.in.fault.map(_.bxx).reduce(_|_)
   val bxx_fault_idx = PriorityEncoder(io.in.fault.map(_.bxx))
+  val rvv_fault = io.in.fault.map(_.rvv.getOrElse(false.B)).reduce(_|_)
+  val rvv_fault_idx = PriorityEncoder(io.in.fault.map(_.rvv.getOrElse(false.B)))
   val instr_access_fault = io.in.memory_fault.valid && io.in.ibus_fault
   val load_fault = io.in.memory_fault.valid && !io.in.memory_fault.bits.write && !io.in.ibus_fault
   val store_fault = io.in.memory_fault.valid && io.in.memory_fault.bits.write && !io.in.ibus_fault
@@ -84,6 +93,7 @@
     (jalr_fault && (jalr_fault_idx === first_fault)) -> 0.U(32.W),
     (bxx_fault && (bxx_fault_idx === first_fault)) -> 0.U(32.W),
     (undef_fault && (undef_fault_idx === first_fault)) -> 2.U(32.W),
+    (rvv_fault && (rvv_fault_idx === first_fault)) -> 2.U(32.W),
   ))
   io.out.bits.mtval := MuxCase(0.U(32.W), Seq(
     load_fault -> io.in.memory_fault.bits.addr,
@@ -94,5 +104,6 @@
     (jalr_fault && (jalr_fault_idx === first_fault)) -> (io.in.jalr(jalr_fault_idx).target & "xFFFFFFFE".U),
     (bxx_fault && (bxx_fault_idx === first_fault)) -> 0.U(32.W),
     (undef_fault && (undef_fault_idx === first_fault)) -> io.in.undef(undef_fault_idx).inst,
+    (rvv_fault && (rvv_fault_idx === first_fault)) -> io.in.undef(rvv_fault_idx).inst,
   ))
 }
\ No newline at end of file
diff --git a/hdl/chisel/src/kelvin/scalar/SCore.scala b/hdl/chisel/src/kelvin/scalar/SCore.scala
index 83f4c54..80cbb86 100644
--- a/hdl/chisel/src/kelvin/scalar/SCore.scala
+++ b/hdl/chisel/src/kelvin/scalar/SCore.scala
@@ -142,6 +142,9 @@
     fault_manager.io.in.fault(i).jalr := dispatch.io.jalrFault(i)
     fault_manager.io.in.fault(i).bxx := dispatch.io.bxxFault(i)
     fault_manager.io.in.fault(i).undef := dispatch.io.undefFault(i)
+    if (p.enableRvv) {
+      fault_manager.io.in.fault(i).rvv.get := dispatch.io.rvvFault.get(i)
+    }
     fault_manager.io.in.pc(i).pc := fetch.io.inst.lanes(i).bits.addr
     fault_manager.io.in.jalr(i).target := regfile.io.target(i).data
     fault_manager.io.in.undef(i).inst := fetch.io.inst.lanes(i).bits.inst
diff --git a/kelvin_test_utils/core_mini_axi_interface.py b/kelvin_test_utils/core_mini_axi_interface.py
index 78c3bc8..162e8b5 100644
--- a/kelvin_test_utils/core_mini_axi_interface.py
+++ b/kelvin_test_utils/core_mini_axi_interface.py
@@ -801,6 +801,15 @@
         timeout_cycles = timeout_cycles - 1
         assert timeout_cycles > 0
 
+  async def wait_for_fault(self, timeout_cycles=1000):
+    cycle_count = 0
+    while self.dut.io_fault.value != 1 and timeout_cycles > 0:
+      await ClockCycles(self.dut.io_aclk, 1)
+      timeout_cycles = timeout_cycles - 1
+      cycle_count += 1
+    assert timeout_cycles > 0
+    return cycle_count
+
   async def watch(self, addr, timeout_cycles=1_000_000):
     elem_addr = addr % 16
     line_addr = addr - elem_addr
diff --git a/kelvin_test_utils/sim_test_fixture.py b/kelvin_test_utils/sim_test_fixture.py
index b387213..ccb2c82 100644
--- a/kelvin_test_utils/sim_test_fixture.py
+++ b/kelvin_test_utils/sim_test_fixture.py
@@ -54,3 +54,10 @@
     async def run_to_halt(self, timeout_cycles=10000):
         await self.core_mini_axi.execute_from(self.entry_point)
         return (await self.core_mini_axi.wait_for_halted(timeout_cycles=timeout_cycles))
+
+    async def run_to_fault(self, timeout_cycles=10000):
+        await self.core_mini_axi.execute_from(self.entry_point)
+        return (await self.core_mini_axi.wait_for_fault(timeout_cycles=timeout_cycles))
+
+    def fault(self):
+        return (self.core_mini_axi.dut.io_fault.value == 1)
\ No newline at end of file
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD
index 3045e5b..1dfea33 100644
--- a/tests/cocotb/BUILD
+++ b/tests/cocotb/BUILD
@@ -190,6 +190,7 @@
 RVV_ARITHMETIC_TESTCASES = [
     "arithmetic_m1_vanilla_ops",
     "reduction_m1_vanilla_ops",
+    "reduction_m1_failure_ops",
     "widen_math_ops_test_impl",
 ]
 # END_TESTCASES_FOR_rvv_arithmetic_cocotb_test
diff --git a/tests/cocotb/rvv/arithmetics/rvv_reduction_template.cc b/tests/cocotb/rvv/arithmetics/rvv_reduction_template.cc
index 780e3e4..e7a06ab 100644
--- a/tests/cocotb/rvv/arithmetics/rvv_reduction_template.cc
+++ b/tests/cocotb/rvv/arithmetics/rvv_reduction_template.cc
@@ -16,16 +16,36 @@
 
 #include <riscv_vector.h>
 
+#define __ATTRIBUTE_IN_DTCM__ \
+    __attribute__((section(".data"))) __attribute__((aligned(16)))
 
-{DTYPE}_t in_buf_1[{IN_DATA_SIZE}] __attribute__((section(".data"))) __attribute__((aligned(16)));
-{DTYPE}_t scalar_input __attribute__((section(".data"))) __attribute__((aligned(16)));
-{DTYPE}_t out_buf __attribute__((section(".data"))) __attribute__((aligned(16)));
+{DTYPE}_t in_buf_1[{IN_DATA_SIZE}] __ATTRIBUTE_IN_DTCM__;
+{DTYPE}_t scalar_input __ATTRIBUTE_IN_DTCM__;
+{DTYPE}_t out_buf __ATTRIBUTE_IN_DTCM__;
+uint32_t vstart __ATTRIBUTE_IN_DTCM__ = 0;
+uint32_t vl __ATTRIBUTE_IN_DTCM__ = {NUM_OPERANDS};
+uint32_t faulted __ATTRIBUTE_IN_DTCM__ = 0;
+uint32_t mcause __ATTRIBUTE_IN_DTCM__ = 0;
+
+// Fault handler to log fault
+extern "C" {
+void kelvin_exception_handler() {
+  faulted = 1;
+  uint32_t local_mcause;
+  asm volatile("csrr %0, mcause" : "=r"(local_mcause));
+  mcause = local_mcause;
+
+  asm volatile("ebreak");
+  while (1) {}
+}
+}
 
 void {REDUCTION_OP}_{SIGN}{SEW}_m1(const {DTYPE}_t* in_buf_1, const {DTYPE}_t scalar_input, {DTYPE}_t* out_buf){
 
-    v{DTYPE}m1_t input_v1 = __riscv_vle{SEW}_v_{SIGN}{SEW}m1(in_buf_1, {NUM_OPERANDS});
-    v{DTYPE}m1_t input_s1 = __riscv_vmv_v_x_{SIGN}{SEW}m1(scalar_input, {NUM_OPERANDS});
-    v{DTYPE}m1_t {REDUCTION_OP}_result = __riscv_v{REDUCTION_OP}_vs_{SIGN}{SEW}m1_{SIGN}{SEW}m1(input_v1, input_s1, {NUM_OPERANDS});
+    v{DTYPE}m1_t input_v1 = __riscv_vle{SEW}_v_{SIGN}{SEW}m1(in_buf_1, vl);
+    v{DTYPE}m1_t input_s1 = __riscv_vmv_v_x_{SIGN}{SEW}m1(scalar_input, vl);
+    asm("csrw vstart, %0" : : "r"(vstart));
+    v{DTYPE}m1_t {REDUCTION_OP}_result = __riscv_v{REDUCTION_OP}_vs_{SIGN}{SEW}m1_{SIGN}{SEW}m1(input_v1, input_s1, vl);
     *out_buf = __riscv_vmv_x_s_{SIGN}{SEW}m1_{SIGN}{SEW}({REDUCTION_OP}_result);
 }
 
diff --git a/tests/cocotb/rvv_arithmetic_cocotb_test.py b/tests/cocotb/rvv_arithmetic_cocotb_test.py
index 607064d..d85d827 100644
--- a/tests/cocotb/rvv_arithmetic_cocotb_test.py
+++ b/tests/cocotb/rvv_arithmetic_cocotb_test.py
@@ -191,6 +191,64 @@
         num_bytes=16)
 
 
+async def reduction_m1_failure_test(dut, dtypes, math_ops: str, num_bytes: int):
+    """RVV reduction test template.
+
+    Each test performs a reduction op loading `in_buf_1` and storing the output to `out_buf`.
+    """
+    m1_failure_op_elfs = [
+        f"rvv_{math_op}_{dtype}_m1.elf" for math_op in math_ops
+        for dtype in dtypes
+    ]
+    pattern_extract = re.compile("rvv_(.*)_(.*)_m1.elf")
+
+    r = runfiles.Create()
+    fixture = await Fixture.Create(dut)
+
+    with tqdm.tqdm(m1_failure_op_elfs) as t:
+        for elf_name in t:
+            t.set_postfix({"binary": os.path.basename(elf_name)})
+            elf_path = r.Rlocation(
+                f"kelvin_hw/tests/cocotb/rvv/arithmetics/{elf_name}")
+            await fixture.load_elf_and_lookup_symbols(
+                elf_path,
+                ['in_buf_1', 'scalar_input', 'out_buf', 'vstart', 'vl',
+                 'faulted', 'mcause'],
+            )
+            math_op, dtype = pattern_extract.match(elf_name).groups()
+            np_type = STR_TO_NP_TYPE[dtype]
+            itemsize = np.dtype(np_type).itemsize
+            num_test_values = int(num_bytes / np.dtype(np_type).itemsize)
+
+            min_value = np.iinfo(np_type).min
+            max_value = np.iinfo(np_type).max + 1  # One above.
+            input_1 = np.random.randint(min_value,
+                                        max_value,
+                                        num_test_values,
+                                        dtype=np_type)
+            input_2 = np.random.randint(min_value, max_value, 1, dtype=np_type)
+
+            await fixture.write('in_buf_1', input_1)
+            await fixture.write('scalar_input', input_2)
+            await fixture.write('vstart', np.array([1], dtype=np.uint32))
+            await fixture.write('out_buf', np.zeros(1, dtype=np_type))
+
+            await fixture.run_to_halt()
+            faulted = (await fixture.read('faulted', 4)).view(np.uint32)
+            mcause = (await fixture.read('mcause', 4)).view(np.uint32)
+            assert(faulted == True)
+            assert(mcause == 0x2) # Invalid instruction
+
+
+@cocotb.test()
+async def reduction_m1_failure_ops(dut):
+    await reduction_m1_failure_test(
+        dut=dut,
+        dtypes=["int8", "int16", "int32", "uint8", "uint16", "uint32"],
+        math_ops=["redsum", "redmin", "redmax"],
+        num_bytes=16)
+
+
 async def _widen_math_ops_test_impl(
     dut,
     dtypes,
diff --git a/toolchain/crt/BUILD b/toolchain/crt/BUILD
index 611960b..70279e8 100644
--- a/toolchain/crt/BUILD
+++ b/toolchain/crt/BUILD
@@ -18,6 +18,7 @@
     name = "crt",
     srcs = [
         "crt.S",
+        "kelvin_exceptions.cc",
         "kelvin_gloss.cc",
         "kelvin_start.S",
     ],
diff --git a/toolchain/crt/kelvin_exceptions.cc b/toolchain/crt/kelvin_exceptions.cc
new file mode 100644
index 0000000..0a7297e
--- /dev/null
+++ b/toolchain/crt/kelvin_exceptions.cc
@@ -0,0 +1,20 @@
+// Copyright 2025 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//      http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+extern "C" {
+void __attribute__((weak)) kelvin_exception_handler() {
+  asm volatile("ebreak");
+  while (1) {}
+}
+}
\ No newline at end of file
diff --git a/toolchain/crt/kelvin_start.S b/toolchain/crt/kelvin_start.S
index d87ffce..dc44a8e 100644
--- a/toolchain/crt/kelvin_start.S
+++ b/toolchain/crt/kelvin_start.S
@@ -16,6 +16,8 @@
 
 // A starting functions for simple kelvin programs.
 
+.extern kelvin_exception_handler
+
 /**
  * Entry point.
  */
@@ -82,7 +84,7 @@
         # simply call ebreak.
         # Users who require real trap handling should
         # install their own trap vector.
-        la t0, failure
+        la t0, kelvin_exception_handler
         csrw mtvec, t0
 
         # Set up sentinel value in _ret