Expand tests for indexed store of same ddtype and idtype. Change-Id: I4d00fd837c8b18308d7bac0bd273403c76c3c213
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD index 7657acc..0888e46 100644 --- a/tests/cocotb/BUILD +++ b/tests/cocotb/BUILD
@@ -194,7 +194,9 @@ "load_store32_unit_m2", "load8_segment2_stride6_m1", "load16_segment2_stride6_m1", - "store8_indexed_m1", + "store8_index8", + "store16_index16", + "store32_index32", "store8_seg_unit", "store16_seg_unit", "store32_seg_unit",
diff --git a/tests/cocotb/rvv/load_store/BUILD b/tests/cocotb/rvv/load_store/BUILD index e0e3b4a..e220ce9 100644 --- a/tests/cocotb/rvv/load_store/BUILD +++ b/tests/cocotb/rvv/load_store/BUILD
@@ -89,8 +89,14 @@ "load32_stride8_m1_partial": { "srcs": ["load32_stride8_m1_partial.cc"], }, - "store8_indexed_m1": { - "srcs": ["store8_indexed_m1.cc"], + "store8_index8": { + "srcs": ["store8_index8.cc"], + }, + "store16_index16": { + "srcs": ["store16_index16.cc"], + }, + "store32_index32": { + "srcs": ["store32_index32.cc"], }, "store8_seg_unit": { "srcs": ["store8_seg_unit.cc"], @@ -130,7 +136,9 @@ ":load32_seg_unit.elf", ":load32_stride8_m1.elf", ":load32_stride8_m1_partial.elf", - ":store8_indexed_m1.elf", + ":store8_index8.elf", + ":store16_index16.elf", + ":store32_index32.elf", ":store8_seg_unit", ":store16_seg_unit", ":store32_seg_unit",
diff --git a/tests/cocotb/rvv/load_store/store16_index16.cc b/tests/cocotb/rvv/load_store/store16_index16.cc new file mode 100644 index 0000000..6487165 --- /dev/null +++ b/tests/cocotb/rvv/load_store/store16_index16.cc
@@ -0,0 +1,100 @@ +// 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 scatter_count = 64; +constexpr size_t buf_size = 16000; +} // namespace + +size_t vl __attribute__((section(".data"))) = 8; +// Indices are always unsigned. +uint16_t index_buf[scatter_count] __attribute__((section(".data"))); +// These instructions don't differentiate signed/unsigned so we only need to +// test one. The types come from intrinsic level. +uint16_t in_buf[scatter_count] __attribute__((section(".data"))); +uint16_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vsuxei16_v_u16mf2() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + auto data = __riscv_vle16_v_u16mf2(in_buf, vl); + __riscv_vsuxei16_v_u16mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u16m1() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + auto data = __riscv_vle16_v_u16m1(in_buf, vl); + __riscv_vsuxei16_v_u16m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u16m2() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + auto data = __riscv_vle16_v_u16m2(in_buf, vl); + __riscv_vsuxei16_v_u16m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u16m4() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + auto data = __riscv_vle16_v_u16m4(in_buf, vl); + __riscv_vsuxei16_v_u16m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u16m8() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + auto data = __riscv_vle16_v_u16m8(in_buf, vl); + __riscv_vsuxei16_v_u16m8(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei16_v_u16mf2() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + auto data = __riscv_vle16_v_u16mf2(in_buf, vl); + __riscv_vsoxei16_v_u16mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u16m1() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + auto data = __riscv_vle16_v_u16m1(in_buf, vl); + __riscv_vsoxei16_v_u16m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u16m2() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + auto data = __riscv_vle16_v_u16m2(in_buf, vl); + __riscv_vsoxei16_v_u16m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u16m4() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + auto data = __riscv_vle16_v_u16m4(in_buf, vl); + __riscv_vsoxei16_v_u16m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u16m8() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + auto data = __riscv_vle16_v_u16m8(in_buf, vl); + __riscv_vsoxei16_v_u16m8(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei16_v_u16m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/store32_index32.cc b/tests/cocotb/rvv/load_store/store32_index32.cc new file mode 100644 index 0000000..07001f5 --- /dev/null +++ b/tests/cocotb/rvv/load_store/store32_index32.cc
@@ -0,0 +1,88 @@ +// 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 scatter_count = 32; +constexpr size_t buf_size = 8000; +} // namespace + +size_t vl __attribute__((section(".data"))) = 8; +// Indices are always unsigned. +uint32_t index_buf[scatter_count] __attribute__((section(".data"))); +// These instructions don't differentiate signed/unsigned so we only need to +// test one. The types come from intrinsic level. +uint32_t in_buf[scatter_count] __attribute__((section(".data"))); +uint32_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vsuxei32_v_u32m1() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + auto data = __riscv_vle32_v_u32m1(in_buf, vl); + __riscv_vsuxei32_v_u32m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u32m2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + auto data = __riscv_vle32_v_u32m2(in_buf, vl); + __riscv_vsuxei32_v_u32m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u32m4() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + auto data = __riscv_vle32_v_u32m4(in_buf, vl); + __riscv_vsuxei32_v_u32m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u32m8() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + auto data = __riscv_vle32_v_u32m8(in_buf, vl); + __riscv_vsuxei32_v_u32m8(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei32_v_u32m1() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + auto data = __riscv_vle32_v_u32m1(in_buf, vl); + __riscv_vsoxei32_v_u32m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u32m2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + auto data = __riscv_vle32_v_u32m2(in_buf, vl); + __riscv_vsoxei32_v_u32m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u32m4() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + auto data = __riscv_vle32_v_u32m4(in_buf, vl); + __riscv_vsoxei32_v_u32m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u32m8() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + auto data = __riscv_vle32_v_u32m8(in_buf, vl); + __riscv_vsoxei32_v_u32m8(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei32_v_u32m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/store8_index8.cc b/tests/cocotb/rvv/load_store/store8_index8.cc new file mode 100644 index 0000000..6d2a7ac --- /dev/null +++ b/tests/cocotb/rvv/load_store/store8_index8.cc
@@ -0,0 +1,112 @@ +// 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 scatter_count = 128; +constexpr size_t buf_size = 512; +} // namespace + +size_t vl __attribute__((section(".data"))) = 16; +// Indices are always unsigned. +uint8_t index_buf[scatter_count] __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[scatter_count] __attribute__((section(".data"))); +uint8_t out_buf[buf_size] __attribute__((section(".data"))); + +extern "C" { +// Unordered +__attribute__((used, retain)) void vsuxei8_v_u8mf4() { + auto indices = __riscv_vle8_v_u8mf4(index_buf, vl); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsuxei8_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei8_v_u8mf2() { + auto indices = __riscv_vle8_v_u8mf2(index_buf, vl); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsuxei8_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei8_v_u8m1() { + auto indices = __riscv_vle8_v_u8m1(index_buf, vl); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsuxei8_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei8_v_u8m2() { + auto indices = __riscv_vle8_v_u8m2(index_buf, vl); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsuxei8_v_u8m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei8_v_u8m4() { + auto indices = __riscv_vle8_v_u8m4(index_buf, vl); + auto data = __riscv_vle8_v_u8m4(in_buf, vl); + __riscv_vsuxei8_v_u8m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei8_v_u8m8() { + auto indices = __riscv_vle8_v_u8m8(index_buf, vl); + auto data = __riscv_vle8_v_u8m8(in_buf, vl); + __riscv_vsuxei8_v_u8m8(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei8_v_u8mf4() { + auto indices = __riscv_vle8_v_u8mf4(index_buf, vl); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsoxei8_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei8_v_u8mf2() { + auto indices = __riscv_vle8_v_u8mf2(index_buf, vl); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsoxei8_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei8_v_u8m1() { + auto indices = __riscv_vle8_v_u8m1(index_buf, vl); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsoxei8_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei8_v_u8m2() { + auto indices = __riscv_vle8_v_u8m2(index_buf, vl); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsoxei8_v_u8m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei8_v_u8m4() { + auto indices = __riscv_vle8_v_u8m4(index_buf, vl); + auto data = __riscv_vle8_v_u8m4(in_buf, vl); + __riscv_vsoxei8_v_u8m4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei8_v_u8m8() { + auto indices = __riscv_vle8_v_u8m8(index_buf, vl); + auto data = __riscv_vle8_v_u8m8(in_buf, vl); + __riscv_vsoxei8_v_u8m8(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei8_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/store8_indexed_m1.cc b/tests/cocotb/rvv/load_store/store8_indexed_m1.cc deleted file mode 100644 index 221ff76..0000000 --- a/tests/cocotb/rvv/load_store/store8_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[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 ac9ab84..f7c6aa6 100644 --- a/tests/cocotb/rvv_load_store_test.py +++ b/tests/cocotb/rvv_load_store_test.py
@@ -180,7 +180,9 @@ async def vector_store_indexed( dut, elf_name: str, - dtype, + cases: list[dict], # keys: impl, vl, out_size. + data_dtype, + index_dtype, ): """RVV load-store test template for indexed stores. @@ -190,36 +192,57 @@ 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({c['impl'] for c in cases}), ) - 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) + for c in tqdm.tqdm(cases): + impl = c['impl'] + vl = c['vl'] + out_size = c['out_size'] + in_bytes = np.dtype(data_dtype).itemsize * vl - await fixture.write('input_data', input_data) - await fixture.write('input_indices', input_indices) - await fixture.write('output_data', original_outputs) + # Don't go beyond the buffer. + index_max = min(np.iinfo(index_dtype).max, out_size) + index_max = index_max - np.dtype(index_dtype).itemsize + # TODO(davidgao): currently assuming the vl is supported. + # We'll eventually want to test unsupported vl. + indices = rng.integers(0, index_max+1, vl, dtype=index_dtype) + # Index is in bytes so input needs to be in bytes. + input_data = rng.integers(0, np.iinfo(data_dtype).max + 1, vl, + dtype=data_dtype) + output_data = np.zeros(out_size, dtype=data_dtype) - expected_outputs = np.copy(original_outputs) - for idx, data in zip(input_indices, input_data): - expected_outputs[idx] = data + 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', output_data) - 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(data_dtype).itemsize)).view(data_dtype) - assert (actual_outputs == expected_outputs).all() + # Compute expected outputs. Note that indices are in bytes for all stores. + expected_outputs = output_data.copy().view(np.uint8) + elem_size = np.dtype(data_dtype).itemsize + input_data_bytes = input_data.view(np.uint8) + indices_bytes = np.repeat(indices, elem_size) + \ + np.tile(np.arange(elem_size), len(indices)) + np.put_along_axis(expected_outputs, indices_bytes, input_data_bytes, None) + expected_outputs = expected_outputs.view(data_dtype) + + debug_msg = str({ + 'impl': impl, + 'input': input_data, + 'indices': indices, + 'expected': expected_outputs, + 'actual': actual_outputs, + }) + + assert (actual_outputs == expected_outputs).all(), debug_msg @cocotb.test() @@ -840,11 +863,128 @@ @cocotb.test() -async def store8_indexed_m1(dut): +async def store8_index8(dut): + """Test vs*xei8_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'out_size': 512, + } + await vector_store_indexed( dut = dut, - elf_name = 'store8_indexed_m1.elf', - dtype = np.uint8, + elf_name = 'store8_index8.elf', + cases = [ + # Unordered + make_test_case('vsuxei8_v_u8mf4', vl = 4), + make_test_case('vsuxei8_v_u8mf4', vl = 3), + make_test_case('vsuxei8_v_u8mf2', vl = 8), + make_test_case('vsuxei8_v_u8mf2', vl = 7), + make_test_case('vsuxei8_v_u8m1', vl = 16), + make_test_case('vsuxei8_v_u8m1', vl = 15), + make_test_case('vsuxei8_v_u8m2', vl = 32), + make_test_case('vsuxei8_v_u8m2', vl = 31), + make_test_case('vsuxei8_v_u8m4', vl = 64), + make_test_case('vsuxei8_v_u8m4', vl = 63), + make_test_case('vsuxei8_v_u8m8', vl = 128), + make_test_case('vsuxei8_v_u8m8', vl = 127), + # Ordered + make_test_case('vsoxei8_v_u8mf2', vl = 4), + make_test_case('vsoxei8_v_u8mf2', vl = 3), + make_test_case('vsoxei8_v_u8mf2', vl = 8), + make_test_case('vsoxei8_v_u8mf2', vl = 7), + make_test_case('vsoxei8_v_u8m1', vl = 16), + make_test_case('vsoxei8_v_u8m1', vl = 15), + make_test_case('vsoxei8_v_u8m2', vl = 32), + make_test_case('vsoxei8_v_u8m2', vl = 31), + make_test_case('vsoxei8_v_u8m4', vl = 64), + make_test_case('vsoxei8_v_u8m4', vl = 63), + make_test_case('vsoxei8_v_u8m8', vl = 128), + make_test_case('vsoxei8_v_u8m8', vl = 127), + ], + data_dtype = np.uint8, + index_dtype = np.uint8, + ) + + +@cocotb.test() +async def store16_index16(dut): + """Test vs*xei16_v_u16 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'out_size': 16000, + } + + await vector_store_indexed( + dut = dut, + elf_name = 'store16_index16.elf', + cases = [ + # Unordered + make_test_case('vsuxei16_v_u16mf2', vl = 4), + make_test_case('vsuxei16_v_u16mf2', vl = 3), + make_test_case('vsuxei16_v_u16m1', vl = 8), + make_test_case('vsuxei16_v_u16m1', vl = 7), + make_test_case('vsuxei16_v_u16m2', vl = 16), + make_test_case('vsuxei16_v_u16m2', vl = 15), + make_test_case('vsuxei16_v_u16m4', vl = 32), + make_test_case('vsuxei16_v_u16m4', vl = 31), + make_test_case('vsuxei16_v_u16m8', vl = 64), + make_test_case('vsuxei16_v_u16m8', vl = 63), + # Ordered + make_test_case('vsoxei16_v_u16mf2', vl = 4), + make_test_case('vsoxei16_v_u16mf2', vl = 3), + make_test_case('vsoxei16_v_u16m1', vl = 8), + make_test_case('vsoxei16_v_u16m1', vl = 7), + make_test_case('vsoxei16_v_u16m2', vl = 16), + make_test_case('vsoxei16_v_u16m2', vl = 15), + make_test_case('vsoxei16_v_u16m4', vl = 32), + make_test_case('vsoxei16_v_u16m4', vl = 31), + make_test_case('vsoxei16_v_u16m8', vl = 64), + make_test_case('vsoxei16_v_u16m8', vl = 63), + ], + data_dtype = np.uint16, + index_dtype = np.uint16, + ) + + +@cocotb.test() +async def store32_index32(dut): + """Test vs*xei32_v_u32 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'out_size': 8000, + } + + await vector_store_indexed( + dut = dut, + elf_name = 'store32_index32.elf', + cases = [ + # Unordered + make_test_case('vsuxei32_v_u32m1', vl = 4), + make_test_case('vsuxei32_v_u32m1', vl = 3), + make_test_case('vsuxei32_v_u32m2', vl = 8), + make_test_case('vsuxei32_v_u32m2', vl = 7), + make_test_case('vsuxei32_v_u32m4', vl = 16), + make_test_case('vsuxei32_v_u32m4', vl = 15), + make_test_case('vsuxei32_v_u32m8', vl = 32), + make_test_case('vsuxei32_v_u32m8', vl = 31), + # Ordered + make_test_case('vsoxei32_v_u32m1', vl = 4), + make_test_case('vsoxei32_v_u32m1', vl = 3), + make_test_case('vsoxei32_v_u32m2', vl = 8), + make_test_case('vsoxei32_v_u32m2', vl = 7), + make_test_case('vsoxei32_v_u32m4', vl = 16), + make_test_case('vsoxei32_v_u32m4', vl = 15), + make_test_case('vsoxei32_v_u32m8', vl = 32), + make_test_case('vsoxei32_v_u32m8', vl = 31), + ], + data_dtype = np.uint32, + index_dtype = np.uint32, )