mirrors/llama.cpp
0
0 Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-10-28 09:45:55 +08:00
Code Activity

Use ggml_need_scale in code-base

Browse Source
This commit is contained in:
Oliver Simons committed 2026-10-27 18:47:47 +01:00
1 parent 2246653403
commit 63e34cf875
9 files changed
+39 -29
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-cpu/ggml-cpu.c
+2 -2
View File
@@ -2951,7 +2951,7 @@ struct ggml_cplan ggml_graph_plan(
case GGML_OP_ADD_ID:
case GGML_OP_ADD1:
{
if (ggml_is_quantized(node->src[0]->type) || node->src[0]->type == GGML_TYPE_F8_E4M3) {
if (ggml_is_quantized(node->src[0]->type) || ggml_needs_scale(node->src[0]->type)) {
cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks;
}
} break;
@@ -2999,7 +2999,7 @@ struct ggml_cplan ggml_graph_plan(
} break;
case GGML_OP_OUT_PROD:
{
if (ggml_is_quantized(node->src[0]->type) || node->src[0]->type == GGML_TYPE_F8_E4M3 ||
if (ggml_is_quantized(node->src[0]->type) || ggml_needs_scale(node->src[0]->type) ||
node->src[0]->type == GGML_TYPE_F16) {
cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks;
}
ggml/src/ggml-cpu/ops.cpp
+2 -2
View File
@@ -615,7 +615,7 @@ static void ggml_compute_forward_add_q_f32(
GGML_ASSERT(nb1 <= nb2);
GGML_ASSERT(nb2 <= nb3);
GGML_ASSERT(ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3);
GGML_ASSERT(ggml_is_quantized(src0->type) || ggml_needs_scale(src0->type));
GGML_ASSERT(src1->type == GGML_TYPE_F32);
GGML_ASSERT(ne00 % ggml_blck_size(type) == 0);
@@ -965,7 +965,7 @@ static void ggml_compute_forward_add1_q_f32(
GGML_ASSERT(nb1 <= nb2);
GGML_ASSERT(nb2 <= nb3);
GGML_ASSERT(ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3);
GGML_ASSERT(ggml_is_quantized(src0->type) || ggml_needs_scale(src0->type));
GGML_ASSERT(dst->type == src0->type);
GGML_ASSERT(src1->type == GGML_TYPE_F32);
ggml/src/ggml-cuda/ggml-cuda.cu
+4 -4
View File
@@ -1683,7 +1683,7 @@ static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml
const int cc = ggml_cuda_info().devices[ctx.device].cc;
const ggml_prec prec = (ggml_prec) ggml_get_op_params_i32(dst, 0);
ggml_type compute_type = src0->type;
if (ggml_is_quantized(compute_type) || compute_type == GGML_TYPE_F8_E4M3) {
if (ggml_is_quantized(compute_type) || ggml_needs_scale(compute_type)) {
compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
compute_type = GGML_TYPE_F32;
@@ -2214,7 +2214,7 @@ static bool ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
GGML_ASSERT(nb2 % nb1 == 0);
const ggml_type type_src1_sorted = (src0->type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc))
|| ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F8_E4M3 ? GGML_TYPE_F32 : src0->type;
|| ggml_is_quantized(src0->type) || ggml_needs_scale(src0->type) ? GGML_TYPE_F32 : src0->type;
const ggml_type type_dst_sorted = GGML_TYPE_F32;
const size_t ts_src1_sorted = ggml_type_size(type_src1_sorted);
const size_t ts_dst_sorted = ggml_type_size(type_dst_sorted);
@@ -4238,7 +4238,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
}
const ggml_tensor * scale = scale_lhs_mm ? scale_node->src[1] : scale_node->src[0];
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
if (!ggml_needs_scale(mm_node->src[0]->type) || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1 ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
@@ -4258,7 +4258,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
}
const ggml_tensor * scale = reshape->src[0];
if ((mm_node->src[0]->type != GGML_TYPE_NVFP4 && mm_node->src[0]->type != GGML_TYPE_F8_E4M3) || scale_node->type != GGML_TYPE_F32 ||
if (!ggml_needs_scale(mm_node->src[0]->type) || scale_node->type != GGML_TYPE_F32 ||
scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm_node->src[0]->ne[2] ||
!ggml_are_same_shape(scale_node, mm_node)) {
return nullptr;
src/llama-graph.cpp
+10 -14
View File
@@ -2991,13 +2991,13 @@ ggml_tensor * llm_graph_context::build_attn_mha(
cur = ggml_mul(ctx0, cur, v_scale);
}
} else {
// A cache without an exported scale stores raw FP8 values.
if ((k->type == GGML_TYPE_F8_E4M3 && !k_scale) || (v->type == GGML_TYPE_F8_E4M3 && !v_scale)) {
// A cache without an exported scale uses a unit scale.
if ((ggml_needs_scale(k->type) && !k_scale) || (ggml_needs_scale(v->type) && !v_scale)) {
ggml_tensor * unit_scale = ggml_fill(ctx0, ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, 1), 1.0f);
if (k->type == GGML_TYPE_F8_E4M3 && !k_scale) {
if (ggml_needs_scale(k->type) && !k_scale) {
k_scale = unit_scale;
}
if (v->type == GGML_TYPE_F8_E4M3 && !v_scale) {
if (ggml_needs_scale(v->type) && !v_scale) {
v_scale = unit_scale;
}
}
@@ -3193,24 +3193,20 @@ llm_graph_input_attn_kv * llm_graph_context::build_attn_inp_kv() const {
return (llm_graph_input_attn_kv *) res->add_input(std::move(inp));
}
static bool is_type_scaled(ggml_type type) {
return type == GGML_TYPE_F8_E4M3;
}
static void build_attn_scale_kv_inputs(
ggml_context * ctx,
const llama_kv_cache_context * mctx,
ggml_tensor * & k,
ggml_tensor * & v,
int32_t il) {
if (is_type_scaled(mctx->type_k()) && k) {
if (ggml_needs_scale(mctx->type_k()) && k) {
ggml_tensor * scale = mctx->get_k_scale(il);
if (scale) {
k = ggml_div(ctx, k, scale);
}
}
if (is_type_scaled(mctx->type_v()) && v) {
if (ggml_needs_scale(mctx->type_v()) && v) {
ggml_tensor * scale = mctx->get_v_scale(il);
if (scale) {
v = ggml_div(ctx, v, scale);
@@ -3278,8 +3274,8 @@ ggml_tensor * llm_graph_context::build_attn(
ggml_tensor * kq_mask = inp->get_kq_mask();
ggml_tensor * cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, 0, kq_scale, il,
is_type_scaled(mctx_cur->type_k()) ? mctx_cur->get_k_scale(il) : nullptr,
is_type_scaled(mctx_cur->type_v()) ? mctx_cur->get_v_scale(il) : nullptr);
ggml_needs_scale(mctx_cur->type_k()) ? mctx_cur->get_k_scale(il) : nullptr,
ggml_needs_scale(mctx_cur->type_v()) ? mctx_cur->get_v_scale(il) : nullptr);
cb(cur, "kqv_out", il);
if (inp->self_v_rot) {
@@ -3556,8 +3552,8 @@ ggml_tensor * llm_graph_context::build_attn(
ggml_tensor * v = use_kv_cur ? v_cur : mctx_cur->get_v(ctx0, il);
ggml_tensor * cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, 0, kq_scale, il,
is_type_scaled(mctx_cur->type_k()) ? mctx_cur->get_k_scale(il) : nullptr,
is_type_scaled(mctx_cur->type_v()) ? mctx_cur->get_v_scale(il) : nullptr);
ggml_needs_scale(mctx_cur->type_k()) ? mctx_cur->get_k_scale(il) : nullptr,
ggml_needs_scale(mctx_cur->type_v()) ? mctx_cur->get_v_scale(il) : nullptr);
cb(cur, "kqv_out", il);
if (v_rot) {
src/llama-kv-cache.cpp
+1 -1
View File
@@ -1952,7 +1952,7 @@ ggml_tensor * llama_kv_cache::build_rope_shift(
: hparams.rope_type;
ggml_tensor * tmp;
if (ggml_is_quantized(cur->type) || cur->type == GGML_TYPE_F8_E4M3) {
if (ggml_is_quantized(cur->type) || ggml_needs_scale(cur->type)) {
// dequantize to f32 -> RoPE -> quantize back
tmp = ggml_cast(ctx, cur, GGML_TYPE_F32);
src/llama-model-loader.cpp
+9
View File
@@ -1230,9 +1230,18 @@ struct ggml_tensor * llama_model_loader::create_tensor(
}
const ggml_tensor * scale_meta = nullptr;
ggml_tensor user_scale_meta = {};
if (op == GGML_OP_MUL_MAT || op == GGML_OP_MUL_MAT_ID) {
const std::string scale_name = LLM_TN_IMPL(tn.arch, tn.tensor, "scale", tn.bid, tn.xid).str();
scale_meta = get_tensor_meta(scale_name.c_str());
if (!scale_meta && files.empty()) {
const int64_t tid = gguf_find_tensor(metadata, scale_name.c_str());
if (tid != -1) {
user_scale_meta.type = gguf_get_tensor_type(metadata, tid);
memcpy(user_scale_meta.ne, gguf_get_tensor_ne(metadata, tid), sizeof(user_scale_meta.ne));
scale_meta = &user_scale_meta;
}
}
}
// select the buffer type for this tensor
src/llama-model.cpp
+6 -1
View File
@@ -1748,7 +1748,12 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
const std::string name = tn_scale.str();
const ggml_tensor * meta = ml.get_tensor_meta(name.c_str());
if (meta == nullptr) {
return nullptr;
const int64_t tid = gguf_find_tensor(ml.metadata, name.c_str());
if (tid == -1) {
return nullptr;
}
const int64_t * ne = gguf_get_tensor_ne(ml.metadata, tid);
return create_tensor(tn_scale, {ne[0], ne[1]}, TENSOR_NOT_REQUIRED);
}
if (meta->ne[1] > 1) {
return create_tensor(tn_scale, {meta->ne[0], meta->ne[1]}, TENSOR_NOT_REQUIRED);
src/llama-quant.cpp
+2 -2
View File
@@ -233,7 +233,7 @@ static void llama_tensor_dequantize_impl(
const size_t nelements, const int nthread
) {
const ggml_type_traits * qtype = ggml_get_type_traits(type);
if (ggml_is_quantized(type) || type == GGML_TYPE_F8_E4M3) {
if (ggml_is_quantized(type) || ggml_needs_scale(type)) {
if (qtype->to_float == NULL) {
throw std::runtime_error(format("type %s unsupported for integer quantization: no dequantization available", ggml_type_name(type)));
}
@@ -247,7 +247,7 @@ static void llama_tensor_dequantize_impl(
ggml_fp16_to_fp32_row((const ggml_fp16_t *)data, f32_output, nelements);
} else if (type == GGML_TYPE_BF16) {
ggml_bf16_to_fp32_row((const ggml_bf16_t *)data, f32_output, nelements);
} else if (ggml_is_quantized(type) || type == GGML_TYPE_F8_E4M3) {
} else if (ggml_is_quantized(type) || ggml_needs_scale(type)) {
qtype->to_float(data, f32_output, nelements);
} else {
GGML_ABORT("fatal error"); // unreachable
tests/test-backend-ops.cpp
+3 -3
View File
@@ -88,7 +88,7 @@ static void init_tensor_uniform(ggml_tensor * tensor, float min = -1.0f, float m
if (tensor->type == GGML_TYPE_F32 || tensor->type == GGML_TYPE_I32) {
ggml_backend_tensor_set(tensor, data.data(), 0, nels * sizeof(float));
} else if (ggml_is_quantized(tensor->type) || tensor->type == GGML_TYPE_F8_E4M3 || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) {
} else if (ggml_is_quantized(tensor->type) || ggml_needs_scale(tensor->type) || tensor->type == GGML_TYPE_F16 || tensor->type == GGML_TYPE_BF16) {
GGML_ASSERT(nels % ggml_blck_size(tensor->type) == 0);
// dummy importance matrix
@@ -262,7 +262,7 @@ static std::vector<float> tensor_to_float(const ggml_tensor * t) {
const auto * tt = ggml_get_type_traits(t->type);
size_t bs = ggml_blck_size(t->type);
std::vector<float> vq(ggml_blck_size(t->type));
bool uses_type_conversion = ggml_is_quantized(t->type) || t->type == GGML_TYPE_F8_E4M3;
bool uses_type_conversion = ggml_is_quantized(t->type) || ggml_needs_scale(t->type);
// access elements by index to avoid gaps in views
for (int64_t i3 = 0; i3 < t->ne[3]; i3++) {
@@ -12464,7 +12464,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
continue;
}
for (bool with_lane_scale : {false, true}) {
if (with_lane_scale && type != GGML_TYPE_NVFP4 && type != GGML_TYPE_F8_E4M3) {
if (with_lane_scale && !ggml_needs_scale(type)) {
continue;
}
test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,