diff --git a/hdl/chisel/src/kelvin/scalar/Lsu.scala b/hdl/chisel/src/kelvin/scalar/Lsu.scala
index 60ae397..e034011 100644
--- a/hdl/chisel/src/kelvin/scalar/Lsu.scala
+++ b/hdl/chisel/src/kelvin/scalar/Lsu.scala
@@ -187,6 +187,30 @@
   }
 }
 
+object ComputeIndexedAddrs {
+  def apply(bytesPerSlot: Int,
+            baseAddr: UInt,
+            indices: UInt,
+            elemWidth: UInt): Vec[UInt] = {
+    val indices8 = UIntToVec(indices, 8)
+    val indices16 = UIntToVec(indices, 16)
+    val indices32 = UIntToVec(indices, 32)
+
+    MuxCase(VecInit.fill(bytesPerSlot)(0.U(32.W)), Seq(
+      // elemWidth validation is done at decode time.
+      // 8-bit indices. Each byte has its own offset.
+      (elemWidth === "b000".U) -> VecInit((0 until bytesPerSlot).map(
+          i => (baseAddr + indices8(i))(31, 0))),
+      // 16-bit indices. Each 2-byte element has an offset.
+      (elemWidth === "b101".U) -> VecInit((0 until bytesPerSlot).map(
+          i => (baseAddr + indices16(i >> 1))(31, 0) + (i & 1).U)),
+      // 32-bit indices. Each 4-byte element has an offset.
+      (elemWidth === "b110".U) -> VecInit((0 until bytesPerSlot).map(
+          i => (baseAddr + indices32(i >> 2))(31, 0) + (i & 3).U))
+    ))
+  }
+}
+
 // bytesPerSlot is the number of bytes in a vector register
 // bytesPerLine is the number of bytes in the AXI bus
 class LsuSlot(bytesPerSlot: Int, bytesPerLine: Int) extends Bundle {
@@ -253,8 +277,11 @@
         op.isOneOf(LsuOp.VLOAD_UNIT, LsuOp.VSTORE_UNIT) ->
             VecInit((0 until bytesPerSlot).map(i => baseAddr + i.U)),
         op.isOneOf(LsuOp.VLOAD_STRIDED, LsuOp.VSTORE_STRIDED) ->
-            ComputeStridedAddrs(bytesPerSlot, baseAddr, stride, elemWidth)
-        // TODO(derekjchow): Support indexed
+            ComputeStridedAddrs(bytesPerSlot, baseAddr, stride, elemWidth),
+        op.isOneOf(LsuOp.VLOAD_OINDEXED, LsuOp.VLOAD_UINDEXED,
+                   LsuOp.VSTORE_OINDEXED, LsuOp.VSTORE_UINDEXED) ->
+            ComputeIndexedAddrs(bytesPerSlot, baseAddr, rvv2lsu.idx.bits.data,
+                                elemWidth)
         // TODO(derekjchow): Support segmented
     ))
     result.elemWidth := elemWidth
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD
index 989df39..d274166 100644
--- a/tests/cocotb/BUILD
+++ b/tests/cocotb/BUILD
@@ -315,6 +315,7 @@
         "//kelvin_test_utils:sim_test_fixture",
     ],
     data = [
+        '//tests/cocotb/rvv/load_store:load8_indexed_m1.elf',
         '//tests/cocotb/rvv/load_store:load8_stride2_m1.elf',
         '//tests/cocotb/rvv/load_store:load8_stride2_m1_partial.elf',
         '//tests/cocotb/rvv/load_store:load8_stride2_mf4.elf',
@@ -324,6 +325,7 @@
         '//tests/cocotb/rvv/load_store:load16_stride4_mf2.elf',
         '//tests/cocotb/rvv/load_store:load32_stride8_m1.elf',
         '//tests/cocotb/rvv/load_store:load32_stride8_m1_partial.elf',
+        '//tests/cocotb/rvv/load_store:store8_indexed_m1.elf',
     ],
     size = "large",
 )
diff --git a/tests/cocotb/rvv/load_store/BUILD b/tests/cocotb/rvv/load_store/BUILD
index f6c3aad..d19281c 100644
--- a/tests/cocotb/rvv/load_store/BUILD
+++ b/tests/cocotb/rvv/load_store/BUILD
@@ -20,6 +20,9 @@
 template_rule(
     kelvin_v2_binary,
     {
+        "load8_indexed_m1": {
+            "srcs": ["load8_indexed_m1.cc"],
+        },
         "load8_stride2_m1": {
             "srcs": ["load8_stride2_m1.cc"],
         },
@@ -47,5 +50,8 @@
         "load32_stride8_m1_partial": {
             "srcs": ["load32_stride8_m1_partial.cc"],
         },
+        "store8_indexed_m1": {
+            "srcs": ["store8_indexed_m1.cc"],
+        },
     },
 )
