diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 1a66f112b..f3f5d4d77 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1225,40 +1225,6 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta return res; } -ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax_part(ggml_metal_library_t lib) { - char base[256]; - char name[256]; - - snprintf(base, 256, "kernel_mul_mm_id_amax_part_f32"); - snprintf(name, 256, "%s", base); - - ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); - if (!res.pipeline) { - res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); - } - - res.smem = 32*sizeof(float); - - return res; -} - -ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax(ggml_metal_library_t lib) { - char base[256]; - char name[256]; - - snprintf(base, 256, "kernel_mul_mm_id_amax_f32"); - snprintf(name, 256, "%s", base); - - ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); - if (!res.pipeline) { - res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); - } - - res.smem = 32*sizeof(float); - - return res; -} - ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_map0(ggml_metal_library_t lib, int ne02, int ne20) { char base[256]; char name[256]; @@ -1301,18 +1267,14 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id(ggml_m const bool bc_inp = op->src[0]->ne[0] % 32 != 0; - // src1 prec [TAG_GGML_PREC] - const bool amax = ggml_get_op_params_i32(op, 3) == GGML_PREC_F32; - snprintf(base, 256, "kernel_mul_mm_id_%s_%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1)); - snprintf(name, 256, "%s_bci=%d_amax=%d", base, bc_inp, amax); + snprintf(name, 256, "%s_bci=%d", base, bc_inp); ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); if (!res.pipeline) { ggml_metal_cv_t cv = ggml_metal_cv_init(); ggml_metal_cv_set_bool(cv, bc_inp, FC_MUL_MM + 0); - ggml_metal_cv_set_bool(cv, amax, FC_MUL_MM + 6); res = ggml_metal_library_compile_pipeline(lib, base, name, cv); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 72a8f92f2..6d77c57ff 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -141,8 +141,6 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_map0 (ggml_metal_library_t lib, int ne02, int ne20); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_zero (ggml_metal_library_t lib); -struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax(ggml_metal_library_t lib); -struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id_amax_part(ggml_metal_library_t lib); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mm_id (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_argmax (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index ea6a677f9..9db266f44 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -14,8 +14,6 @@ #define N_MM_SIMD_GROUP_X 2 #define N_MM_SIMD_GROUP_Y 2 -#define N_MM_NPART_AMAX 256 - #define N_ARGMAX_NPART_MAX 256 #define N_ARGMAX_CHUNK 4096 @@ -579,14 +577,6 @@ typedef struct { uint64_t nb21; } ggml_metal_kargs_mul_mm_id_zero; -typedef struct { - int32_t ne00; - int32_t ne01; - int32_t ne02; - uint64_t nb01; - uint64_t nb02; -} ggml_metal_kargs_mul_mm_id_amax; - typedef struct { int32_t ne00; int32_t ne02; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 589f7fafa..0cebbe5d6 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2812,15 +2812,6 @@ size_t ggml_metal_op_mul_mat_id_extra_ids(const ggml_tensor * op) { return ggml_type_size(GGML_TYPE_I32)*(ne20*ne21 + ne02); } -size_t ggml_metal_op_mul_mat_id_extra_amax(const ggml_tensor * op) { - assert(op->op == GGML_OP_MUL_MAT_ID); - - GGML_UNUSED(op); - - // 2 scaling factors (8 bytes) + N_MM_NPART_AMAX per-threadgroup scales for stage-1 - return 8 + N_MM_NPART_AMAX*sizeof(float); -} - int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); @@ -2872,39 +2863,6 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) { ggml_metal_buffer_id bid_ids = bid_tpe; bid_ids.offs += ggml_metal_op_mul_mat_id_extra_tpe(op); - ggml_metal_buffer_id bid_amax = bid_ids; - bid_amax.offs += ggml_metal_op_mul_mat_id_extra_ids(op); - - // src1 prec [TAG_GGML_PREC] - const bool use_amax = ggml_get_op_params_i32(op, 3) == GGML_PREC_F32; - - // src1 rescale factors, computed before the matmul - // ref: https://github.com/ggml-org/llama.cpp/pull/26223 - if (use_amax) { - ggml_metal_kargs_mul_mm_id_amax args = { - /*.ne00 =*/ ne10, - /*.ne01 =*/ ne11, - /*.ne02 =*/ ne12, - /*.nb01 =*/ nb11, - /*.nb02 =*/ nb12, - }; - - auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_amax_part(lib); - - const size_t smem = pipeline.smem; - - GGML_ASSERT(smem <= props_dev->max_theadgroup_memory_size); - - ggml_metal_encoder_set_pipeline(enc, pipeline); - ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); - ggml_metal_encoder_set_buffer (enc, bid_src1, 1); - ggml_metal_encoder_set_buffer (enc, bid_amax, 2); - - ggml_metal_encoder_set_threadgroup_memory_size(enc, smem, 0); - - ggml_metal_encoder_dispatch_threadgroups(enc, N_MM_NPART_AMAX, 1, 1, 256, 1, 1); - } - { ggml_metal_kargs_mul_mm_id_map0 args = { ne10, @@ -2954,20 +2912,9 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_dispatch_threadgroups(enc, ne20, ne21, 1, 32, 1, 1); } + // this barrier is always needed because the next kernel has to wait for the id maps to be computed ggml_metal_op_concurrency_reset(ctx); - if (use_amax) { - auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id_amax(lib); - - ggml_metal_encoder_set_pipeline(enc, pipeline); - ggml_metal_encoder_set_buffer (enc, bid_amax, 0); - - ggml_metal_encoder_dispatch_threadgroups(enc, 1, 1, 1, 32, 1, 1); - - // the next kernel has to wait for the amax data - ggml_metal_op_concurrency_reset(ctx); - } - { auto pipeline = ggml_metal_library_get_pipeline_mul_mm_id(lib, op); @@ -2996,7 +2943,6 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_set_buffer (enc, bid_tpe, 3); ggml_metal_encoder_set_buffer (enc, bid_ids, 4); ggml_metal_encoder_set_buffer (enc, bid_dst, 5); - ggml_metal_encoder_set_buffer (enc, bid_amax, 6); const size_t smem = pipeline.smem; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index b8296bf2b..e14a2dd5f 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -36,7 +36,6 @@ size_t ggml_metal_op_mul_mat_id_extra_tpe(const struct ggml_tensor * op); // id map [n_tokens, n_expert] size_t ggml_metal_op_mul_mat_id_extra_ids(const struct ggml_tensor * op); -size_t ggml_metal_op_mul_mat_id_extra_amax(const struct ggml_tensor * op); // return true if we should use the FA vector kernel for this op bool ggml_metal_op_flash_attn_ext_use_vec(const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp index 074a9d8e6..d0771eb49 100644 --- a/ggml/src/ggml-metal/ggml-metal.cpp +++ b/ggml/src/ggml-metal/ggml-metal.cpp @@ -226,7 +226,6 @@ static size_t ggml_backend_metal_buffer_type_get_alloc_size(ggml_backend_buffer_ { res += ggml_metal_op_mul_mat_id_extra_tpe(tensor); res += ggml_metal_op_mul_mat_id_extra_ids(tensor); - res += ggml_metal_op_mul_mat_id_extra_amax(tensor); } break; case GGML_OP_FLASH_ATTN_EXT: { diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal index 4acc24e5a..2c8492037 100644 --- a/ggml/src/ggml-metal/kernels/mul_mm.metal +++ b/ggml/src/ggml-metal/kernels/mul_mm.metal @@ -7,7 +7,6 @@ constant short FC_mul_mm_ne12 [[function_constant(FC_MUL_MM + 2)]]; constant short FC_mul_mm_ne13 [[function_constant(FC_MUL_MM + 3)]]; constant short FC_mul_mm_r2 [[function_constant(FC_MUL_MM + 4)]]; constant short FC_mul_mm_r3 [[function_constant(FC_MUL_MM + 5)]]; -constant bool FC_mul_mm_id_amax [[function_constant(FC_MUL_MM + 6)]]; constant int32_t FC_mul_mm_id_map0_ne02 [[function_constant(FC_MUL_MM + 7)]]; // each block_q contains 16*nl weights @@ -532,85 +531,6 @@ kernel void kernel_mul_mm_id_zero( } } -kernel void kernel_mul_mm_id_amax_part_f32( - constant ggml_metal_kargs_mul_mm_id_amax & args, - device const char * src1, - device char * dst, - threadgroup char * shmem [[threadgroup(0)]], - uint tgpig[[threadgroup_position_in_grid]], - ushort tiitg[[thread_index_in_threadgroup]], - ushort tiisg[[thread_index_in_simdgroup]], - ushort sgitg[[simdgroup_index_in_threadgroup]], - ushort ntg[[threads_per_threadgroup]]) { - const int nrow = args.ne01*args.ne02; - - float lmax = 0.0f; - - for (int ir = tgpig; ir < nrow; ir += N_MM_NPART_AMAX) { - const int i01 = ir % args.ne01; - const int i02 = ir / args.ne01; - - device const float * row = (device const float *) (src1 + i02*args.nb02 + i01*args.nb01); - - for (int i00 = tiitg; i00 < args.ne00; i00 += ntg) { - lmax = max(lmax, fabs(row[i00])); - } - } - - float amax = simd_max(lmax); - - threadgroup float * shared_amax = (threadgroup float *) shmem; - - if (ntg > N_SIMDWIDTH) { - if (sgitg == 0) { - shared_amax[tiisg] = 0.0f; - } - threadgroup_barrier(mem_flags::mem_threadgroup); - - if (tiisg == 0) { - shared_amax[sgitg] = amax; - } - threadgroup_barrier(mem_flags::mem_threadgroup); - - amax = shared_amax[tiisg]; - amax = simd_max(amax); - } - - if (tiitg == 0) { - ((device float *) (dst + 8))[tgpig] = amax; - } -} - -kernel void kernel_mul_mm_id_amax_f32( - device char * dst, - ushort tiitg[[thread_index_in_threadgroup]]) { - device const float * part = (device const float *) (dst + 8); - - float amax = 0.0f; - - for (int i = tiitg; i < N_MM_NPART_AMAX; i += N_SIMDWIDTH) { - amax = max(amax, part[i]); - } - - amax = simd_max(amax); - - if (tiitg == 0) { - // leave a comfortable margin below the f16 max of 65504 - float scale = 1.0f; - - // isfinite: src1 already inf/nan is not ours to fix - keep the - // scale at 1.0 instead of turning it into a different failure - if (isfinite(amax) && amax > 32768.0f) { - scale = exp2(ceil(log2(amax)) - 15.0f); - } - - device float * d = (device float *) dst; - - d[0] = 1.0f/scale; // exact: scale is a power of two - d[1] = scale; - } -} - typedef decltype(kernel_mul_mm_id_map0<1>) kernel_mul_mm_id_map0_t; template [[host_name("kernel_mul_mm_id_map0_ne20_1" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<1>; @@ -633,7 +553,6 @@ kernel void kernel_mul_mm_id( device const char * htpe, device const char * hids, device char * dst, - device const char * amax, threadgroup char * shmem [[threadgroup(0)]], uint3 tgpig[[threadgroup_position_in_grid]], ushort tiitg[[thread_index_in_threadgroup]], @@ -689,10 +608,6 @@ kernel void kernel_mul_mm_id( // skip the upper half of the token tile when the expert did not fill it constexpr short NR1H = NR1/2; - // power-of-two rescaling - const float s1_inv = FC_mul_mm_id_amax ? ((device const float *) amax)[0] : 1.0f; - const float s1_scale = FC_mul_mm_id_amax ? ((device const float *) amax)[1] : 1.0f; - #ifndef GGML_METAL_HAS_TENSOR S0_8x8 ma[4]; S1_8x8 mb[2]; @@ -798,7 +713,7 @@ kernel void kernel_mul_mm_id( const short ib = 4*sx + sy; - *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) (*((device T1 *) y + i) * (T1) s1_inv) : 0; + *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0; } } else { const short sx = (tiitg%NL1); @@ -811,7 +726,7 @@ kernel void kernel_mul_mm_id( const short ib = 4*sx + sy; - *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)((*((device T1_2x4 *) y)) * (T1) s1_inv); + *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)(*((device T1_2x4 *) y)); } #else // load data and store to threadgroup memory @@ -859,7 +774,7 @@ kernel void kernel_mul_mm_id( //const short lx = (tiitg/NL1)%8; //const short ly = i; - *(sb + NK*(8*sy + ly) + 8*sx + lx) = loop_k + iy + i < args.ne00 ? (S1) (*((device T1 *) y + i) * (T1) s1_inv) : 0; + *(sb + NK*(8*sy + ly) + 8*sx + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0; } } else { const short sx = (tiitg%NL1); @@ -870,7 +785,7 @@ kernel void kernel_mul_mm_id( //const short lx = (tiitg/NL1)%8; //const short ly = i; - *(threadgroup S1_2x4 *)(sb + NK*(8*sy + ly) + 8*sx) = (S1_2x4)((*((device T1_2x4 *) y)) * (T1) s1_inv); + *(threadgroup S1_2x4 *)(sb + NK*(8*sy + ly) + 8*sx) = (S1_2x4)(*((device T1_2x4 *) y)); } #endif @@ -961,12 +876,12 @@ kernel void kernel_mul_mm_id( int i = tiisg; for (; i < nr0/4; i += 32) { - *(D4 + i) = *(C4 + i) * s1_scale; + *(D4 + i) = *(C4 + i); } i = (4*(nr0/4)) + tiisg; for (; i < nr0; i += 32) { - *(D + i) = *(C + i) * s1_scale; + *(D + i) = *(C + i); } } } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 636c61be2..5b583aa4c 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5617,14 +5617,9 @@ static void init_mul_mat_id_ids(ggml_context * ctx, int n_mats, bool skip = fals } } -static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats, bool skip = false, float amax = 1.0f) { +static void init_mul_mat_id_tensors(ggml_context * ctx, int n_mats, bool skip = false) { for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) { - if (t->type == GGML_TYPE_I32) { - continue; - } - if (amax != 1.0f && t->type == GGML_TYPE_F32) { - init_tensor_uniform(t, -amax, amax); - } else { + if (t->type != GGML_TYPE_I32) { init_tensor_uniform(t); } } @@ -5641,13 +5636,12 @@ struct test_mul_mat_id : public test_case { const int64_t m; const int64_t n; const int64_t k; - const float amax; // magnitude of src1 const int64_t m_v; // rows of as in memory, the experts of as are strided for m_v > m, no view for m_v == 0 const bool ids_offset; const bool skip; // set some ids to a negative value std::string vars() override { - return VARS_TO_STR12(type_a, type_b, n_mats, n_used, b, m, n, k, amax, m_v, ids_offset, skip); + return VARS_TO_STR11(type_a, type_b, n_mats, n_used, b, m, n, k, m_v, ids_offset, skip); } double max_nmse_err() override { @@ -5672,9 +5666,9 @@ struct test_mul_mat_id : public test_case { test_mul_mat_id(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32, int n_mats = 8, int n_used = 2, bool b = false, int64_t m = 32, int64_t n = 32, int64_t k = 32, - float amax = 1.0f, int64_t m_v = 0, bool ids_offset = false, bool skip = false) + int64_t m_v = 0, bool ids_offset = false, bool skip = false) : type_a(type_a), type_b(type_b), n_mats(n_mats), n_used(n_used), b(b), - m(m), n(n), k(k), amax(amax), m_v(m_v), ids_offset(ids_offset), skip(skip) { + m(m), n(n), k(k), m_v(m_v), ids_offset(ids_offset), skip(skip) { GGML_ASSERT(n_used <= n_mats); GGML_ASSERT(m_v == 0 || m_v > m); } @@ -5700,16 +5694,11 @@ struct test_mul_mat_id : public test_case { ggml_tensor * out = ggml_mul_mat_id(ctx, as, b, ids); ggml_set_name(out, "out"); - if (amax > 65504.0f) { - // src1 exceeds F16 range - ggml_prec_set_src(out, GGML_PREC_F32, 1); - } - return out; } void initialize_tensors(ggml_context * ctx) override { - init_mul_mat_id_tensors(ctx, n_mats, skip, amax); + init_mul_mat_id_tensors(ctx, n_mats, skip); } void reinit_perf_iter(ggml_context * ctx) override { @@ -11337,8 +11326,8 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 8, 1, 64, {8, 1}, {1, 1}, {0, 1, 2, 3}, 0, 1, false, 16, 16)); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 8, 16, 64, {8, 1}, {1, 1}, {0, 1, 2, 3}, 0, 1, false, 16, 16)); // as is a view whose experts are strided by more rows than it uses - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 4, 2, false, 8, 1, 64, 1.0f, 64)); - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 2, false, 8, 16, 64, 1.0f, 64)); + test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 4, 2, false, 8, 1, 64, 64)); + test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 2, false, 8, 16, 64, 64)); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 576, 512, 576, {1,1}, {1,1})); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 1, 2048, 8192, {1, 1}, {1, 1})); @@ -11452,7 +11441,7 @@ static std::vector> make_test_cases_eval() { } for (ggml_type ta : {GGML_TYPE_F16, GGML_TYPE_Q4_0}) { - test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 1025, 10, false, 64, 33, 256, 1.0f, 0, true)); + test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 1025, 10, false, 64, 33, 256, 0, true)); test_cases.emplace_back(new test_mul_mat_id_dup(ta, GGML_TYPE_F32, 2050, 10, true, 64, 33, 256)); } #endif @@ -11467,12 +11456,6 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat_id(type_a, GGML_TYPE_F32, 4, 4, false, 16, 10, 256)); } - // test src1 f16 overflow - for (int n : {16, 32, 64}) { - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q4_K, GGML_TYPE_F32, 128, 4, false, 4096, n, 2048, 1e5f)); - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8, 2, false, 512, n, 256, 1e5f)); - } - for (ggml_type type_a : base_types) { for (ggml_type type_b : {GGML_TYPE_F32 /*, GGML_TYPE_F16 */}) { for (int n_mats : {4, 8}) { @@ -11510,14 +11493,14 @@ static std::vector> make_test_cases_eval() { for (int n_used : {1, 2, 4}) { for (bool b : {false, true}) { for (int n : {1, 4, 33, 129}) { - test_cases.emplace_back(new test_mul_mat_id(type_a, GGML_TYPE_F32, 8, n_used, b, 512, n, 256, 1.0f, 0, false, true)); + test_cases.emplace_back(new test_mul_mat_id(type_a, GGML_TYPE_F32, 8, n_used, b, 512, n, 256, 0, false, true)); } } } } // n_used == n_mats, so every expert can be skipped - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 4, false, 512, 64, 256, 1.0f, 0, false, true)); - test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, 1, 1.0f, 0, false, true)); + test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 4, 4, false, 512, 64, 256, 0, false, true)); + test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, 1, 0, false, true)); test_cases.emplace_back(new test_mul_mat_id_fusion(GGML_TYPE_F16, GGML_TYPE_F32, 16, 16, false, 32, 32, 32, 3, false, true)); for (int bs : {1, 4, 512}) {