metal : remove mul_mm_id amax rescaling (#30924)
The src1 power-of-two rescale in kernel_mul_mm_id is superseded by the per expert-token amax scaling done in the compute graph (ref: https://github.com/ggml-org/llama.cpp/pull/30861), which keeps the MoE down-projection activations within f16 range before the op runs. Drop the two-stage reduction kernels, their scratch buffer and the GGML_PREC_F32 dispatch added in https://github.com/ggml-org/llama.cpp/pull/26223, as well as the test cases that drive MUL_MAT_ID src1 out of f16 range. Assisted-by: pi:llama.cpp/ggml-org/Qwen3.8-Flash-Next-GGUF (MacBook, M5 Max, 128GB)
This commit is contained in:
Georgi Gerganov authored and 
GitHub committed
1 parent
9a0fde9967
commit
3f03b17724
8 files changed
+20
-228
No files matched your search
@@ -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);
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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:
|
||||
{
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
+12
-29
@@ -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<std::unique_ptr<test_case>> 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<std::unique_ptr<test_case>> 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<std::unique_ptr<test_case>> 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<std::unique_ptr<test_case>> 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}) {
|
||||
|
||||
Reference in new issue
Block a user