Matmul rvv intrinsics in cpp This add matmul test in cpp along with ablity to count cycles utilized. Change-Id: I01d802519b2e6c4854d63ad858f0466c8de693d2
diff --git a/kelvin_test_utils/core_mini_axi_interface.py b/kelvin_test_utils/core_mini_axi_interface.py index 4586f0e..78c3bc8 100644 --- a/kelvin_test_utils/core_mini_axi_interface.py +++ b/kelvin_test_utils/core_mini_axi_interface.py
@@ -776,10 +776,13 @@ self.dut.io_irq.value = 0 async def wait_for_halted(self, timeout_cycles=1000): + cycle_count = 0 while self.dut.io_halted.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 wait_for_halted_semihost(self, elf, timeout_cycles=1000000): tohost = self.lookup_symbol(elf, "tohost")
diff --git a/kelvin_test_utils/sim_test_fixture.py b/kelvin_test_utils/sim_test_fixture.py index ae85ad0..b387213 100644 --- a/kelvin_test_utils/sim_test_fixture.py +++ b/kelvin_test_utils/sim_test_fixture.py
@@ -53,4 +53,4 @@ async def run_to_halt(self, timeout_cycles=10000): await self.core_mini_axi.execute_from(self.entry_point) - await self.core_mini_axi.wait_for_halted(timeout_cycles=timeout_cycles) + return (await self.core_mini_axi.wait_for_halted(timeout_cycles=timeout_cycles))
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD index 5259869..05cac01 100644 --- a/tests/cocotb/BUILD +++ b/tests/cocotb/BUILD
@@ -435,10 +435,10 @@ "//kelvin_test_utils:sim_test_fixture", "@bazel_tools//tools/python/runfiles", ], - "data": ["//tests/cocotb/rvv/ml_ops:rvv_matmul_assembly"], + "data": ["//tests/cocotb/rvv/ml_ops:rvv_mlop_tests"], "size": "large", }, - vcs_data = ["//tests/cocotb/rvv/ml_ops:rvv_matmul_assembly"] + [":coverage_exclude.cfg"], + vcs_data = ["//tests/cocotb/rvv/ml_ops:rvv_mlop_tests"] + [":coverage_exclude.cfg"], vcs_build_args = VCS_BUILD_ARGS, vcs_test_args = VCS_TEST_ARGS, vcs_defines = VCS_DEFINES,
diff --git a/tests/cocotb/rvv/ml_ops/BUILD b/tests/cocotb/rvv/ml_ops/BUILD index e6f4ff4..62b479b 100644 --- a/tests/cocotb/rvv/ml_ops/BUILD +++ b/tests/cocotb/rvv/ml_ops/BUILD
@@ -23,5 +23,16 @@ "rvv_matmul_assembly": { "srcs": ["rvv_matmul_assembly.cc"], }, + "rvv_matmul": { + "srcs": ["rvv_matmul.cc"], + }, }, ) + +filegroup( + name = "rvv_mlop_tests", + srcs = [ + ":rvv_matmul_assembly.elf", + ":rvv_matmul.elf" + ], +) \ No newline at end of file
diff --git a/tests/cocotb/rvv/ml_ops/rvv_matmul.cc b/tests/cocotb/rvv/ml_ops/rvv_matmul.cc new file mode 100644 index 0000000..e828318 --- /dev/null +++ b/tests/cocotb/rvv/ml_ops/rvv_matmul.cc
@@ -0,0 +1,51 @@ +#include <riscv_vector.h> +#include <stdint.h> + +constexpr size_t kLhsRows = 16; +constexpr size_t kRhsCols = 16; +constexpr size_t kInner = 48; + +int8_t lhs_input[kLhsRows * kInner] __attribute__((section(".data"))) +__attribute__((aligned(16))); +int8_t rhs_input[kInner * kRhsCols] __attribute__((section(".data"))) +__attribute__((aligned(16))); +int32_t result_output[kLhsRows * kRhsCols] __attribute__((section(".data"))) +__attribute__((aligned(16))); + +// Assume rhs is column major. +void MatMul(size_t lhs_rows, size_t inner, size_t rhs_cols, const int8_t* lhs, + const int8_t* rhs, int32_t* result) { + const size_t vlenb = __riscv_vlenb(); + + for (size_t r = 0; r < lhs_rows; r++) { + const int8_t* lhs_data = lhs + (r * inner); + int32_t* result_row = result + (r * rhs_cols); + for (size_t c = 0; c < rhs_cols; c++) { + const int8_t* rhs_data = rhs + (c * inner); + // Reset accumulators + vint32m1_t vacc = __riscv_vmv_v_x_i32m1(0, 1); + + // Inner dot product loop + size_t k = 0; + size_t vl = vlenb; + while (k < inner) { + if (inner - k < vl) { + vl = inner - k; + } + // Load weights/activations + vint8m1_t vlhs_data = __riscv_vle8_v_i8m1(lhs_data + k, vl); + vint8m1_t vrhs_data = + __riscv_vle8_v_i8m1(rhs_data + k, vl); // input rhs is transposed + vint16m2_t vmul_16 = __riscv_vwmul_vv_i16m2(vlhs_data, vrhs_data, vl); + vacc = __riscv_vwredsum_vs_i16m2_i32m1(vmul_16, vacc, vlenb); + k += vl; + } + __riscv_vse32_v_i32m1(result_row + c, vacc, 1); + } + } +} + +int main() { + MatMul(kLhsRows, kInner, kRhsCols, lhs_input, rhs_input, result_output); + return 0; +} \ No newline at end of file
diff --git a/tests/cocotb/rvv/ml_ops/rvv_matmul_assembly.cc b/tests/cocotb/rvv/ml_ops/rvv_matmul_assembly.cc index d50981d..0f420c2 100644 --- a/tests/cocotb/rvv/ml_ops/rvv_matmul_assembly.cc +++ b/tests/cocotb/rvv/ml_ops/rvv_matmul_assembly.cc
@@ -3,20 +3,25 @@ constexpr size_t kLhsRows = 16; constexpr size_t kRhsCols = 16; -constexpr size_t kInner = 24; +constexpr size_t kInner = 48; -int8_t lhs_input[kLhsRows*kInner] __attribute__((section(".data"))) __attribute__((aligned(16))); -int8_t rhs_input[kInner*kRhsCols] __attribute__((section(".data"))) __attribute__((aligned(16))); -int32_t result_output[kLhsRows*kRhsCols] __attribute__((section(".data"))) __attribute__((aligned(16))); +int8_t lhs_input[kLhsRows * kInner] __attribute__((section(".data"))) +__attribute__((aligned(16))); +int8_t rhs_input[kInner * kRhsCols] __attribute__((section(".data"))) +__attribute__((aligned(16))); +int32_t result_output[kLhsRows * kRhsCols] __attribute__((section(".data"))) +__attribute__((aligned(16))); // Assume rhs is column major. -void MatMul(size_t lhs_rows, size_t inner, size_t rhs_cols, - const int8_t* lhs, const int8_t* rhs, int32_t* result) { +void MatMul(size_t lhs_rows, size_t inner, size_t rhs_cols, const int8_t* lhs, + const int8_t* rhs, int32_t* result) { const size_t vlenb = __riscv_vlenb(); // Create zero register for vredsum asm("vsetvli zero, %0, e32, m4, ta, ma;" - "vmv.v.i v0, 0;" : : "r" (vlenb)); + "vmv.v.i v0, 0;" + : + : "r"(vlenb)); for (size_t r = 0; r < lhs_rows; r++) { const int8_t* lhs_data = lhs + (r * inner); @@ -25,7 +30,7 @@ const int8_t* rhs_data = rhs + (c * inner); // Reset accumulators - asm("vsetvli zero, %0, e32, m4, ta, ma" : : "r" (vlenb)); + asm("vsetvli zero, %0, e32, m4, ta, ma" : : "r"(vlenb)); asm("vmv.v.i v8, 0"); // Inner dot product loop @@ -36,26 +41,32 @@ vl = inner - k; } // Load weights/activations - asm("vsetvli zero, %0, e8, m1, ta, ma" : : "r" (vl)); - asm("vle8.v v14, (%0)" : : "r" (lhs_data + k)); - asm("vle8.v v15, (%0)" : : "r" (rhs_data + k)); + asm("vsetvli zero, %0, e8, m1, ta, ma" : : "r"(vl)); + asm("vle8.v v14, (%0)" : : "r"(lhs_data + k)); + asm("vle8.v v15, (%0)" : : "r"(rhs_data + k)); // Multiply-accumulate asm("vsetvli zero, %0, e8, m1, ta, ma;" "vwmul.vv v12, v14, v15;" "vsetvli zero, %0, e16, m2, ta, ma;" - "vwadd.wv v8, v8, v12;" : : "r" (vl)); + "vwadd.wv v8, v8, v12;" + : + : "r"(vl)); k += vl; } // Reduction asm("vsetvli zero, %0, e32, m4, ta, ma;" - "vredsum.vs v8, v8, v0;" : : "r" (vlenb)); + "vredsum.vs v8, v8, v0;" + : + : "r"(vlenb)); // Store asm("vsetivli zero, 1, e32, m1, ta, ma;" - "vse32.v v8, (%0);" : : "r" (result_row + c)); + "vse32.v v8, (%0);" + : + : "r"(result_row + c)); } } }
diff --git a/tests/cocotb/rvv_ml_ops_cocotb_test.py b/tests/cocotb/rvv_ml_ops_cocotb_test.py index 22f8f8b..919ec97 100644 --- a/tests/cocotb/rvv_ml_ops_cocotb_test.py +++ b/tests/cocotb/rvv_ml_ops_cocotb_test.py
@@ -16,32 +16,33 @@ LHS_ROWS = 16 RHS_COLS = 16 - INNER = 24 + INNER = 48 fixture = await Fixture.Create(dut) r = runfiles.Create() - await fixture.load_elf_and_lookup_symbols( - r.Rlocation( - 'kelvin_hw/tests/cocotb/rvv/ml_ops/rvv_matmul_assembly.elf'), - ['lhs_input', 'rhs_input', 'result_output']) - np_type = np.int8 - min_value = np.iinfo(np_type).min - max_value = np.iinfo(np_type).max + 1 # One above. - lhs_data = np.random.randint(min_value, - max_value, [LHS_ROWS, INNER], - dtype=np_type) - rhs_data = np.random.randint(min_value, - max_value, [INNER, RHS_COLS], - dtype=np_type) - result_data = np.matmul(lhs_data.astype(np.int32), - rhs_data.astype(np.int32)) + elf_files = ['rvv_matmul.elf', 'rvv_matmul_assembly.elf'] + for elf_file in elf_files: - await fixture.write('lhs_input', lhs_data.flatten()) - await fixture.write('rhs_input', rhs_data.transpose().flatten()) + await fixture.load_elf_and_lookup_symbols( + r.Rlocation('kelvin_hw/tests/cocotb/rvv/ml_ops/' + elf_file), + ['lhs_input', 'rhs_input', 'result_output']) + np_type = np.int8 + min_value = np.iinfo(np_type).min + max_value = np.iinfo(np_type).max + 1 # One above. + lhs_data = np.random.randint(min_value, + max_value, [LHS_ROWS, INNER], + dtype=np_type) + rhs_data = np.random.randint(min_value, + max_value, [INNER, RHS_COLS], + dtype=np_type) + result_data = np.matmul(lhs_data.astype(np.int32), + rhs_data.astype(np.int32)) - await fixture.run_to_halt(timeout_cycles=1000000) + await fixture.write('lhs_input', lhs_data.flatten()) + await fixture.write('rhs_input', rhs_data.transpose().flatten()) + await fixture.run_to_halt(timeout_cycles=1000000) + output_matmul_result = (await fixture.read( + 'result_output', LHS_ROWS * RHS_COLS * + 4)).view(dtype=np.int32).reshape([LHS_ROWS, RHS_COLS]) - output_matmul_result = (await fixture.read( - 'result_output', LHS_ROWS * RHS_COLS * - 4)).view(dtype=np.int32).reshape([LHS_ROWS, RHS_COLS]) - assert ((result_data == output_matmul_result).all()) + assert ((result_data == output_matmul_result).all())