Improve support for mixed dtype indexed loads/stores. Better support data dtype < index dtype cases. Change-Id: I607af046bfed91d321d68bf50813a165e96676c9
diff --git a/hdl/chisel/src/kelvin/scalar/Lsu.scala b/hdl/chisel/src/kelvin/scalar/Lsu.scala index e3ef588..5ef23ef 100644 --- a/hdl/chisel/src/kelvin/scalar/Lsu.scala +++ b/hdl/chisel/src/kelvin/scalar/Lsu.scala
@@ -104,6 +104,16 @@ LsuOp.VSTORE_OINDEXED, LsuOp.VSTORE_UINDEXED) } + def isIndexedVector(op: LsuOp.Type): Bool = { + op.isOneOf(LsuOp.VLOAD_OINDEXED, LsuOp.VLOAD_UINDEXED, + LsuOp.VSTORE_OINDEXED, LsuOp.VSTORE_UINDEXED) + } + + def isNonindexedVector(op: LsuOp.Type): Bool = { + op.isOneOf(LsuOp.VLOAD_UNIT, LsuOp.VLOAD_STRIDED, + LsuOp.VSTORE_UNIT, LsuOp.VSTORE_STRIDED) + } + def isFlush(op: LsuOp.Type): Bool = { op.isOneOf(LsuOp.FENCEI, LsuOp.FLUSHAT, LsuOp.FLUSHALL) } @@ -172,7 +182,7 @@ // 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 lmul = Option.when(p.enableRvv) { UInt(3.W) } val nfields = Option.when(p.enableRvv) { UInt(3.W) } override def toPrintable: Printable = { @@ -203,13 +213,9 @@ } if (p.enableRvv) { result.elemWidth.get := cmd.elemWidth.get - val effectiveLmul = MuxCase(rvvState.get.bits.lmul(1, 0), Seq( - // Treat fractional LMULs as LMUL=1 - (rvvState.get.bits.lmul(2)) -> 0.U(2.W), - // If mask operation, always force LMUL=1. - (cmd.isMaskOperation()) -> 0.U(2.W), - )) - result.lmul.get := 1.U(1.W) << effectiveLmul + // If mask operation, always make LMUL=1. + result.lmul.get := Mux(cmd.isMaskOperation(), 0.U, rvvState.get.bits.lmul) + // 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 @@ -243,20 +249,18 @@ indices: UInt, indexWidth: UInt, sew: UInt): Vec[UInt] = { - val indices8 = UIntToVec(indices, 8) - val indices16 = UIntToVec(indices, 16).padTo(bytesPerSlot, 0.U) - val indices32 = UIntToVec(indices, 32).padTo(bytesPerSlot, 0.U) + val indices8 = UIntToVec(indices, 8).map(x => Cat(0.U(24.W), x)) + val indices16 = UIntToVec(indices, 16).map(x => Cat(0.U(16.W), x)) + val indices32 = UIntToVec(indices, 32) 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)))), + (indexWidth === "b000".U) -> VecInit(indices8), // 16-bit indices. - (indexWidth === "b101".U) -> VecInit((0 until bytesPerSlot).map( - i => Cat(0.U(16.W), indices16(i)))), + (indexWidth === "b101".U) -> VecInit(indices16 ++ indices16), // 32-bit indices. - (indexWidth === "b110".U) -> VecInit((0 until bytesPerSlot).map( - i => indices32(i))), + (indexWidth === "b110".U) -> VecInit( + indices32 ++ indices32 ++ indices32 ++ indices32), )) MuxCase(VecInit.fill(bytesPerSlot)(0.U(32.W)), Seq( @@ -289,6 +293,8 @@ val data = Vec(bytesPerSlot, UInt(8.W)) val pendingVector = Bool() val pendingWriteback = Bool() + val subvector = UInt(3.W) + val subvectors = UInt(3.W) 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 @@ -311,6 +317,7 @@ active.reduce(_||_) || // Active transaction pendingWriteback || // Send result back to regfile (nfields =/= segment) || // No pending segments + (subvector =/= subvectors) || // No pending subvectors (LsuOp.isVector(op) && (lmul =/= 0.U)) // More operations in progress ) } @@ -346,6 +353,8 @@ result.store := store result.pc := pc result.pendingWriteback := pendingWriteback + result.subvector := subvector + result.subvectors := subvectors result.lmul := lmul result.baseAddr := baseAddr result.elemStride := elemStride @@ -368,11 +377,16 @@ result.elemWidth := elemWidth result.sew := sew - result.data := Mux(updated && LsuOp.isVector(op) && rvv2lsu.vregfile.valid, + val shouldUpdate = updated && ( + LsuOp.isNonindexedVector(op) || + (subvectors === 1.U) || + rvv2lsu.idx.valid) + + result.data := Mux(shouldUpdate && LsuOp.isVector(op) && rvv2lsu.vregfile.valid, UIntToVec(rvv2lsu.vregfile.bits.data, 8), data) - result.active := Mux(updated && LsuOp.isVector(op) && rvv2lsu.mask.valid, + result.active := Mux(shouldUpdate && LsuOp.isVector(op) && rvv2lsu.mask.valid, VecInit(rvv2lsu.mask.bits.asBools), active) - result.pendingVector := pendingVector && !updated + result.pendingVector := pendingVector && !shouldUpdate result.nfields := nfields result.segment := segment @@ -402,6 +416,8 @@ result.baseAddr := baseAddr result.addrs := addrs result.pendingWriteback := pendingWriteback + result.subvector := subvector + result.subvectors := subvectors result.pendingVector := pendingVector result.lmul := lmul result.active := (0 until bytesPerSlot).map( @@ -445,17 +461,26 @@ val vectorWriteback = writeback && LsuOp.isVector(op) + val subvectorDone = ((subvector + 1.U) === subvectors) + val subvectorNext = Mux(subvectorDone, 0.U, subvector + 1.U) + result.subvectors := subvectors + result.subvector := MuxCase(subvectorNext, Seq( + (!vectorWriteback) -> subvector, + // Final subvector, final segment, final lmul, don't reset subvector. + ((segment === nfields) && (lmul === 1.U) && ((subvector + 1.U) === subvectors)) -> (subvector + 1.U), + )) + val segmentNext = MuxCase(segment, Seq( // Final segment, No next LMUL: don't reset segment - (vectorWriteback && (segment === nfields) && (lmul === 1.U)) -> segment, + (subvectorDone && (segment === nfields) && (lmul === 1.U)) -> segment, // Final segment, next LMUL: reset segment - (vectorWriteback && (segment === nfields)) -> 0.U, + (subvectorDone && (segment === nfields)) -> 0.U, // Next segment - vectorWriteback -> (segment + 1.U), + subvectorDone -> (segment + 1.U), )) result.segment := segmentNext - val lmulUpdate = vectorWriteback && (segment === nfields) + val lmulUpdate = subvectorDone && (segment === nfields) val lmulNext = MuxCase(lmul, Seq( // Don't decrease below 0! (lmul === 0.U) -> 0.U, @@ -488,7 +513,7 @@ // Finished one lmul, start from segment 0 again lmulUpdate -> (rd - nextLmulVectorRewind + 1.U), // Finished one segment, proceed to next - vectorWriteback -> (rd + nextSegmentVectorOffset), + subvectorDone -> (rd + nextSegmentVectorOffset), )) result @@ -511,6 +536,8 @@ result.pc := pc result.pendingWriteback := pendingWriteback result.pendingVector := pendingVector + result.subvector := subvector + result.subvectors := subvectors result.active := (0 until bytesPerSlot).map(i => active(i) & ~selected(i)) result.baseAddr := baseAddr result.addrs := addrs @@ -551,6 +578,7 @@ cf" $i: ${active(i)}, 0x${addrs(i)}%x, 0x${data(i)}%x\n") cf"store: $store\n op: ${op}\n pendingVector: ${pendingVector}\n" + cf" pendingWriteback: ${pendingWriteback}\n lmul: ${lmul}\n" + + cf" subvector: ${subvector}\n subvectors: ${subvectors}\n" + cf" nfields: ${nfields}\n segment: ${segment}\n" + cf" elemWidth: 0b${elemWidth}%b elemStride: ${elemStride}\n" + lines.reduce(_+_) @@ -569,23 +597,50 @@ result.store := uop.store result.pc := uop.pc if (p.enableRvv) { - val lmul = uop.lmul.getOrElse(0.U) - result.lmul := lmul + val effectiveLmul = MuxCase(uop.lmul.getOrElse(0.U)(1, 0), Seq( + // Treat fractional LMULs as LMUL=1 + (uop.lmul.getOrElse(0.U)(2)) -> 0.U(2.W), + )) + result.lmul := 1.U(1.W) << effectiveLmul + val nfields = Mux(LsuOp.isVector(uop.op), uop.nfields.get, 0.U) result.nfields := nfields result.segment := 0.U - result.nextSegmentVectorOffset := lmul - result.nextLmulVectorRewind := nfields * lmul + result.nextSegmentVectorOffset := result.lmul + result.nextLmulVectorRewind := nfields * result.lmul } else { result.lmul := 0.U result.nextSegmentVectorOffset := 0.U result.nextLmulVectorRewind := 0.U } + // Determine number of rvv2lsu interactions required for one vector for + // indexed loads. This occurs when the index dtype is greater than data + // dtype. + val elemWidth = uop.elemWidth.get + val elemMultiplier = MuxCase(1.U, Seq( + // 8-bit data, 16-bit indices + ((elemWidth === "b101".U) && (uop.sew.get === 0.U)) -> 2.U, + // 8-bit data, 32-bit indices + ((elemWidth === "b110".U) && (uop.sew.get === 0.U)) -> 4.U, + // 16-bit data, 32-bit indices + ((elemWidth === "b110".U) && (uop.sew.get === 1.U)) -> 2.U, + )) + val subvectors = MuxCase(1.U, Seq( + ((elemMultiplier === 2.U) && (uop.lmul.get.asSInt >= 0.S)) -> 2.U, + ((elemMultiplier === 4.U) && (uop.lmul.get.asSInt >= 0.S)) -> 4.U, + ((elemMultiplier === 4.U) && (uop.lmul.get.asSInt === -1.S)) -> 2.U, + )) + // All vector ops require writeback. Lsu needs to inform RVV core store uop // has completed. result.pendingWriteback := !uop.store || LsuOp.isVector(uop.op) result.pendingVector := LsuOp.isVector(uop.op) + result.subvector := 0.U + result.subvectors := MuxCase(0.U, Seq( + LsuOp.isIndexedVector(uop.op) -> subvectors, + LsuOp.isNonindexedVector(uop.op) -> 1.U, + )) val active = MuxCase(0.U(bytesPerSlot.W), Seq( uop.op.isOneOf(LsuOp.LB, LsuOp.LBU, LsuOp.SB) -> "b1".U(bytesPerSlot.W),
diff --git a/tests/cocotb/rvv/load_store/load8_index16.cc b/tests/cocotb/rvv/load_store/load8_index16.cc index f6decff..94b0bbe 100644 --- a/tests/cocotb/rvv/load_store/load8_index16.cc +++ b/tests/cocotb/rvv/load_store/load8_index16.cc
@@ -44,21 +44,51 @@ } __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); + asm volatile("vsetvli zero, %0, e16, m2, ta, ma;" + "vle16.v v2, (%1);" + "vsetvli zero, %0, e8, m1, ta, ma;" + "vluxei16.v v2, (%2), v2;" + "vse8.v v2, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v2"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e16, m4, ta, ma;" + "vle16.v v4, (%1);" + "vsetvli zero, %0, e8, m2, ta, ma;" + "vluxei16.v v4, (%2), v4;" + "vse8.v v4, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v4"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e16, m8, ta, ma;" + "vle16.v v8, (%1);" + "vsetvli zero, %0, e8, m4, ta, ma;" + "vluxei16.v v8, (%2), v8;" + "vse8.v v8, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v8"); + + // TODO: Revert once compiler bug is eliminated + // 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 @@ -75,21 +105,51 @@ } __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); + asm volatile("vsetvli zero, %0, e16, m2, ta, ma;" + "vle16.v v2, (%1);" + "vsetvli zero, %0, e8, m1, ta, ma;" + "vloxei16.v v2, (%2), v2;" + "vse8.v v2, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v2"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e16, m4, ta, ma;" + "vle16.v v4, (%1);" + "vsetvli zero, %0, e8, m2, ta, ma;" + "vloxei16.v v4, (%2), v4;" + "vse8.v v4, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v4"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e16, m8, ta, ma;" + "vle16.v v8, (%1);" + "vsetvli zero, %0, e8, m4, ta, ma;" + "vloxei16.v v8, (%2), v8;" + "vse8.v v8, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v8"); + + // TODO: Revert once compiler bug is eliminated + // 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); } }
diff --git a/tests/cocotb/rvv/load_store/load8_index32.cc b/tests/cocotb/rvv/load_store/load8_index32.cc index 524e86d..071cca0 100644 --- a/tests/cocotb/rvv/load_store/load8_index32.cc +++ b/tests/cocotb/rvv/load_store/load8_index32.cc
@@ -38,21 +38,51 @@ } __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); + asm volatile("vsetvli zero, %0, e32, m2, ta, ma;" + "vle32.v v2, (%1);" + "vsetvli zero, %0, e8, mf2, ta, ma;" + "vluxei32.v v2, (%2), v2;" + "vse8.v v2, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v2"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e32, m4, ta, ma;" + "vle32.v v4, (%1);" + "vsetvli zero, %0, e8, m1, ta, ma;" + "vluxei32.v v4, (%2), v4;" + "vse8.v v4, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v4"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e32, m8, ta, ma;" + "vle32.v v8, (%1);" + "vsetvli zero, %0, e8, m2, ta, ma;" + "vluxei32.v v8, (%2), v8;" + "vse8.v v8, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v8"); + + // TODO: Revert once compiler bug is eliminated + // 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 @@ -63,21 +93,51 @@ } __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); + asm volatile("vsetvli zero, %0, e32, m2, ta, ma;" + "vle32.v v2, (%1);" + "vsetvli zero, %0, e8, mf2, ta, ma;" + "vloxei32.v v2, (%2), v2;" + "vse8.v v2, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v2"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e32, m4, ta, ma;" + "vle32.v v4, (%1);" + "vsetvli zero, %0, e8, m1, ta, ma;" + "vloxei32.v v4, (%2), v4;" + "vse8.v v4, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v4"); + + // TODO: Revert once compiler bug is eliminated + // 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); + asm volatile("vsetvli zero, %0, e32, m8, ta, ma;" + "vle32.v v8, (%1);" + "vsetvli zero, %0, e8, m2, ta, ma;" + "vluxei32.v v8, (%2), v8;" + "vse8.v v8, (%3);" + : + : "r" (vl), "r" (index_buf), "r" (in_buf), "r" (out_buf) + : "v8"); + + // TODO: Revert once compiler bug is eliminated + // 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); } }
diff --git a/tests/cocotb/rvv_load_store_test.py b/tests/cocotb/rvv_load_store_test.py index ea99784..bc7b0ab 100644 --- a/tests/cocotb/rvv_load_store_test.py +++ b/tests/cocotb/rvv_load_store_test.py
@@ -552,23 +552,23 @@ 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), + 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), + 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, @@ -593,25 +593,21 @@ # 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), + 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), # 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), + 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), ], dtype = np.uint8, index_dtype = np.uint32,