Fix wrong results for mixed prec indexed loads Adding a bunch of indexed load tests too. The fix is tested on data8 index16 loads below m1 and index32 below mf2. Pending another fix for larger lmul. TODO: test store TODO: test load data16/32 Change-Id: I83b986339ff0900449fa4b57fa2eb895ca4e26f8
diff --git a/hdl/chisel/src/kelvin/scalar/Lsu.scala b/hdl/chisel/src/kelvin/scalar/Lsu.scala index 832db87..e3ef588 100644 --- a/hdl/chisel/src/kelvin/scalar/Lsu.scala +++ b/hdl/chisel/src/kelvin/scalar/Lsu.scala
@@ -166,7 +166,12 @@ val pc = UInt(32.W) val addr = UInt(32.W) val data = UInt(32.W) // Doubles as rs2 + // This aligns with "width" in the spec. It controls index width in + // indexed loads/stores and data width otherwise. val elemWidth = Option.when(p.enableRvv) { UInt(3.W) } + // This is the sew from vtype. It controls data width in indexed + // loads/stores and is unused in other ops. + val sew = Option.when(p.enableRvv) { UInt(3.W) } val lmul = Option.when(p.enableRvv) { UInt(4.W) } val nfields = Option.when(p.enableRvv) { UInt(3.W) } @@ -207,6 +212,7 @@ result.lmul.get := 1.U(1.W) << effectiveLmul // If mask operation, force fields to zero result.nfields.get := Mux(cmd.isMaskOperation(), 0.U, cmd.nfields.get) + result.sew.get := rvvState.get.bits.sew } result @@ -235,22 +241,35 @@ def apply(bytesPerSlot: Int, baseAddr: UInt, indices: UInt, - elemWidth: UInt): Vec[UInt] = { + indexWidth: UInt, + sew: UInt): Vec[UInt] = { val indices8 = UIntToVec(indices, 8) - val indices16 = UIntToVec(indices, 16) - val indices32 = UIntToVec(indices, 32) + val indices16 = UIntToVec(indices, 16).padTo(bytesPerSlot, 0.U) + val indices32 = UIntToVec(indices, 32).padTo(bytesPerSlot, 0.U) + + val indices_v = MuxCase(VecInit.fill(bytesPerSlot)(0.U(32.W)), Seq( + // 8-bit indices. + (indexWidth === "b000".U) -> VecInit((0 until bytesPerSlot).map( + i => Cat(0.U(24.W), indices8(i)))), + // 16-bit indices. + (indexWidth === "b101".U) -> VecInit((0 until bytesPerSlot).map( + i => Cat(0.U(16.W), indices16(i)))), + // 32-bit indices. + (indexWidth === "b110".U) -> VecInit((0 until bytesPerSlot).map( + i => indices32(i))), + )) 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))), + (sew === "b000".U) -> VecInit((0 until bytesPerSlot).map( + i => (baseAddr + indices_v(i)))), // 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)), + (sew === "b001".U) -> VecInit((0 until bytesPerSlot).map( + i => (baseAddr + indices_v(i >> 1) + (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)) + (sew === "b010".U) -> VecInit((0 until bytesPerSlot).map( + i => (baseAddr + indices_v(i >> 2) + (i & 3).U))) )) } } @@ -273,7 +292,12 @@ val lmul = UInt(4.W) val elemStride = UInt(32.W) // Stride between lanes in a vector val segmentStride = UInt(32.W) // Stride between base addr between segments + // This aligns with "width" in the spec. It controls index width in + // indexed loads/stores and data width otherwise. val elemWidth = UInt(3.W) + // This controls data width in indexed loads/stores and is unused in + // other ops. + val sew = UInt(3.W) val nfields = UInt(3.W) val segment = UInt(3.W) // Add this to find the next segment. @@ -339,9 +363,10 @@ op.isOneOf(LsuOp.VLOAD_OINDEXED, LsuOp.VLOAD_UINDEXED, LsuOp.VSTORE_OINDEXED, LsuOp.VSTORE_UINDEXED) -> ComputeIndexedAddrs(bytesPerSlot, baseAddr, rvv2lsu.idx.bits.data, - elemWidth) + elemWidth, sew), )) result.elemWidth := elemWidth + result.sew := sew result.data := Mux(updated && LsuOp.isVector(op) && rvv2lsu.vregfile.valid, UIntToVec(rvv2lsu.vregfile.bits.data, 8), data) @@ -386,6 +411,7 @@ result.elemStride := elemStride result.segmentStride := segmentStride result.elemWidth := elemWidth + result.sew := sew result.nfields := nfields result.segment := segment result.nextSegmentVectorOffset := nextSegmentVectorOffset @@ -412,6 +438,7 @@ result.elemStride := elemStride result.segmentStride := segmentStride result.elemWidth := elemWidth + result.sew := sew result.nfields := nfields result.nextSegmentVectorOffset := nextSegmentVectorOffset result.nextLmulVectorRewind := nextLmulVectorRewind @@ -492,6 +519,7 @@ result.elemStride := elemStride result.segmentStride := segmentStride result.elemWidth := elemWidth + result.sew := sew result.nfields := nfields result.segment := segment result.nextSegmentVectorOffset := nextSegmentVectorOffset @@ -571,6 +599,7 @@ // Compute addrs result.baseAddr := uop.addr result.elemWidth := uop.elemWidth.getOrElse(0.U(3.W)) + result.sew := uop.sew.getOrElse(0.U(3.W)) result.addrs := Mux( uop.op.isOneOf(LsuOp.VLOAD_STRIDED, LsuOp.VSTORE_STRIDED), ComputeStridedAddrs(bytesPerSlot, uop.addr, uop.data, uop.elemWidth.getOrElse(0.U(3.W))),
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD index ee7bec8..e0d7820 100644 --- a/tests/cocotb/BUILD +++ b/tests/cocotb/BUILD
@@ -173,6 +173,9 @@ # BEGIN_TESTCASES_FOR_rvv_load_store_test RVV_LOAD_STORE_TESTCASES = [ "load_store_bits", + "load8_index8", + "load8_index16", + "load8_index32", "load8_seg_unit", "load8_stride2_m1", "load8_stride2_m1_partial", @@ -189,7 +192,6 @@ "load_store32_unit_m2", "load8_segment2_stride6_m1", "load16_segment2_stride6_m1", - "load8_indexed_m1", "store8_indexed_m1", "store8_seg_unit", "store16_seg_unit", @@ -377,6 +379,7 @@ "deps": [ "//kelvin_test_utils:sim_test_fixture", "@bazel_tools//tools/python/runfiles", + requirement("tqdm"), ], "data": ["//tests/cocotb/rvv/load_store:rvv_load_store_tests"], "size": "large",
diff --git a/tests/cocotb/rvv/load_store/BUILD b/tests/cocotb/rvv/load_store/BUILD index c5a79d3..4e07588 100644 --- a/tests/cocotb/rvv/load_store/BUILD +++ b/tests/cocotb/rvv/load_store/BUILD
@@ -23,8 +23,14 @@ "load_store_bits": { "srcs": ["load_store_bits.cc"], }, - "load8_indexed_m1": { - "srcs": ["load8_indexed_m1.cc"], + "load8_index8": { + "srcs": ["load8_index8.cc"], + }, + "load8_index16": { + "srcs": ["load8_index16.cc"], + }, + "load8_index32": { + "srcs": ["load8_index32.cc"], }, "load8_seg_unit": { "srcs": ["load8_seg_unit.cc"], @@ -96,7 +102,9 @@ name = "rvv_load_store_tests", srcs = [ ":load_store_bits.elf", - ":load8_indexed_m1.elf", + ":load8_index8.elf", + ":load8_index16.elf", + ":load8_index32.elf", ":load8_seg_unit.elf", ":load8_segment2_stride6_m1.elf", ":load16_segment2_stride6_m1.elf",
diff --git a/tests/cocotb/rvv/load_store/load8_index16.cc b/tests/cocotb/rvv/load_store/load8_index16.cc new file mode 100644 index 0000000..f6decff --- /dev/null +++ b/tests/cocotb/rvv/load_store/load8_index16.cc
@@ -0,0 +1,101 @@ +// 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> + +namespace { +constexpr size_t lut_size = 32000; // DTCM is 32KB. +// Double sized so we can check trailing regions are not read/written. +constexpr size_t buf_size = 128; +} // namespace + +size_t vl __attribute__((section(".data"))) = 8; +// Indices are always unsigned. +uint16_t index_buf[buf_size] __attribute__((section(".data"))); +// These instructions don't differentiate signed/unsigned so we only need to +// test one. The types come from intrinsic level. +uint8_t in_buf[lut_size] __attribute__((section(".data"))); +uint8_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vluxei16_v_u8mf4() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + auto data = __riscv_vluxei16_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei16_v_u8mf2() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + auto data = __riscv_vluxei16_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei16_v_u8m1() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + auto data = __riscv_vluxei16_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei16_v_u8m2() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + auto data = __riscv_vluxei16_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei16_v_u8m4() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + auto data = __riscv_vluxei16_v_u8m4(in_buf, indices, vl); + __riscv_vse8_v_u8m4(out_buf, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vloxei16_v_u8mf4() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + auto data = __riscv_vloxei16_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei16_v_u8mf2() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + auto data = __riscv_vloxei16_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei16_v_u8m1() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + auto data = __riscv_vloxei16_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei16_v_u8m2() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + auto data = __riscv_vloxei16_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei16_v_u8m4() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + auto data = __riscv_vloxei16_v_u8m4(in_buf, indices, vl); + __riscv_vse8_v_u8m4(out_buf, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vluxei16_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/load8_index32.cc b/tests/cocotb/rvv/load_store/load8_index32.cc new file mode 100644 index 0000000..524e86d --- /dev/null +++ b/tests/cocotb/rvv/load_store/load8_index32.cc
@@ -0,0 +1,89 @@ +// 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> + +namespace { +constexpr size_t lut_size = 32000; // DTCM is 32KB. +// Double sized so we can check trailing regions are not read/written. +constexpr size_t buf_size = 64; +} // namespace + +size_t vl __attribute__((section(".data"))) = 8; +// Indices are always unsigned. +uint32_t index_buf[buf_size] __attribute__((section(".data"))); +// These instructions don't differentiate signed/unsigned so we only need to +// test one. The types come from intrinsic level. +uint8_t in_buf[lut_size] __attribute__((section(".data"))); +uint8_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vluxei32_v_u8mf4() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + auto data = __riscv_vluxei32_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei32_v_u8mf2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + auto data = __riscv_vluxei32_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei32_v_u8m1() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + auto data = __riscv_vluxei32_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei32_v_u8m2() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + auto data = __riscv_vluxei32_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vloxei32_v_u8mf4() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + auto data = __riscv_vloxei32_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei32_v_u8mf2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + auto data = __riscv_vloxei32_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei32_v_u8m1() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + auto data = __riscv_vloxei32_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei32_v_u8m2() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + auto data = __riscv_vloxei32_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vluxei32_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/load8_index8.cc b/tests/cocotb/rvv/load_store/load8_index8.cc new file mode 100644 index 0000000..6a6d75f --- /dev/null +++ b/tests/cocotb/rvv/load_store/load8_index8.cc
@@ -0,0 +1,113 @@ +// 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> + +namespace { +constexpr size_t lut_size = 256; +// Double sized so we can check trailing regions are not read/written. +constexpr size_t buf_size = 256; +} // namespace + +size_t vl __attribute__((section(".data"))) = 16; +// Indices are always unsigned. +uint8_t index_buf[buf_size] __attribute__((section(".data"))); +// These instructions don't differentiate signed/unsigned so we only need to +// test one. The types come from intrinsic level. +uint8_t in_buf[lut_size] __attribute__((section(".data"))); +uint8_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vluxei8_v_u8mf4() { + auto indices = __riscv_vle8_v_u8mf4(index_buf, vl); + auto data = __riscv_vluxei8_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei8_v_u8mf2() { + auto indices = __riscv_vle8_v_u8mf2(index_buf, vl); + auto data = __riscv_vluxei8_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei8_v_u8m1() { + auto indices = __riscv_vle8_v_u8m1(index_buf, vl); + auto data = __riscv_vluxei8_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei8_v_u8m2() { + auto indices = __riscv_vle8_v_u8m2(index_buf, vl); + auto data = __riscv_vluxei8_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei8_v_u8m4() { + auto indices = __riscv_vle8_v_u8m4(index_buf, vl); + auto data = __riscv_vluxei8_v_u8m4(in_buf, indices, vl); + __riscv_vse8_v_u8m4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vluxei8_v_u8m8() { + auto indices = __riscv_vle8_v_u8m8(index_buf, vl); + auto data = __riscv_vluxei8_v_u8m8(in_buf, indices, vl); + __riscv_vse8_v_u8m8(out_buf, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vloxei8_v_u8mf4() { + auto indices = __riscv_vle8_v_u8mf4(index_buf, vl); + auto data = __riscv_vloxei8_v_u8mf4(in_buf, indices, vl); + __riscv_vse8_v_u8mf4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei8_v_u8mf2() { + auto indices = __riscv_vle8_v_u8mf2(index_buf, vl); + auto data = __riscv_vloxei8_v_u8mf2(in_buf, indices, vl); + __riscv_vse8_v_u8mf2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei8_v_u8m1() { + auto indices = __riscv_vle8_v_u8m1(index_buf, vl); + auto data = __riscv_vloxei8_v_u8m1(in_buf, indices, vl); + __riscv_vse8_v_u8m1(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei8_v_u8m2() { + auto indices = __riscv_vle8_v_u8m2(index_buf, vl); + auto data = __riscv_vloxei8_v_u8m2(in_buf, indices, vl); + __riscv_vse8_v_u8m2(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei8_v_u8m4() { + auto indices = __riscv_vle8_v_u8m4(index_buf, vl); + auto data = __riscv_vloxei8_v_u8m4(in_buf, indices, vl); + __riscv_vse8_v_u8m4(out_buf, data, vl); +} + +__attribute__((used, retain)) void vloxei8_v_u8m8() { + auto indices = __riscv_vle8_v_u8m8(index_buf, vl); + auto data = __riscv_vloxei8_v_u8m8(in_buf, indices, vl); + __riscv_vse8_v_u8m8(out_buf, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vluxei8_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/load8_indexed_m1.cc b/tests/cocotb/rvv/load_store/load8_indexed_m1.cc deleted file mode 100644 index 8c6c9a7..0000000 --- a/tests/cocotb/rvv/load_store/load8_indexed_m1.cc +++ /dev/null
@@ -1,28 +0,0 @@ -// 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_test.py b/tests/cocotb/rvv_load_store_test.py index 2f8d9d3..843e6e2 100644 --- a/tests/cocotb/rvv_load_store_test.py +++ b/tests/cocotb/rvv_load_store_test.py
@@ -14,6 +14,7 @@ import cocotb import numpy as np +import tqdm from bazel_tools.tools.python.runfiles import runfiles from kelvin_test_utils.sim_test_fixture import Fixture @@ -75,13 +76,14 @@ r = runfiles.Create() await fixture.load_elf_and_lookup_symbols( r.Rlocation('kelvin_hw/tests/cocotb/rvv/load_store/' + elf_name), - ['impl', 'vl', 'in_buf', 'out_buf'] + [c['impl'] for c in cases], + ['impl', 'vl', 'in_buf', 'out_buf'] + + list(set([c['impl'] for c in cases])), ) min_value = np.iinfo(dtype).min max_value = np.iinfo(dtype).max + 1 # One above. rng = np.random.default_rng() - for c in cases: + for c in tqdm.tqdm(cases): impl = c['impl'] vl = c['vl'] in_size = c['in_size'] @@ -111,10 +113,13 @@ }) assert (actual_outputs == expected_outputs).all(), debug_msg + async def vector_load_indexed( dut, elf_name: str, + cases: list[dict], # keys: impl, vl, in_size, out_size. dtype, + index_dtype, ): """RVV load-store test template for indexed loads. @@ -124,32 +129,47 @@ r = runfiles.Create() await fixture.load_elf_and_lookup_symbols( r.Rlocation('kelvin_hw/tests/cocotb/rvv/load_store/' + elf_name), - ['input_indices', 'input_data', 'output_data'], + ['impl', 'vl', 'in_buf', 'out_buf', 'index_buf'] + + list(set([c['impl'] for c in cases])), ) - 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) + for c in tqdm.tqdm(cases): + impl = c['impl'] + vl = c['vl'] + in_size = c['in_size'] + out_size = c['out_size'] - expected_outputs = np.take(input_data, input_indices) + # TODO(davidgao): currently assuming the vl is supported. + # We'll eventually want to test unsupported vl. + indices = rng.integers(0, in_size, vl, dtype=index_dtype) + input_data = rng.integers(min_value, max_value, in_size, dtype=dtype) + expected_outputs = input_data[indices[:vl]] + sbz = np.zeros(out_size - vl, dtype=dtype) + expected_outputs = np.concat((expected_outputs, sbz)) - 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.write_ptr('impl', impl) + await fixture.write_word('vl', vl) + await fixture.write('index_buf', indices) + await fixture.write('in_buf', input_data) + await fixture.write('out_buf', np.zeros([out_size], dtype=dtype)) - await fixture.run_to_halt() + await fixture.run_to_halt() - actual_outputs = (await fixture.read( - 'output_data', out_data_count * np.dtype(dtype).itemsize)).view(dtype) + actual_outputs = (await fixture.read( + 'out_buf', out_size * np.dtype(dtype).itemsize)).view(dtype) - assert (actual_outputs == expected_outputs).all() + debug_msg = str({ + 'impl': impl, + 'input': input_data, + 'indices': indices, + 'expected': expected_outputs, + 'actual': actual_outputs, + }) + assert (actual_outputs == expected_outputs).all(), debug_msg + async def vector_store_indexed( dut, @@ -223,7 +243,8 @@ await fixture.load_elf_and_lookup_symbols( r.Rlocation( 'kelvin_hw/tests/cocotb/rvv/load_store/load_store_bits.elf'), - ['vl', 'in_buf', 'out_buf', 'impl'] + [c['impl'] for c in cases], + ['vl', 'in_buf', 'out_buf', 'impl'] + + list(set([c['impl'] for c in cases])), ) rng = np.random.default_rng() for c in cases: @@ -460,6 +481,139 @@ @cocotb.test() +async def load8_index8(dut): + """Test vl*xei8_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'in_size': 256, + 'out_size': vl * 2, + } + + await vector_load_indexed( + dut = dut, + elf_name = 'load8_index8.elf', + cases = [ + # Unordered + make_test_case('vluxei8_v_u8mf4', vl = 4), + make_test_case('vluxei8_v_u8mf4', vl = 3), + make_test_case('vluxei8_v_u8mf2', vl = 8), + make_test_case('vluxei8_v_u8mf2', vl = 7), + make_test_case('vluxei8_v_u8m1', vl = 16), + make_test_case('vluxei8_v_u8m1', vl = 15), + make_test_case('vluxei8_v_u8m2', vl = 32), + make_test_case('vluxei8_v_u8m2', vl = 31), + make_test_case('vluxei8_v_u8m4', vl = 64), + make_test_case('vluxei8_v_u8m4', vl = 63), + make_test_case('vluxei8_v_u8m8', vl = 128), + make_test_case('vluxei8_v_u8m8', vl = 127), + # Ordered + make_test_case('vloxei8_v_u8mf4', vl = 4), + make_test_case('vloxei8_v_u8mf4', vl = 3), + make_test_case('vloxei8_v_u8mf2', vl = 8), + make_test_case('vloxei8_v_u8mf2', vl = 7), + make_test_case('vloxei8_v_u8m1', vl = 16), + make_test_case('vloxei8_v_u8m1', vl = 15), + make_test_case('vloxei8_v_u8m2', vl = 32), + make_test_case('vloxei8_v_u8m2', vl = 31), + make_test_case('vloxei8_v_u8m4', vl = 64), + make_test_case('vloxei8_v_u8m4', vl = 63), + make_test_case('vloxei8_v_u8m8', vl = 128), + make_test_case('vloxei8_v_u8m8', vl = 127), + ], + dtype = np.uint8, + index_dtype = np.uint8, + ) + + +@cocotb.test() +async def load8_index16(dut): + """Test vl*xei16_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'in_size': 32000, # DTCM is 32KB + 'out_size': vl * 2, + } + + await vector_load_indexed( + dut = dut, + elf_name = 'load8_index16.elf', + cases = [ + # Unordered + make_test_case('vluxei16_v_u8mf4', vl = 4), + make_test_case('vluxei16_v_u8mf4', vl = 3), + make_test_case('vluxei16_v_u8mf2', vl = 8), + make_test_case('vluxei16_v_u8mf2', vl = 7), + # make_test_case('vluxei16_v_u8m1', vl = 16), + # make_test_case('vluxei16_v_u8m1', vl = 15), + # make_test_case('vluxei16_v_u8m2', vl = 32), + # make_test_case('vluxei16_v_u8m2', vl = 31), + # make_test_case('vluxei16_v_u8m4', vl = 64), + # make_test_case('vluxei16_v_u8m4', vl = 63), + # Ordered + make_test_case('vloxei16_v_u8mf4', vl = 4), + make_test_case('vloxei16_v_u8mf4', vl = 3), + make_test_case('vloxei16_v_u8mf2', vl = 8), + make_test_case('vloxei16_v_u8mf2', vl = 7), + # make_test_case('vloxei16_v_u8m1', vl = 16), + # make_test_case('vloxei16_v_u8m1', vl = 15), + # make_test_case('vloxei16_v_u8m2', vl = 32), + # make_test_case('vloxei16_v_u8m2', vl = 31), + # make_test_case('vloxei16_v_u8m4', vl = 64), + # make_test_case('vloxei16_v_u8m4', vl = 63), + ], + dtype = np.uint8, + index_dtype = np.uint16, + ) + + +@cocotb.test() +async def load8_index32(dut): + """Test vl*xei32_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'in_size': 32000, # DTCM is 32KB + 'out_size': vl * 2, + } + + await vector_load_indexed( + dut = dut, + elf_name = 'load8_index32.elf', + cases = [ + # Unordered + make_test_case('vluxei32_v_u8mf4', vl = 4), + make_test_case('vluxei32_v_u8mf4', vl = 3), + # make_test_case('vluxei32_v_u8mf2', vl = 8), + # make_test_case('vluxei32_v_u8mf2', vl = 7), + # make_test_case('vluxei32_v_u8m1', vl = 16), + # make_test_case('vluxei32_v_u8m1', vl = 15), + # make_test_case('vluxei32_v_u8m2', vl = 32), + # make_test_case('vluxei32_v_u8m2', vl = 31), + # make_test_case('vluxei32_v_u8m4', vl = 64), + # make_test_case('vluxei32_v_u8m4', vl = 63), + # Ordered + make_test_case('vloxei32_v_u8mf4', vl = 4), + make_test_case('vloxei32_v_u8mf4', vl = 3), + # make_test_case('vloxei32_v_u8mf2', vl = 8), + # make_test_case('vloxei32_v_u8mf2', vl = 7), + # make_test_case('vloxei32_v_u8m1', vl = 16), + # make_test_case('vloxei32_v_u8m1', vl = 15), + # make_test_case('vloxei32_v_u8m2', vl = 32), + # make_test_case('vloxei32_v_u8m2', vl = 31), + # make_test_case('vloxei32_v_u8m4', vl = 64), + # make_test_case('vloxei32_v_u8m4', vl = 63), + ], + dtype = np.uint8, + index_dtype = np.uint32, + ) + + +@cocotb.test() async def load16_seg_unit(dut): """Test vlseg*e16 usage accessible from intrinsics.""" def make_test_case(impl: str, vl: int, n_segs: int): @@ -602,14 +756,6 @@ @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,