Finish intrinsic coverage for indexed stores Change-Id: I2b655b4f7a0588eaf6a6b553b9a43818a16382f1
diff --git a/tests/cocotb/BUILD b/tests/cocotb/BUILD index e33dc35..6c0da8b 100644 --- a/tests/cocotb/BUILD +++ b/tests/cocotb/BUILD
@@ -222,8 +222,11 @@ "store_strided_all_vtypes_test", "store8_index8", "store8_index8_seg", + "store8_index16", + "store8_index32", "store16_index8", "store16_index16", + "store16_index32", "store32_index8", "store32_index16", "store32_index32",
diff --git a/tests/cocotb/rvv/load_store/BUILD b/tests/cocotb/rvv/load_store/BUILD index f129395..a2aa111 100644 --- a/tests/cocotb/rvv/load_store/BUILD +++ b/tests/cocotb/rvv/load_store/BUILD
@@ -134,12 +134,21 @@ "store8_index8_seg": { "srcs": ["store8_index8_seg.cc"], }, + "store8_index16": { + "srcs": ["store8_index16.cc"], + }, + "store8_index32": { + "srcs": ["store8_index32.cc"], + }, "store16_index8": { "srcs": ["store16_index8.cc"], }, "store16_index16": { "srcs": ["store16_index16.cc"], }, + "store16_index32": { + "srcs": ["store16_index32.cc"], + }, "store32_index8": { "srcs": ["store32_index8.cc"], }, @@ -211,8 +220,11 @@ ":load32_stride8_m1_partial.elf", ":store8_index8.elf", ":store8_index8_seg.elf", + ":store8_index16.elf", + ":store8_index32.elf", ":store16_index8.elf", ":store16_index16.elf", + ":store16_index32.elf", ":store32_index8.elf", ":store32_index16.elf", ":store32_index32.elf",
diff --git a/tests/cocotb/rvv/load_store/store16_index16.cc b/tests/cocotb/rvv/load_store/store16_index16.cc index 6487165..dcffd60 100644 --- a/tests/cocotb/rvv/load_store/store16_index16.cc +++ b/tests/cocotb/rvv/load_store/store16_index16.cc
@@ -17,7 +17,7 @@ namespace { constexpr size_t scatter_count = 64; -constexpr size_t buf_size = 16000; +constexpr size_t buf_size = 15000; } // namespace size_t vl __attribute__((section(".data"))) = 8;
diff --git a/tests/cocotb/rvv/load_store/store16_index32.cc b/tests/cocotb/rvv/load_store/store16_index32.cc new file mode 100644 index 0000000..fb199b4 --- /dev/null +++ b/tests/cocotb/rvv/load_store/store16_index32.cc
@@ -0,0 +1,104 @@ +// 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 = 15000; +} // 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. +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 vsuxei32_v_u16mf2() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16mf2(in_buf, vl); + __riscv_vsuxei32_v_u16mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u16m1() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m1(in_buf, vl); + __riscv_vsuxei32_v_u16m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u16m2() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m2(in_buf, vl); + __riscv_vsuxei32_v_u16m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u16m4() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m4(in_buf, vl); + __riscv_vsuxei32_v_u16m4(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei32_v_u16mf2() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16mf2(in_buf, vl); + __riscv_vsoxei32_v_u16mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u16m1() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m1(in_buf, vl); + __riscv_vsoxei32_v_u16m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u16m2() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m2(in_buf, vl); + __riscv_vsoxei32_v_u16m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u16m4() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle16_v_u16m4(in_buf, vl); + __riscv_vsoxei32_v_u16m4(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei32_v_u16m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/store32_index16.cc b/tests/cocotb/rvv/load_store/store32_index16.cc index 4c081f3..5b3f1e2 100644 --- a/tests/cocotb/rvv/load_store/store32_index16.cc +++ b/tests/cocotb/rvv/load_store/store32_index16.cc
@@ -17,7 +17,7 @@ namespace { constexpr size_t scatter_count = 32; -constexpr size_t buf_size = 4000; +constexpr size_t buf_size = 7000; } // namespace size_t vl __attribute__((section(".data"))) = 4;
diff --git a/tests/cocotb/rvv/load_store/store32_index32.cc b/tests/cocotb/rvv/load_store/store32_index32.cc index 07001f5..01ce3b8 100644 --- a/tests/cocotb/rvv/load_store/store32_index32.cc +++ b/tests/cocotb/rvv/load_store/store32_index32.cc
@@ -17,7 +17,7 @@ namespace { constexpr size_t scatter_count = 32; -constexpr size_t buf_size = 8000; +constexpr size_t buf_size = 7000; } // namespace size_t vl __attribute__((section(".data"))) = 8;
diff --git a/tests/cocotb/rvv/load_store/store32_index8.cc b/tests/cocotb/rvv/load_store/store32_index8.cc index 465994f..19d5546 100644 --- a/tests/cocotb/rvv/load_store/store32_index8.cc +++ b/tests/cocotb/rvv/load_store/store32_index8.cc
@@ -17,7 +17,7 @@ namespace { constexpr size_t scatter_count = 32; -constexpr size_t buf_size = 257; +constexpr size_t buf_size = 512; } // namespace size_t vl __attribute__((section(".data"))) = 4;
diff --git a/tests/cocotb/rvv/load_store/store8_index16.cc b/tests/cocotb/rvv/load_store/store8_index16.cc new file mode 100644 index 0000000..412e3e1 --- /dev/null +++ b/tests/cocotb/rvv/load_store/store8_index16.cc
@@ -0,0 +1,120 @@ +// 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 = 30000; +} // namespace + +size_t vl __attribute__((section(".data"))) = 16; +// 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. +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 vsuxei16_v_u8mf4() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsuxei16_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u8mf2() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsuxei16_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u8m1() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsuxei16_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u8m2() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsuxei16_v_u8m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei16_v_u8m4() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m4(in_buf, vl); + __riscv_vsuxei16_v_u8m4(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei16_v_u8mf4() { + auto indices = __riscv_vle16_v_u16mf2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsoxei16_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u8mf2() { + auto indices = __riscv_vle16_v_u16m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsoxei16_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u8m1() { + auto indices = __riscv_vle16_v_u16m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsoxei16_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u8m2() { + auto indices = __riscv_vle16_v_u16m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsoxei16_v_u8m2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei16_v_u8m4() { + auto indices = __riscv_vle16_v_u16m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m4(in_buf, vl); + __riscv_vsoxei16_v_u8m4(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei16_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv/load_store/store8_index32.cc b/tests/cocotb/rvv/load_store/store8_index32.cc new file mode 100644 index 0000000..2cd6d35 --- /dev/null +++ b/tests/cocotb/rvv/load_store/store8_index32.cc
@@ -0,0 +1,104 @@ +// 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 = 30000; +} // namespace + +size_t vl __attribute__((section(".data"))) = 16; +// 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. +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 vsuxei32_v_u8mf4() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsuxei32_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u8mf2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsuxei32_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u8m1() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsuxei32_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsuxei32_v_u8m2() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsuxei32_v_u8m2(out_buf, indices, data, vl); +} + +// Ordered +__attribute__((used, retain)) void vsoxei32_v_u8mf4() { + auto indices = __riscv_vle32_v_u32m1(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf4(in_buf, vl); + __riscv_vsoxei32_v_u8mf4(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u8mf2() { + auto indices = __riscv_vle32_v_u32m2(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8mf2(in_buf, vl); + __riscv_vsoxei32_v_u8mf2(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u8m1() { + auto indices = __riscv_vle32_v_u32m4(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m1(in_buf, vl); + __riscv_vsoxei32_v_u8m1(out_buf, indices, data, vl); +} + +__attribute__((used, retain)) void vsoxei32_v_u8m2() { + auto indices = __riscv_vle32_v_u32m8(index_buf, vl); + // TODO(davidgao): Remove once compiler bug is eliminated + asm volatile("" ::: "vtype"); + auto data = __riscv_vle8_v_u8m2(in_buf, vl); + __riscv_vsoxei32_v_u8m2(out_buf, indices, data, vl); +} +} + +void (*impl)() __attribute__((section(".data"))) = &vsuxei32_v_u8m1; + +int main(int argc, char** argv) { + impl(); + return 0; +}
diff --git a/tests/cocotb/rvv_load_store_test.py b/tests/cocotb/rvv_load_store_test.py index 424a247..c3c2f79 100644 --- a/tests/cocotb/rvv_load_store_test.py +++ b/tests/cocotb/rvv_load_store_test.py
@@ -214,7 +214,9 @@ struct_bytes = np.dtype(data_dtype).itemsize * segments # Don't go beyond the buffer. - index_max = min(np.iinfo(index_dtype).max, out_size) + index_max = min( + np.iinfo(index_dtype).max, + out_size * np.dtype(data_dtype).itemsize) assert vl * struct_bytes <= index_max index_max = index_max - struct_bytes # TODO(davidgao): currently assuming the vl is supported. @@ -2229,6 +2231,88 @@ @cocotb.test() +async def store8_index16(dut): + """Test vs*xei16_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'segments': 1, + 'out_size': 30000, + } + + await vector_store_segmented_indexed( + dut = dut, + elf_name = 'store8_index16.elf', + cases = [ + # Unordered + make_test_case('vsuxei16_v_u8mf4', vl = 4), + make_test_case('vsuxei16_v_u8mf4', vl = 3), + make_test_case('vsuxei16_v_u8mf2', vl = 8), + make_test_case('vsuxei16_v_u8mf2', vl = 7), + make_test_case('vsuxei16_v_u8m1', vl = 16), + make_test_case('vsuxei16_v_u8m1', vl = 15), + make_test_case('vsuxei16_v_u8m2', vl = 32), + make_test_case('vsuxei16_v_u8m2', vl = 31), + make_test_case('vsuxei16_v_u8m4', vl = 64), + make_test_case('vsuxei16_v_u8m4', vl = 63), + # Ordered + make_test_case('vsoxei16_v_u8mf2', vl = 4), + make_test_case('vsoxei16_v_u8mf2', vl = 3), + make_test_case('vsoxei16_v_u8mf2', vl = 8), + make_test_case('vsoxei16_v_u8mf2', vl = 7), + make_test_case('vsoxei16_v_u8m1', vl = 16), + make_test_case('vsoxei16_v_u8m1', vl = 15), + make_test_case('vsoxei16_v_u8m2', vl = 32), + make_test_case('vsoxei16_v_u8m2', vl = 31), + make_test_case('vsoxei16_v_u8m4', vl = 64), + make_test_case('vsoxei16_v_u8m4', vl = 63), + ], + data_dtype = np.uint8, + index_dtype = np.uint16, + ) + + +@cocotb.test() +async def store8_index32(dut): + """Test vs*xei32_v_u8 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'segments': 1, + 'out_size': 30000, + } + + await vector_store_segmented_indexed( + dut = dut, + elf_name = 'store8_index32.elf', + cases = [ + # Unordered + make_test_case('vsuxei32_v_u8mf4', vl = 4), + make_test_case('vsuxei32_v_u8mf4', vl = 3), + make_test_case('vsuxei32_v_u8mf2', vl = 8), + make_test_case('vsuxei32_v_u8mf2', vl = 7), + make_test_case('vsuxei32_v_u8m1', vl = 16), + make_test_case('vsuxei32_v_u8m1', vl = 15), + make_test_case('vsuxei32_v_u8m2', vl = 32), + make_test_case('vsuxei32_v_u8m2', vl = 31), + # Ordered + make_test_case('vsoxei32_v_u8mf2', vl = 4), + make_test_case('vsoxei32_v_u8mf2', vl = 3), + make_test_case('vsoxei32_v_u8mf2', vl = 8), + make_test_case('vsoxei32_v_u8mf2', vl = 7), + make_test_case('vsoxei32_v_u8m1', vl = 16), + make_test_case('vsoxei32_v_u8m1', vl = 15), + make_test_case('vsoxei32_v_u8m2', vl = 32), + make_test_case('vsoxei32_v_u8m2', vl = 31), + ], + data_dtype = np.uint8, + index_dtype = np.uint32, + ) + + +@cocotb.test() async def store16_index8(dut): """Test vs*xei8_v_u16 usage accessible from intrinsics.""" def make_test_case(impl: str, vl: int): @@ -2279,7 +2363,7 @@ 'impl': impl, 'vl': vl, 'segments': 1, - 'out_size': 16000, + 'out_size': 15000, } await vector_store_segmented_indexed( @@ -2315,6 +2399,45 @@ @cocotb.test() +async def store16_index32(dut): + """Test vs*xei32_v_u16 usage accessible from intrinsics.""" + def make_test_case(impl: str, vl: int): + return { + 'impl': impl, + 'vl': vl, + 'segments': 1, + 'out_size': 15000, + } + + await vector_store_segmented_indexed( + dut = dut, + elf_name = 'store16_index32.elf', + cases = [ + # Unordered + make_test_case('vsuxei32_v_u16mf2', vl = 4), + make_test_case('vsuxei32_v_u16mf2', vl = 3), + make_test_case('vsuxei32_v_u16m1', vl = 8), + make_test_case('vsuxei32_v_u16m1', vl = 7), + make_test_case('vsuxei32_v_u16m2', vl = 16), + make_test_case('vsuxei32_v_u16m2', vl = 15), + make_test_case('vsuxei32_v_u16m4', vl = 32), + make_test_case('vsuxei32_v_u16m4', vl = 31), + # Ordered + make_test_case('vsoxei32_v_u16mf2', vl = 4), + make_test_case('vsoxei32_v_u16mf2', vl = 3), + make_test_case('vsoxei32_v_u16m1', vl = 8), + make_test_case('vsoxei32_v_u16m1', vl = 7), + make_test_case('vsoxei32_v_u16m2', vl = 16), + make_test_case('vsoxei32_v_u16m2', vl = 15), + make_test_case('vsoxei32_v_u16m4', vl = 32), + make_test_case('vsoxei32_v_u16m4', vl = 31), + ], + data_dtype = np.uint16, + index_dtype = np.uint32, + ) + + +@cocotb.test() async def store32_index8(dut): """Test vs*xei8_v_u32 usage accessible from intrinsics.""" def make_test_case(impl: str, vl: int): @@ -2322,7 +2445,7 @@ 'impl': impl, 'vl': vl, 'segments': 1, - 'out_size': 257, + 'out_size': 512, } await vector_store_segmented_indexed( @@ -2361,7 +2484,7 @@ 'impl': impl, 'vl': vl, 'segments': 1, - 'out_size': 4000, + 'out_size': 7000, } await vector_store_segmented_indexed( @@ -2400,7 +2523,7 @@ 'impl': impl, 'vl': vl, 'segments': 1, - 'out_size': 8000, + 'out_size': 7000, } await vector_store_segmented_indexed(