diff --git a/tests/cocotb/rvv/load_store/load8_indexed_m1.cc b/tests/cocotb/rvv/load_store/load8_indexed_m1.cc
new file mode 100644
index 0000000..8c6c9a7
--- /dev/null
+++ b/tests/cocotb/rvv/load_store/load8_indexed_m1.cc
@@ -0,0 +1,28 @@
+// 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.
+
+#include <riscv_vector.h>
+#include <stdint.h>
+
+uint8_t input_indices[128] __attribute__((section(".data")));
+uint8_t input_data[4096] __attribute__((section(".data")));
+uint8_t output_data[128] __attribute__((section(".data")));
+
+int main(int argc, char **argv) {
+  vuint8m1_t indices = __riscv_vle8_v_u8m1(input_indices, /*vl=*/16);
+  vuint8m1_t data = __riscv_vloxei8_v_u8m1(input_data, indices, /*vl=*/16);
+  __riscv_vse8_v_u8m1(output_data, data, /*vl=*/16);
+
+  return 0;
+}
diff --git a/tests/cocotb/rvv/load_store/store8_indexed_m1.cc b/tests/cocotb/rvv/load_store/store8_indexed_m1.cc
new file mode 100644
index 0000000..221ff76
--- /dev/null
+++ b/tests/cocotb/rvv/load_store/store8_indexed_m1.cc
@@ -0,0 +1,28 @@
+// 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.
+
+#include <riscv_vector.h>
+#include <stdint.h>
+
+uint8_t input_indices[128] __attribute__((section(".data")));
+uint8_t input_data[128] __attribute__((section(".data")));
+uint8_t output_data[4096] __attribute__((section(".data")));
+
+int main(int argc, char **argv) {
+  vuint8m1_t indices = __riscv_vle8_v_u8m1(input_indices, /*vl=*/16);
+  vuint8m1_t data = __riscv_vle8_v_u8m1(input_data, /*vl=*/16);
+  __riscv_vsoxei8_v_u8m1(output_data, indices, data, /*vl=*/16);
+
+  return 0;
+}
diff --git a/tests/cocotb/rvv_load_store_test.py b/tests/cocotb/rvv_load_store_test.py
index 6e4b481..640d1f2 100644
--- a/tests/cocotb/rvv_load_store_test.py
+++ b/tests/cocotb/rvv_load_store_test.py
@@ -59,6 +59,88 @@
 
     assert (actual_outputs == expected_outputs).all(), debug_msg
 
+async def vector_load_indexed(
+        dut,
+        elf_name: str,
+        dtype,
+):
+    """RVV load-store test template for indexed loads.
+
+    Each test performs a gather operation and writes the result to an output.
+    """
+    fixture = await Fixture.Create(dut)
+    await fixture.load_elf_and_lookup_symbols(
+        '../tests/cocotb/rvv/load_store/' + elf_name,
+        ['input_indices', 'input_data', 'output_data'],
+    )
+
+    indices_count = 16 // np.dtype(dtype).itemsize
+    in_data_count = 4096 // np.dtype(dtype).itemsize
+    out_data_count = 16 // np.dtype(dtype).itemsize
+
+    min_value = np.iinfo(dtype).min
+    max_value = np.iinfo(dtype).max + 1  # One above.
+    rng = np.random.default_rng()
+    input_data = rng.integers(min_value, max_value, in_data_count, dtype=dtype)
+    input_indices = rng.integers(
+        0, min(max_value, in_data_count-1), indices_count, dtype=dtype)
+
+    expected_outputs = np.take(input_data, input_indices)
+
+    await fixture.write('input_data', input_data)
+    await fixture.write('input_indices', input_indices)
+    await fixture.write('output_data', np.zeros([out_data_count], dtype=dtype))
+
+    await fixture.run_to_halt()
+
+    actual_outputs = (await fixture.read(
+        'output_data', out_data_count * np.dtype(dtype).itemsize)).view(dtype)
+
+    assert (actual_outputs == expected_outputs).all()
+
+async def vector_store_indexed(
+        dut,
+        elf_name: str,
+        dtype,
+):
+    """RVV load-store test template for indexed stores.
+
+    Each test loads indices and data and performs a scatter operation.
+    """
+    fixture = await Fixture.Create(dut)
+    await fixture.load_elf_and_lookup_symbols(
+        '../tests/cocotb/rvv/load_store/' + elf_name,
+        ['input_indices', 'input_data', 'output_data'],
+    )
+
+    indices_count = 16 // np.dtype(dtype).itemsize
+    in_data_count = 16 // np.dtype(dtype).itemsize
+    out_data_count = 4096 // np.dtype(dtype).itemsize
+
+    min_value = np.iinfo(dtype).min
+    max_value = np.iinfo(dtype).max + 1  # One above.
+    rng = np.random.default_rng()
+    input_data = rng.integers(min_value, max_value, in_data_count, dtype=dtype)
+    input_indices = rng.integers(
+        0, min(max_value, out_data_count-1), indices_count, dtype=dtype)
+    original_outputs = rng.integers(
+        min_value, max_value, out_data_count, dtype=dtype)
+
+    await fixture.write('input_data', input_data)
+    await fixture.write('input_indices', input_indices)
+    await fixture.write('output_data', original_outputs)
+
+    expected_outputs = np.copy(original_outputs)
+    for idx, data in zip(input_indices, input_data):
+      expected_outputs[idx] = data
+
+    await fixture.run_to_halt()
+
+    actual_outputs = (await fixture.read(
+        'output_data', out_data_count * np.dtype(dtype).itemsize)).view(dtype)
+
+    assert (actual_outputs == expected_outputs).all()
+
 @cocotb.test()
 async def load8_stride2_m1(dut):
     await vector_load_store(
@@ -157,3 +239,19 @@
         out_size = 64,
         pattern = list(range(0, 32)),
     )
+
+@cocotb.test()
+async def load8_indexed_m1(dut):
+    await vector_load_indexed(
+        dut = dut,
+        elf_name = 'load8_indexed_m1.elf',
+        dtype = np.uint8,
+    )
+
+@cocotb.test()
+async def store8_indexed_m1(dut):
+    await vector_store_indexed(
+        dut = dut,
+        elf_name = 'store8_indexed_m1.elf',
+        dtype = np.uint8,
+    )