From 126fa16e7a9096aeb32ad8027849d201a6b5f878 Mon Sep 17 00:00:00 2001 From: Andgihat Date: Fri, 18 Sep 2026 01:07:12 +0300 Subject: [PATCH 1/3] llama: support prism.hadamard folded weights (PrismML Bonsai 2) Bonsai 2 ships weights that were folded with a normalized Sylvester Walsh-Hadamard rotation, so the activation must be rotated by the same transform before every folded matmul. Without it the model loads and silently produces garbage. - llama-model: parse the prism.hadamard.* metadata block, validate it, and build one F32 rotation per (block size, buffer type) plus the per-width sign vectors in the same buffer as the folded weights - llama-graph: apply sign flip and rotation to the activation in build_lora_mm / build_lora_mm_id, memoized per (activation, rotation), and apply the inverse after the token-embedding lookup - llama-context: one-time graph check that every folded weight is consumed through its transform, turning a silent wrong-math path into a load error The generic pieces (GGML_HINT_SRC0_IS_HADAMARD, llama_mul_mat_hadamard, the CUDA FWHT kernel) were already present upstream; only the model-side glue was missing. Verified on Ternary-Bonsai-2-27B Q2_0 (group 64, 402 folded weights) on an RTX 5060 Ti: coherent output and 46 t/s decode, where the same file on the previous build returns gibberish. --- src/llama-context.cpp | 71 ++++++++++ src/llama-context.h | 3 + src/llama-graph.cpp | 65 ++++++++- src/llama-graph.h | 30 +++++ src/llama-model.cpp | 297 ++++++++++++++++++++++++++++++++++++++++++ src/llama-model.h | 12 ++ 6 files changed, 476 insertions(+), 2 deletions(-) diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 96ff3f55c18e..92a10874f685 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -205,6 +205,68 @@ uint64_t llama_kv_tail_planner_timing_ns(const llama_context * ctx) { // llama_context // +// Verify that every Hadamard-folded weight consumed by the graph receives its +// activation-side transform, and every latent lookup table gets the inverse. +// An architecture whose matmul path bypasses the transform helpers would +// otherwise load cleanly and silently compute wrong results. +static void llama_verify_hadamard_graph( + ggml_cgraph * gf, + const llama_hadamard_rotations & rotations, + const llama_hadamard_rotations & inverses) { + auto unwrap = [](const ggml_tensor * t) { + while (t && (t->op == GGML_OP_RESHAPE || t->op == GGML_OP_VIEW)) { + t = t->src[0]; + } + return t; + }; + + std::map lookups; // get_rows results of latent tables + + for (int i = 0; i < ggml_graph_n_nodes(gf); ++i) { + const ggml_tensor * node = ggml_graph_node(gf, i); + + if (node->op == GGML_OP_GET_ROWS && inverses.count(node->src[0])) { + lookups.emplace(node, false); + continue; + } + + if (node->op != GGML_OP_MUL_MAT && node->op != GGML_OP_MUL_MAT_ID) { + continue; + } + + if (node->op == GGML_OP_MUL_MAT && ((const int32_t *) node->op_params)[1] == GGML_HINT_SRC0_IS_HADAMARD) { + const auto lk = lookups.find(unwrap(node->src[1])); + if (lk != lookups.end()) { + lk->second = true; + } + continue; + } + + const auto it = rotations.find(node->src[0]); + if (it == rotations.end()) { + continue; + } + const ggml_tensor * src = unwrap(node->src[1]); + const bool transformed = src && src->op == GGML_OP_MUL_MAT && + ((const int32_t *) src->op_params)[1] == GGML_HINT_SRC0_IS_HADAMARD && + src->src[0] == it->second.rot; + if (!transformed) { + throw std::runtime_error(format( + "Hadamard-folded weight '%s' is consumed without its activation transform; " + "this graph's matmul path does not support prism.hadamard folding", + node->src[0]->name)); + } + } + + for (const auto & [node, ok] : lookups) { + if (!ok) { + throw std::runtime_error(format( + "Hadamard-latent table '%s' is read without the inverse transform", + node->src[0]->name)); + } + } +} + static llm_graph_type ctx_type_to_graph_type(llama_context_type ctx_type) { switch (ctx_type) { case LLAMA_CONTEXT_TYPE_DEFAULT: return LLM_GRAPH_TYPE_DEFAULT; @@ -3018,6 +3080,13 @@ ggml_cgraph * llama_context::graph_reserve( auto * gf = model.build_graph(gparams); + // verify transform coverage on the pristine graph: after scheduling, + // cross-backend copies break the producer chain the check follows + if (!hadamard_verified && gf && (!model.hadamard_rotations.empty() || !model.hadamard_inverses.empty())) { + llama_verify_hadamard_graph(gf, model.hadamard_rotations, model.hadamard_inverses); + hadamard_verified = true; + } + this->n_outputs = save_n_outputs; // initialize scheduler with the specified graph @@ -3053,6 +3122,8 @@ llm_graph_params llama_context::graph_params( /*.loras =*/ loras.get(), /*.mctx =*/ mctx, /*.cross =*/ &cross, + /*.hadamard_rotations =*/ &model.hadamard_rotations, + /*.hadamard_inverses =*/ &model.hadamard_inverses, /*.samplers =*/ sampling.samplers, /*.n_outputs =*/ n_outputs, /*.cb =*/ graph_get_cb(), diff --git a/src/llama-context.h b/src/llama-context.h index c5c121a525d1..b0f960a7eab8 100644 --- a/src/llama-context.h +++ b/src/llama-context.h @@ -373,6 +373,9 @@ struct llama_context { llm_graph_result_ptr gf_res_prev; llm_graph_result_ptr gf_res_reserve; + // one-time Hadamard transform-coverage check on the first built graph + bool hadamard_verified = false; + // host buffer for the model output (logits and embeddings) ggml_backend_buffer_ptr buf_output; diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 5c523d09931b..2680dcf24478 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1880,6 +1880,8 @@ llm_graph_context::llm_graph_context(const llm_graph_params & params) : loras (params.loras), mctx (params.mctx), cross (params.cross), + hadamard_rotations(params.hadamard_rotations), + hadamard_inverses (params.hadamard_inverses), samplers (params.samplers), cb_func (params.cb), res (params.res), @@ -1902,11 +1904,56 @@ ggml_tensor * llm_graph_context::build_cvec( return cvec->apply_to(ctx0, cur, il); } +ggml_tensor * llm_graph_context::build_hadamard_activation( + ggml_tensor * w, + ggml_tensor * cur) const { + if (!hadamard_rotations) { + return cur; + } + + const auto it = hadamard_rotations->find(w); + if (it == hadamard_rotations->end()) { + return cur; + } + + const auto & t = it->second; + + // another folded weight on this same activation already built the transform + const auto memo_key = std::make_pair((const ggml_tensor *) cur, (const ggml_tensor *) t.rot); + const auto memo_it = hadamard_memo.find(memo_key); + if (memo_it != hadamard_memo.end()) { + return memo_it->second; + } + + ggml_tensor * cur_mm = cur; + + if (t.perm_rep > 1) { + // tiled [hd, nk, rep] -> grouped [hd, rep, nk] feature order + ggml_tensor * x = ggml_is_contiguous(cur_mm) ? cur_mm : ggml_cont(ctx0, cur_mm); + const int64_t ne1 = x->ne[1], ne2 = x->ne[2], ne3 = x->ne[3]; + x = ggml_reshape_4d(ctx0, x, t.perm_hd, t.perm_nk, t.perm_rep, ne1*ne2*ne3); + x = ggml_cont(ctx0, ggml_permute(ctx0, x, 0, 2, 1, 3)); + cur_mm = ggml_reshape_4d(ctx0, x, t.perm_hd*t.perm_nk*t.perm_rep, ne1, ne2, ne3); + } + + if (t.signs) { + cur_mm = ggml_mul(ctx0, cur_mm, t.signs); + } + + cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, t.rot); + + hadamard_memo[memo_key] = cur_mm; + + return cur_mm; +} + ggml_tensor * llm_graph_context::build_lora_mm( ggml_tensor * w, ggml_tensor * cur, ggml_tensor * w_s) const { - ggml_tensor * res = ggml_mul_mat(ctx0, w, cur); + ggml_tensor * cur_mm = build_hadamard_activation(w, cur); + + ggml_tensor * res = ggml_mul_mat(ctx0, w, cur_mm); if (w_s) { res = ggml_mul(ctx0, res, w_s); @@ -1938,7 +1985,9 @@ ggml_tensor * llm_graph_context::build_lora_mm_id( ggml_tensor * cur, // ggml_tensor * b ggml_tensor * ids, ggml_tensor * w_s) const { - ggml_tensor * res = ggml_mul_mat_id(ctx0, w, cur, ids); + ggml_tensor * cur_mm = build_hadamard_activation(w, cur); + + ggml_tensor * res = ggml_mul_mat_id(ctx0, w, cur_mm, ids); if (w_s) { const int64_t n_expert = w_s->ne[0]; @@ -2773,6 +2822,18 @@ ggml_tensor * llm_graph_context::build_inp_embd(ggml_tensor * tok_embd) const { cur = ggml_get_rows(ctx0, tok_embd, inp->tokens); + // a Hadamard-latent embedding table stores rotated rows; restore the + // primal basis right after the lookup: h = s * (H z) + if (hadamard_inverses) { + const auto it = hadamard_inverses->find(tok_embd); + if (it != hadamard_inverses->end()) { + cur = llama_mul_mat_hadamard(ctx0, cur, it->second.rot); + if (it->second.signs) { + cur = ggml_mul(ctx0, cur, it->second.signs); + } + } + } + // apply lora for embedding tokens if needed for (const auto & lora : *loras) { llama_adapter_lora_weight * lw = lora.first->get_weight(tok_embd); diff --git a/src/llama-graph.h b/src/llama-graph.h index 32cafe4eeeaa..152785f8ca50 100644 --- a/src/llama-graph.h +++ b/src/llama-graph.h @@ -11,11 +11,27 @@ #include #include #include +#include struct ggml_cgraph; struct ggml_context; struct ggml_tensor; +// Maps a folded model weight to the activation-side transform applied +// immediately before the matmul: optional sign flip, then the normalized +// blockwise Hadamard rotation. +struct llama_hadamard_transform { + ggml_tensor * rot; + ggml_tensor * signs; // nullptr for identity sign mode + // when perm_rep > 1 the activation arrives with its feature axis in tiled + // head order [hd, nk, rep] and must be permuted to the grouped order + // [hd, rep, nk] the fold was computed in, before signs and rotation + int64_t perm_hd = 0; + int64_t perm_nk = 0; + int64_t perm_rep = 0; +}; +using llama_hadamard_rotations = std::unordered_map; + struct llama_cparams; struct llama_layer; @@ -866,6 +882,8 @@ struct llm_graph_params { const llama_adapter_loras * loras; const llama_memory_context_i * mctx; const llama_cross * cross; + const llama_hadamard_rotations * hadamard_rotations; + const llama_hadamard_rotations * hadamard_inverses; std::map samplers; @@ -1106,6 +1124,12 @@ struct llm_graph_context { const llama_adapter_loras * loras; const llama_memory_context_i * mctx; const llama_cross * cross; + const llama_hadamard_rotations * hadamard_rotations; + const llama_hadamard_rotations * hadamard_inverses; + + // Transforms shared by folded weights on the same activation. Key is (input, rotation); + // both must match. Valid for one graph build only. + mutable std::map, ggml_tensor *> hadamard_memo; std::map samplers; @@ -1130,6 +1154,12 @@ struct llm_graph_context { int il) const; // do mat_mul, while optionally apply lora and per-tensor scale + // if w is a Hadamard-folded weight, return the activation with its + // transform applied (sign flip, then rotation); otherwise return it as is + ggml_tensor * build_hadamard_activation( + ggml_tensor * w, + ggml_tensor * cur) const; + ggml_tensor * build_lora_mm( ggml_tensor * w, ggml_tensor * cur, diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 083f2db0dba4..8de9beefb9c7 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1237,6 +1237,146 @@ void llama_model_base::load_hparams(llama_model_loader & ml) { gguf_kv.emplace(name, value); } + uint32_t hadamard_version = 0; + if (ml.get_key("prism.hadamard.version", hadamard_version, false)) { + if (hadamard_version != 1) { + throw std::runtime_error(format("unsupported prism.hadamard.version: %u", hadamard_version)); + } + + uint32_t block_size = 0; + std::string transform; + std::string axis; + std::string sign_mode; + std::vector weight_names; + + ml.get_key("prism.hadamard.block_size", block_size); + ml.get_key("prism.hadamard.transform", transform); + ml.get_key("prism.hadamard.axis", axis); + ml.get_key("prism.hadamard.sign_mode", sign_mode); + ml.get_arr("prism.hadamard.weight_names", weight_names); + + if (block_size == 0 || (block_size & (block_size - 1)) != 0) { + throw std::runtime_error(format("invalid prism.hadamard.block_size: %u", block_size)); + } + if (transform != "normalized-sylvester-walsh-hadamard") { + throw std::runtime_error(format("unsupported prism.hadamard.transform: %s", transform.c_str())); + } + if (axis != "input-last-dimension") { + throw std::runtime_error(format("unsupported prism.hadamard.axis: %s", axis.c_str())); + } + if (sign_mode != "identity" && sign_mode != "explicit") { + throw std::runtime_error(format("unsupported prism.hadamard.sign_mode: %s", sign_mode.c_str())); + } + if (weight_names.empty()) { + throw std::runtime_error("prism.hadamard.weight_names is empty"); + } + + if (sign_mode == "explicit") { + std::vector sign_widths; + std::vector sign_values; + ml.get_arr("prism.hadamard.sign_widths", sign_widths); + ml.get_arr("prism.hadamard.sign_values", sign_values); + // explicit mode with no widths would leave the sign table empty, which reads + // as identity later and silently changes the model function + if (sign_widths.empty()) { + throw std::runtime_error("prism.hadamard.sign_mode is explicit but sign_widths is empty"); + } + size_t off = 0; + for (const int32_t width : sign_widths) { + if (width <= 0 || (uint32_t) width % block_size != 0 || off + width > sign_values.size()) { + throw std::runtime_error(format("invalid prism.hadamard sign width: %d", width)); + } + auto & vec = hadamard_sign_data[width]; + vec.assign(sign_values.begin() + off, sign_values.begin() + off + width); + for (const int32_t v : vec) { + if (v != 1 && v != -1) { + throw std::runtime_error("prism.hadamard sign values must be +/-1"); + } + } + off += width; + } + if (off != sign_values.size()) { + throw std::runtime_error("prism.hadamard.sign_values length mismatch"); + } + } + + ml.get_key("prism.hadamard.gdn_v_grouped", hadamard_gdn_v_grouped, false); + + // the activation-side transform is applied only by build_lora_mm/build_lora_mm_id; + // refuse to load folded weights for architectures or tensor kinds that are not + // verified to route every matmul through those helpers, rather than run wrong math + switch (arch) { + case LLM_ARCH_LLAMA: + case LLM_ARCH_QWEN3: + case LLM_ARCH_QWEN3MOE: + case LLM_ARCH_QWEN35: + case LLM_ARCH_QWEN35MOE: + case LLM_ARCH_QWEN3NEXT: + break; + default: + throw std::runtime_error(format( + "prism.hadamard: arch '%s' is not verified to apply the activation transform to all folded weights", + llm_arch_name(arch))); + } + + const auto is_foldable_weight = [](const std::string & name) { + static const char * kinds[] = { + "attn_q", "attn_k", "attn_v", "attn_qkv", "attn_gate", "attn_output", + "ffn_gate", "ffn_up", "ffn_down", + "ffn_gate_exps", "ffn_up_exps", "ffn_down_exps", "ffn_gate_up_exps", + "ffn_gate_shexp", "ffn_up_shexp", "ffn_down_shexp", + "ssm_out", + }; + if (name == "output.weight") { + return true; // the output head is built through build_lora_mm in every arch + } + if (name.compare(0, 4, "blk.") != 0) { + return false; + } + size_t pos = 4; + while (pos < name.size() && isdigit((unsigned char) name[pos])) { + pos++; + } + if (pos == 4 || pos >= name.size() || name[pos] != '.') { + return false; + } + pos++; + for (const char * kind : kinds) { + const std::string suffix = std::string(kind) + ".weight"; + if (name.compare(pos, std::string::npos, suffix) == 0) { + return true; + } + } + return false; + }; + + for (const auto & weight_name : weight_names) { + if (!is_foldable_weight(weight_name)) { + throw std::runtime_error(format( + "prism.hadamard: weight '%s' is not on a verified Hadamard-aware matmul path", weight_name.c_str())); + } + if (!hadamard_weight_blocks.emplace(weight_name, block_size).second) { + throw std::runtime_error(format("duplicate prism.hadamard weight: %s", weight_name.c_str())); + } + } + + // tensors consumed by row lookup store latent rows and need the + // inverse transform applied to the lookup result instead + std::vector inverse_names; + ml.get_arr("prism.hadamard.inverse_weight_names", inverse_names, false); + for (const auto & name : inverse_names) { + // the graph applies the inverse only to the token-embedding lookup; any + // other latent table would load and silently stay rotated + if (name != "token_embd.weight") { + throw std::runtime_error(format( + "prism.hadamard: weight '%s' is not a verified inverse-after-lookup table", name.c_str())); + } + if (hadamard_weight_blocks.count(name) || !hadamard_inverse_blocks.emplace(name, block_size).second) { + throw std::runtime_error(format("duplicate prism.hadamard inverse weight: %s", name.c_str())); + } + } + } + // get general kv ml.get_key(LLM_KV_GENERAL_NAME, name, false); @@ -1862,6 +2002,163 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { } } + if (!hadamard_weight_blocks.empty() || !hadamard_inverse_blocks.empty()) { + struct hadamard_rotation { + uint32_t block_size; + ggml_backend_buffer_type_t buft; + ggml_tensor * tensor; + }; + + std::vector rotations; + std::map, ggml_tensor *> sign_tensors; + + const std::pair *, llama_hadamard_rotations *> groups[] = { + { &hadamard_weight_blocks, &hadamard_rotations }, + { &hadamard_inverse_blocks, &hadamard_inverses }, + }; + // inverse (lookup-side) transforms must not inherit a host buffer + // type from a CPU-mapped table: the per-token transform would then + // ping-pong across the PCIe boundary. Prefer the buffer type the + // forward rotations live on (the GPU when layers are offloaded). + ggml_backend_buffer_type_t preferred_buft = nullptr; + + for (const auto & [blocks, target] : groups) + for (const auto & entry : *blocks) { + const std::string & weight_name = entry.first; + const uint32_t block_size = entry.second; + const ggml_tensor * weight = get_tensor(weight_name.c_str()); + if (weight == nullptr) { + throw std::runtime_error(format("prism.hadamard weight not found: %s", weight_name.c_str())); + } + if (weight->ne[0] % block_size != 0) { + throw std::runtime_error(format( + "prism.hadamard block size %u does not divide input dimension %lld for %s", + block_size, (long long) weight->ne[0], weight_name.c_str())); + } + if (weight->buffer == nullptr) { + throw std::runtime_error(format("prism.hadamard weight has no buffer: %s", weight_name.c_str())); + } + + ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(weight->buffer); + if (target == &hadamard_rotations) { + preferred_buft = buft; + } else if (preferred_buft) { + buft = preferred_buft; + } + auto it = std::find_if(rotations.begin(), rotations.end(), + [block_size, buft](const hadamard_rotation & rotation) { + return rotation.block_size == block_size && rotation.buft == buft; + }); + + if (it == rotations.end()) { + ggml_init_params params = { + /*.mem_size =*/ ggml_tensor_overhead(), + /*.mem_buffer =*/ NULL, + /*.no_alloc =*/ true, + }; + ggml_context_ptr ctx { ggml_init(params) }; + if (!ctx) { + throw std::runtime_error("failed to create Hadamard rotation context"); + } + + ggml_tensor * rotation = ggml_new_tensor_2d(ctx.get(), GGML_TYPE_F32, block_size, block_size); + char rotation_name[GGML_MAX_NAME]; + snprintf(rotation_name, sizeof(rotation_name), "prism.hadamard.%u", block_size); + ggml_set_name(rotation, rotation_name); + + ggml_backend_buffer_ptr buffer { ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft) }; + if (!buffer) { + throw std::runtime_error(format("unable to allocate %s Hadamard rotation buffer", ggml_backend_buft_name(buft))); + } + ggml_backend_buffer_set_usage(buffer.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + + std::vector data((size_t) block_size * block_size); + const float scale = 1.0f / sqrtf((float) block_size); + for (uint32_t row = 0; row < block_size; ++row) { + for (uint32_t col = 0; col < block_size; ++col) { + uint32_t parity = row & col; + parity ^= parity >> 16; + parity ^= parity >> 8; + parity ^= parity >> 4; + parity ^= parity >> 2; + parity ^= parity >> 1; + data[(size_t) row * block_size + col] = (parity & 1) ? -scale : scale; + } + } + ggml_backend_tensor_set(rotation, data.data(), 0, data.size() * sizeof(float)); + + std::vector buffers; + buffers.emplace_back(std::move(buffer)); + pimpl->ctxs_bufs.emplace_back(std::move(ctx), std::move(buffers)); + rotations.push_back({ block_size, buft, rotation }); + it = std::prev(rotations.end()); + } + + ggml_tensor * sign_tensor = nullptr; + if (!hadamard_sign_data.empty()) { + const uint32_t width = (uint32_t) weight->ne[0]; + const auto sd = hadamard_sign_data.find(width); + if (sd == hadamard_sign_data.end()) { + throw std::runtime_error(format( + "prism.hadamard has no sign vector for width %u (%s)", width, weight_name.c_str())); + } + const auto key = std::make_pair(width, buft); + auto st = sign_tensors.find(key); + if (st == sign_tensors.end()) { + ggml_init_params params = { + /*.mem_size =*/ ggml_tensor_overhead(), + /*.mem_buffer =*/ NULL, + /*.no_alloc =*/ true, + }; + ggml_context_ptr ctx { ggml_init(params) }; + if (!ctx) { + throw std::runtime_error("failed to create Hadamard sign context"); + } + + ggml_tensor * signs = ggml_new_tensor_1d(ctx.get(), GGML_TYPE_F32, width); + char sign_name[GGML_MAX_NAME]; + snprintf(sign_name, sizeof(sign_name), "prism.hadamard.signs.%u", width); + ggml_set_name(signs, sign_name); + + ggml_backend_buffer_ptr buffer { ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft) }; + if (!buffer) { + throw std::runtime_error(format("unable to allocate %s Hadamard sign buffer", ggml_backend_buft_name(buft))); + } + ggml_backend_buffer_set_usage(buffer.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + + std::vector data(width); + for (uint32_t i = 0; i < width; ++i) { + data[i] = (float) sd->second[i]; + } + ggml_backend_tensor_set(signs, data.data(), 0, data.size() * sizeof(float)); + + std::vector buffers; + buffers.emplace_back(std::move(buffer)); + pimpl->ctxs_bufs.emplace_back(std::move(ctx), std::move(buffers)); + st = sign_tensors.emplace(key, signs).first; + } + sign_tensor = st->second; + } + + llama_hadamard_transform transform { it->tensor, sign_tensor }; + if (hadamard_gdn_v_grouped && weight_name.find(".ssm_out.") != std::string::npos) { + const int64_t n_v = hparams.ssm_dt_rank; + const int64_t n_k = hparams.ssm_n_group; + if (n_k <= 0 || n_v <= 0 || n_v % n_k != 0 || weight->ne[0] % n_v != 0) { + throw std::runtime_error(format("prism.hadamard: bad GDN head geometry for %s", weight_name.c_str())); + } + transform.perm_hd = weight->ne[0] / n_v; + transform.perm_nk = n_k; + transform.perm_rep = n_v / n_k; + } + target->emplace(weight, transform); + } + + LLAMA_LOG_INFO("%s: loaded %zu Hadamard-folded weight(s) (%zu inverse-lookup) using %zu rotation(s) and %zu sign vector(s)\n", + __func__, hadamard_rotations.size() + hadamard_inverses.size(), hadamard_inverses.size(), + rotations.size(), sign_tensors.size()); + } + return true; } diff --git a/src/llama-model.h b/src/llama-model.h index 1230f5244b13..e4198c5e7f65 100644 --- a/src/llama-model.h +++ b/src/llama-model.h @@ -697,6 +697,18 @@ struct llama_model { // gguf metadata std::unordered_map gguf_kv; + // Hadamard-folded GGUF weights are matched with persistent model tensors + // containing the activation-side transform. The string map is populated + // from GGUF metadata while loading hparams; the pointer map is populated + // after model buffers have been allocated. In explicit sign mode the + // per-width sign vectors come from GGUF metadata as well. + std::unordered_map hadamard_weight_blocks; + std::unordered_map hadamard_inverse_blocks; + std::map> hadamard_sign_data; + bool hadamard_gdn_v_grouped = false; + llama_hadamard_rotations hadamard_rotations; + llama_hadamard_rotations hadamard_inverses; + // list of devices used in this model std::vector devices; From dd614c2907d28b7f6f3c208fed50229a80924ab3 Mon Sep 17 00:00:00 2001 From: Andgihat Date: Fri, 18 Sep 2026 15:30:38 +0300 Subject: [PATCH 2/3] ggml : add PQ2_0 and PTQ1_0 quant types (PrismML Bonsai 2) The Bonsai 2 ggufs published on the main PrismML repo use two quant types that are not in upstream ggml: PQ2_0 (type 142, group 128, 2.13 bpw) and PTQ1_0 (type 143, group 128, 1.75 bpw, trits packed base-3 five per byte). Only the dev-repo Q2_0 file (type 42, group 64) loaded before this change. Ported from the PrismML fork: - block layouts, traits and reference quantize/dequantize - CPU vec_dot: generic path plus the x86 VNNI and ARM SIMD variants, with the usual arch-fallback aliases for every other target - CUDA: dequantize, convert, get_rows, MMVQ and MMQ kernels including the PTQ1_0 multi-column path that shares one 128-trit decode across up to three columns, plus the Ampere MMQ tuning table - ftype plumbing: gguf, model loader, quantize tool and gguf-py PTQ1_0 has no SIMD vec_dot on any arch upstream, so it aliases the generic one everywhere; PQ2_0 keeps the x86/ARM implementations. gguf.cpp also learns to recognise a Q2_0 file that is really in the legacy group-128 layout and point at the PQ2_0 build instead of failing with a bare tensor-offset mismatch. Not ported: the DGX Spark (GB10) tile double-buffering and L2 prefetch paths, which have no counterpart here, and the Metal/Vulkan/CDNA/RDNA backends. Verified on a 5060 Ti (sm_120, CUDA 12.8) with Ternary-Bonsai-2-27B-PQ2_0: test-backend-ops MUL_MAT 94/94 and GET_ROWS 8/8 against the CPU reference, correct generation at 47.3 t/s, and no regression on the existing type-42 file. --- ggml/include/ggml.h | 8 +- ggml/src/ggml-common.h | 28 +++ ggml/src/ggml-cpu/arch-fallback.h | 43 ++++ ggml/src/ggml-cpu/arch/arm/quants.c | 71 +++++++ ggml/src/ggml-cpu/arch/x86/quants.c | 76 +++++++ ggml/src/ggml-cpu/ggml-cpu.c | 12 ++ ggml/src/ggml-cpu/ops.cpp | 14 ++ ggml/src/ggml-cpu/quants.c | 124 ++++++++++++ ggml/src/ggml-cpu/quants.h | 6 + ggml/src/ggml-cuda/common.cuh | 41 ++++ ggml/src/ggml-cuda/convert.cu | 117 +++++++++++ ggml/src/ggml-cuda/dequantize.cuh | 28 +++ ggml/src/ggml-cuda/getrows.cu | 8 + ggml/src/ggml-cuda/ggml-cuda.cu | 4 + ggml/src/ggml-cuda/mmq-config-ampere.cuh | 29 +++ ggml/src/ggml-cuda/mmq-load-tiles.cuh | 180 +++++++++++++++++ ggml/src/ggml-cuda/mmq.cu | 27 +++ ggml/src/ggml-cuda/mmq.cuh | 40 ++++ ggml/src/ggml-cuda/mmvq.cu | 60 ++++++ .../template-instances/generate_cu_files.py | 1 + .../template-instances/mmq-instance-pq2_0.cu | 5 + .../template-instances/mmq-instance-ptq1_0.cu | 7 + ggml/src/ggml-cuda/vecdotq.cuh | 190 ++++++++++++++++++ ggml/src/ggml-quants.c | 179 +++++++++++++++++ ggml/src/ggml-quants.h | 6 + ggml/src/ggml.c | 20 ++ ggml/src/gguf.cpp | 19 ++ gguf-py/gguf/constants.py | 6 + include/llama.h | 3 + src/llama-model-loader.cpp | 6 + src/llama-quant.cpp | 6 +- tests/test-backend-ops.cpp | 3 + tools/quantize/quantize.cpp | 2 + 33 files changed, 1367 insertions(+), 2 deletions(-) create mode 100644 ggml/src/ggml-cuda/template-instances/mmq-instance-pq2_0.cu create mode 100644 ggml/src/ggml-cuda/template-instances/mmq-instance-ptq1_0.cu diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 215e8ec42203..00e08d74c333 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -438,7 +438,11 @@ extern "C" { GGML_TYPE_Q3_1 = 46, GGML_TYPE_Q2_0S = 47, GGML_TYPE_Q2_1 = 48, - GGML_TYPE_COUNT = 49, + // Prism-private formats from PrismML's fork, kept at their upstream-fork ids so + // Bonsai 2 files load unchanged. type_traits is sized to COUNT, 49..141 stay unused. + GGML_TYPE_PQ2_0 = 142, // ternary, group 128 (upstream Q2_0 is group 64) + GGML_TYPE_PTQ1_0 = 143, // ternary, group 128, trits packed base-3, five per byte + GGML_TYPE_COUNT = 144, }; // precision @@ -499,6 +503,8 @@ extern "C" { GGML_FTYPE_MOSTLY_NVFP4 = 26, // except 1d tensors GGML_FTYPE_MOSTLY_Q1_0 = 27, // except 1d tensors GGML_FTYPE_MOSTLY_Q2_0 = 28, // except 1d tensors + GGML_FTYPE_MOSTLY_PQ2_0 = 128, // except 1d tensors (Prism-private group-128 Q2_0) + GGML_FTYPE_MOSTLY_PTQ1_0 = 129, // except 1d tensors (Prism-private group-128 ternary) }; // available tensor operations: diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 049e842daf4f..5f06fca9d390 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -99,6 +99,11 @@ typedef sycl::half2 ggml_half2; #define QI2_0 (QK2_0 / 32) #define QR2_0 1 +#define QI_PQ2_0 (QK_PQ2_0 / 32) +#define QR_PQ2_0 1 +#define QI_PTQ1_0 (QK_PTQ1_0 / 32) +#define QR_PTQ1_0 1 + #define QI4_0 (QK4_0 / (4 * QR4_0)) #define QR4_0 2 @@ -191,6 +196,29 @@ typedef sycl::half2 ggml_half2; #ifdef _MSC_VER #define GGML_EXTENSION +// PQ2_0: Prism-private Q2_0 at group size 128. Same 2-bit codec as Q2_0 +// (group 64) but one fp16 scale per 128 weights (~5% smaller). Distinct ggml +// type (142) so it coexists with upstream's group-64 Q2_0 (type 42). +#define QK_PQ2_0 128 +typedef struct { + ggml_half d; // delta (scale) + uint8_t qs[QK_PQ2_0 / 4]; // 2 bits per element +} block_pq2_0; +static_assert(sizeof(block_pq2_0) == sizeof(ggml_half) + QK_PQ2_0 / 4, "wrong pq2_0 block size/padding"); + +// PTQ1_0: Prism-private ternary at group size 128. Same base-3 trit packing as +// upstream TQ1_0 (type 34) but one fp16 scale per 128 weights instead of per 256. +// 1.75 bpw vs PQ2_0's 2.125, and lossless for checkpoints that are already ternary +// at group 128 -- TQ1_0 cannot represent those, because a 256-wide scale has to +// discard one of the two group scales it straddles. +#define QK_PTQ1_0 128 +typedef struct { + uint8_t qs[(QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5]; // 24 B, 5 trits per byte -> 120 values + uint8_t qh[QK_PTQ1_0/64]; // 2 B, 4 trits per byte -> 8 values + ggml_half d; // scale +} block_ptq1_0; +static_assert(sizeof(block_ptq1_0) == sizeof(ggml_half) + QK_PTQ1_0/64 + (QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5, "wrong ptq1_0 block size/padding"); + #else // _MSC_VER #define GGML_EXTENSION __extension__ #endif // _MSC_VER diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index 152e0bac99b0..878ca10c4b27 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -17,6 +17,8 @@ #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K @@ -55,6 +57,9 @@ #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 #define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0 #define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0 #define ggml_gemm_q4_0_8x8_q8_0_generic ggml_gemm_q4_0_8x8_q8_0 @@ -72,6 +77,8 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // repack.cpp #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8 @@ -82,6 +89,8 @@ #define ggml_gemm_mxfp4_8x8_q8_0_generic ggml_gemm_mxfp4_8x8_q8_0 #define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K #elif defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 // repack.cpp @@ -110,6 +119,9 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__POWERPC__) || defined(__powerpc__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 // ref: https://github.com/ggml-org/llama.cpp/pull/14146#issuecomment-2972561679 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K @@ -140,6 +152,9 @@ #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 #define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0 #define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0 #define ggml_gemm_q4_0_8x8_q8_0_generic ggml_gemm_q4_0_8x8_q8_0 @@ -157,6 +172,8 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__loongarch64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K @@ -165,6 +182,7 @@ #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 @@ -187,6 +205,9 @@ #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 #define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0 #define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0 #define ggml_gemm_q4_0_8x8_q8_0_generic ggml_gemm_q4_0_8x8_q8_0 @@ -204,6 +225,9 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__riscv) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 @@ -228,6 +252,9 @@ #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 #define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0 #define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0 #define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K @@ -244,6 +271,9 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__s390x__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 @@ -277,6 +307,10 @@ #define ggml_gemv_iq4_nl_4x4_q8_0_generic ggml_gemv_iq4_nl_4x4_q8_0 #define ggml_gemv_iq4_nl_8x8_q8_0_generic ggml_gemv_iq4_nl_8x8_q8_0 #define ggml_gemv_mxfp4_4x4_q8_0_generic ggml_gemv_mxfp4_4x4_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 +#define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0 #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0 @@ -297,6 +331,8 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__wasm__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K @@ -308,6 +344,7 @@ #define ggml_vec_dot_iq1_s_q8_K_generic ggml_vec_dot_iq1_s_q8_K #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K #define ggml_vec_dot_iq4_nl_q8_0_generic ggml_vec_dot_iq4_nl_q8_0 +#define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 #define ggml_vec_dot_iq4_xs_q8_K_generic ggml_vec_dot_iq4_xs_q8_K #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 @@ -330,6 +367,9 @@ #define ggml_gemv_q6_K_8x8_q8_K_generic ggml_gemv_q6_K_8x8_q8_K #define ggml_gemv_iq4_nl_4x4_q8_0_generic ggml_gemv_iq4_nl_4x4_q8_0 #define ggml_gemv_iq4_nl_8x8_q8_0_generic ggml_gemv_iq4_nl_8x8_q8_0 +#define ggml_gemv_q1_0_4x4_q8_0_generic ggml_gemv_q1_0_4x4_q8_0 +#define ggml_gemv_q1_0_4x8_q8_0_generic ggml_gemv_q1_0_4x8_q8_0 +#define ggml_gemv_pq2_0_4x8_q8_0_generic ggml_gemv_pq2_0_4x8_q8_0 #define ggml_gemv_mxfp4_4x4_q8_0_generic ggml_gemv_mxfp4_4x4_q8_0 #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0 @@ -346,6 +386,9 @@ #define ggml_gemm_q6_K_8x8_q8_K_generic ggml_gemm_q6_K_8x8_q8_K #define ggml_gemm_iq4_nl_4x4_q8_0_generic ggml_gemm_iq4_nl_4x4_q8_0 #define ggml_gemm_iq4_nl_8x8_q8_0_generic ggml_gemm_iq4_nl_8x8_q8_0 +#define ggml_gemm_q1_0_4x4_q8_0_generic ggml_gemm_q1_0_4x4_q8_0 +#define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 +#define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #define ggml_gemm_mxfp4_4x4_q8_0_generic ggml_gemm_mxfp4_4x4_q8_0 #define ggml_gemm_mxfp4_8x8_q8_0_generic ggml_gemm_mxfp4_8x8_q8_0 #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 diff --git a/ggml/src/ggml-cpu/arch/arm/quants.c b/ggml/src/ggml-cpu/arch/arm/quants.c index b988abf9963a..b90afd897776 100644 --- a/ggml/src/ggml-cpu/arch/arm/quants.c +++ b/ggml/src/ggml-cpu/arch/arm/quants.c @@ -219,6 +219,77 @@ void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const voi #endif } +void ggml_vec_dot_pq2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PQ2_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_pq2_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + +#if defined(__ARM_NEON) + // same 2-bit codec as Q2_0, one fp16 scale per 128 elements (4 q8_0 sub-blocks) + // replicate pattern: each byte repeated 4 times + static const uint8_t tbl_idx_lo[16] = {0,0,0,0, 1,1,1,1, 2,2,2,2, 3,3,3,3}; + static const uint8_t tbl_idx_hi[16] = {4,4,4,4, 5,5,5,5, 6,6,6,6, 7,7,7,7}; + // right-shift amounts: 0,2,4,6 repeated for each group of 4 + static const int8_t shift_vals[16] = {0,-2,-4,-6, 0,-2,-4,-6, 0,-2,-4,-6, 0,-2,-4,-6}; + + const uint8x16_t idx_lo = vld1q_u8(tbl_idx_lo); + const uint8x16_t idx_hi = vld1q_u8(tbl_idx_hi); + const int8x16_t shifts = vld1q_s8(shift_vals); + const uint8x16_t mask2 = vdupq_n_u8(0x03); + const int8x16_t one = vdupq_n_s8(1); + + float32x4_t sumv = vdupq_n_f32(0.0f); + + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + + const uint8x8_t raw = vld1_u8(&x[i].qs[k * 8]); + const uint8x16_t raw16 = vcombine_u8(raw, raw); + + uint8x16_t bytes0 = vqtbl1q_u8(raw16, idx_lo); + int8x16_t qv0 = vsubq_s8( + vreinterpretq_s8_u8(vandq_u8(vshlq_u8(bytes0, shifts), mask2)), + one); + + uint8x16_t bytes1 = vqtbl1q_u8(raw16, idx_hi); + int8x16_t qv1 = vsubq_s8( + vreinterpretq_s8_u8(vandq_u8(vshlq_u8(bytes1, shifts), mask2)), + one); + + const int8x16_t y0 = vld1q_s8(yb->qs); + const int8x16_t y1 = vld1q_s8(yb->qs + 16); + + int32x4_t p0 = ggml_vdotq_s32(vdupq_n_s32(0), qv0, y0); + int32x4_t p1 = ggml_vdotq_s32(p0, qv1, y1); + + sumv = vmlaq_n_f32(sumv, vcvtq_f32_s32(p1), d0 * d1); + } + } + + sumf = vaddvq_f32(sumv); +#else + ggml_vec_dot_pq2_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc); + return; +#endif + + *s = sumf; +} + void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK2_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c index ea54cfe44ce4..6cdc754b27f3 100644 --- a/ggml/src/ggml-cpu/arch/x86/quants.c +++ b/ggml/src/ggml-cpu/arch/x86/quants.c @@ -698,6 +698,82 @@ void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const voi #endif } +#if defined(__AVX512VNNI__) && defined(__AVX512VL__) +# define GGML_DPBUSD_256 _mm256_dpbusd_epi32 +#elif defined(__AVXVNNI__) +# define GGML_DPBUSD_256 _mm256_dpbusd_avx_epi32 +#endif + +void ggml_vec_dot_pq2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PQ2_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_pq2_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + +#if (defined(__AVX512VNNI__) && defined(__AVX512VL__)) || defined(__AVXVNNI__) + // AVX-VNNI / AVX-512-VNNI: unpack 2-bit codes c in {0,1,2,3} (value = c-1), then + // dot((c-1), qy) = dpbusd(c, qy) - dpbusd(1, qy). Group 128 = four q8_0 sub-blocks. + const __m256i ones = _mm256_set1_epi8(1); + const __m128i idxlo = _mm_setr_epi8(0,0,0,0,1,1,1,1,2,2,2,2,3,3,3,3); + const __m128i idxhi = _mm_setr_epi8(4,4,4,4,5,5,5,5,6,6,6,6,7,7,7,7); + const __m256i mul = _mm256_setr_epi16(64,16,4,1, 64,16,4,1, 64,16,4,1, 64,16,4,1); // <<(6-2c) + const __m256i three = _mm256_set1_epi16(3); + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + float sumi = 0.0f; + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + const __m256i qy = _mm256_loadu_si256((const __m256i *) yb->qs); + const __m128i src = _mm_loadl_epi64((const __m128i *) &x[i].qs[k * 8]); // 8 bytes + const __m256i rep = _mm256_set_m128i(_mm_shuffle_epi8(src, idxhi), _mm_shuffle_epi8(src, idxlo)); + __m256i r0 = _mm256_cvtepu8_epi16(_mm256_castsi256_si128(rep)); + __m256i r1 = _mm256_cvtepu8_epi16(_mm256_extracti128_si256(rep, 1)); + r0 = _mm256_and_si256(_mm256_srli_epi16(_mm256_mullo_epi16(r0, mul), 6), three); + r1 = _mm256_and_si256(_mm256_srli_epi16(_mm256_mullo_epi16(r1, mul), 6), three); + __m256i codes = _mm256_permute4x64_epi64(_mm256_packus_epi16(r0, r1), 0xD8); + const int dp = hsum_i32_8(GGML_DPBUSD_256(_mm256_setzero_si256(), codes, qy)); + const int sy = hsum_i32_8(GGML_DPBUSD_256(_mm256_setzero_si256(), ones, qy)); + sumi += d1 * (float)(dp - sy); + } + sumf += d0 * sumi; + } +#else + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + float sumi = 0.0f; + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + int sumi_block = 0; + const uint8_t * GGML_RESTRICT qs = &x[i].qs[k * 8]; + const int8_t * GGML_RESTRICT qy = yb->qs; + for (int b = 0; b < 8; ++b) { + const uint8_t byte = qs[b]; + sumi_block += ((int)((byte >> 0) & 3) - 1) * qy[b*4 + 0]; + sumi_block += ((int)((byte >> 2) & 3) - 1) * qy[b*4 + 1]; + sumi_block += ((int)((byte >> 4) & 3) - 1) * qy[b*4 + 2]; + sumi_block += ((int)((byte >> 6) & 3) - 1) * qy[b*4 + 3]; + } + sumi += d1 * sumi_block; + } + sumf += d0 * sumi; + } +#endif + + *s = sumf; +} + void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 0858afccd6c1..27abc4a76a92 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -237,6 +237,18 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_0, .nrows = 1, }, + [GGML_TYPE_PQ2_0] = { + .from_float = quantize_row_pq2_0, + .vec_dot = ggml_vec_dot_pq2_0_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, + [GGML_TYPE_PTQ1_0] = { + .from_float = quantize_row_ptq1_0, + .vec_dot = ggml_vec_dot_ptq1_0_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_Q6_0] = { .from_float = quantize_row_q6_0, .vec_dot = ggml_vec_dot_q6_0_q8_0, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 83311671c936..609f55bf664f 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -669,6 +669,8 @@ void ggml_compute_forward_add( } break; case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1120,6 +1122,8 @@ void ggml_compute_forward_add1( } break; case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1251,6 +1255,8 @@ void ggml_compute_forward_acc( case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4663,6 +4669,8 @@ void ggml_compute_forward_out_prod( switch (src0->type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4945,6 +4953,8 @@ void ggml_compute_forward_set( case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5244,6 +5254,8 @@ void ggml_compute_forward_get_rows( switch (src0->type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -6040,6 +6052,8 @@ void ggml_compute_forward_clamp( case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 74b7f07b4358..8afe7155e4b6 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -30,6 +30,14 @@ void quantize_row_q2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in quantize_row_q2_0_ref(x, y, k); } +void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_pq2_0_ref(x, y, k); +} + +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_ptq1_0_ref(x, y, k); +} + void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { quantize_row_q4_0_ref(x, y, k); } @@ -246,6 +254,122 @@ void ggml_vec_dot_q2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, c *s = sumf; } +// PQ2_0: 128 weights per block = four Q8_0 blocks (4 * 32). No arch defines +// a SIMD variant yet, so this scalar path is the symbol referenced by the traits. +void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PQ2_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_pq2_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + + float sumi = 0.0f; + + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + int sumi_block = 0; + + const uint8_t * GGML_RESTRICT qs = &x[i].qs[k * 8]; + const int8_t * GGML_RESTRICT qy = yb->qs; + + for (int b = 0; b < 8; ++b) { + const uint8_t byte = qs[b]; + sumi_block += ((int)((byte >> 0) & 3) - 1) * qy[b*4 + 0]; + sumi_block += ((int)((byte >> 2) & 3) - 1) * qy[b*4 + 1]; + sumi_block += ((int)((byte >> 4) & 3) - 1) * qy[b*4 + 2]; + sumi_block += ((int)((byte >> 6) & 3) - 1) * qy[b*4 + 3]; + } + + sumi += d1 * sumi_block; + } + + sumf += d0 * sumi; + } + + *s = sumf; +} + +// PTQ1_0 x Q8_0. The trits are stored base-3 interleaved rather than in element +// order, so decode a block into element order first using the same traversal as +// dequantize_row_ptq1_0 -- that keeps the two provably in step. Four Q8_0 blocks +// cover one 128-wide PTQ1_0 block. +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PTQ1_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_ptq1_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + static const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + static const size_t stages[3] = {32, 16, 8}; + + float sumf = 0.0f; + + for (int i = 0; i < nb; i++) { + int8_t q[QK_PTQ1_0]; + int o = 0; + + size_t j = 0; + for (size_t st = 0; st < 3; ++st) { + const size_t c = stages[st]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t nn = 0; nn < 5; ++nn) { + for (size_t m = 0; m < c; ++m) { + const uint8_t v = x[i].qs[j + m] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + } + } + for (size_t nn = 0; nn < 4; ++nn) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + const uint8_t v = x[i].qh[h] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + assert(o == QK_PTQ1_0); + + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + float sumi = 0.0f; + + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + int sumi_block = 0; + for (int b = 0; b < 32; ++b) { + sumi_block += (int) q[k*32 + b] * (int) yb->qs[b]; + } + sumi += d1 * sumi_block; + } + + sumf += d0 * sumi; + } + + *s = sumf; +} + void ggml_vec_dot_q4_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index 4de5b93ff67b..c788b17b3237 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -14,6 +14,8 @@ extern "C" { // Quantization void quantize_row_q1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -40,6 +42,8 @@ void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_tq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void ggml_vec_dot_pq2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -73,6 +77,8 @@ void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo void ggml_vec_dot_iq2_xxs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_iq2_xs_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_iq2_s_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_iq3_xxs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_iq1_s_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_iq1_m_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 54f82d86aef3..b156526d77a4 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -992,6 +992,47 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI2_0; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_PQ2_0; + static constexpr int qr = QR_PQ2_0; + static constexpr int qi = QI_PQ2_0; +}; + +// PTQ1_0 packs trits base-3, five per byte, so element order is not positional. It +// follows the CPU codec exactly: a 16-byte chunk of qs, then an 8-byte chunk, then qh +// at four trits per byte. Trits come out by the base-3 remainder recurrence +// t = (v*3)>>8 with v = (v*3)&0xFF, two integer ops per step and no table. +static __device__ __forceinline__ int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; + int n; + if (e < 80) { // qs[0..15], chunk of 16 + b = x->qs[e & 15]; n = e >> 4; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + b = x->qs[16 + (t & 7)]; n = t >> 3; + } else { // qh[0..1], four trits per byte + const int t = e - 120; + b = x->qh[t & 1]; n = t >> 1; + } + + uint32_t v = b; +#pragma unroll + for (int i = 0; i < 4; ++i) { + if (i < n) { + v = (v * 3) & 0xFF; + } + } + return (int) ((v * 3) >> 8) - 1; +} + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_PTQ1_0; + static constexpr int qr = QR_PTQ1_0; + static constexpr int qi = QI_PTQ1_0; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK4_0; diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 69a4d5fc5783..63bbfc771d85 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -136,6 +136,87 @@ static __global__ void dequantize_block_q4_1(const void * __restrict__ vx, dst_t } } +#if !defined(GGML_USE_HIP) +template +static __device__ +__forceinline__ void dequantize_ptq1_0_qs4(uint32_t packed, float d, dst_t * __restrict__ y, int base, int stride) { + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const uint32_t q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); +# pragma unroll + for (int b = 0; b < 4; ++b) { + const int trit = (int8_t) (q >> (8 * b)); + y[base + t * stride + b] = ggml_cuda_cast(d * trit); + } + } +} + +template +static __global__ void dequantize_block_ptq1_0(const block_ptq1_0 * __restrict__ x, + dst_t * __restrict__ y, + int64_t nb) { + constexpr int threads_per_quant_block = 8; + constexpr int quant_blocks_per_cuda_block = CUDA_DEQUANTIZE_BLOCK_SIZE / threads_per_quant_block; + // The shared transpose keeps packed-byte ownership while making the CTA stores contiguous. + __shared__ dst_t dequantized[quant_blocks_per_cuda_block * QK_PTQ1_0]; + + const int64_t ib0 = int64_t(blockIdx.x) * quant_blocks_per_cuda_block; + const int64_t ib = ib0 + threadIdx.x / threads_per_quant_block; + const int lane = threadIdx.x % threads_per_quant_block; + const bool valid = ib < nb; + + const block_ptq1_0 * bq = valid ? x + ib : x; + float d = valid && lane == 7 ? (float) bq->d : 0.0f; + d = __shfl_sync(__activemask(), d, 7, threads_per_quant_block); + + dst_t * out = dequantized + (threadIdx.x / threads_per_quant_block) * QK_PTQ1_0; + if (valid && lane < 4) { + const uint32_t packed = ((const uint32_t *) bq->qs)[lane]; + dequantize_ptq1_0_qs4(packed, d, out, 4 * lane, 16); + } else if (valid && lane < 6) { + const int g = lane - 4; + const uint32_t packed = ((const uint32_t *) (bq->qs + 16))[g]; + dequantize_ptq1_0_qs4(packed, d, out, 80 + 4 * g, 8); + } else if (valid && lane == 6) { + uint32_t v = (uint32_t) bq->qh[0] | ((uint32_t) bq->qh[1] << 16); +# pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + v = w & 0x00FF00FF; + out[120 + 2 * t + 0] = ggml_cuda_cast(d * ((int) ((w >> 8) & 0xFF) - 1)); + out[120 + 2 * t + 1] = ggml_cuda_cast(d * ((int) ((w >> 24) & 0xFF) - 1)); + } + } + + __syncthreads(); + const int nblocks = nb - ib0 < quant_blocks_per_cuda_block ? nb - ib0 : quant_blocks_per_cuda_block; + const int nelements = nblocks * QK_PTQ1_0; + for (int i = threadIdx.x; i < nelements; i += CUDA_DEQUANTIZE_BLOCK_SIZE) { + y[ib0 * QK_PTQ1_0 + i] = dequantized[i]; + } +} + +template +static void dequantize_row_ptq1_0_cuda(const void * __restrict__ vx, + dst_t * __restrict__ y, + const int64_t k, + cudaStream_t stream) { + GGML_ASSERT(k % QK_PTQ1_0 == 0); + constexpr int quant_blocks_per_cuda_block = CUDA_DEQUANTIZE_BLOCK_SIZE / 8; + const int64_t nb = k / QK_PTQ1_0; + const int num_blocks = (nb + quant_blocks_per_cuda_block - 1) / quant_blocks_per_cuda_block; + dequantize_block_ptq1_0<<>>((const block_ptq1_0 *) vx, y, nb); +} +#endif + //================================== k-quants template @@ -510,6 +591,14 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: +#if !defined(GGML_USE_HIP) + return dequantize_row_ptq1_0_cuda; +#else + return dequantize_block_cont_cuda; +#endif case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -567,6 +656,14 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: +#if !defined(GGML_USE_HIP) + return dequantize_row_ptq1_0_cuda; +#else + return dequantize_block_cont_cuda; +#endif case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -639,6 +736,14 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: +#if !defined(GGML_USE_HIP) + return dequantize_row_ptq1_0_cuda; +#else + return dequantize_block_cont_cuda; +#endif case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -710,6 +815,10 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -749,6 +858,10 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -786,6 +899,10 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_PQ2_0: + return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index 4a5e8b62fff7..60b81dd6f20b 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -52,6 +52,34 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in v.y = (c1 - 1) * d; } +static __device__ __forceinline__ void dequantize_pq2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_pq2_0 * x = (const block_pq2_0 *) vx; + + const float d = x[ib].d; + + // Same 2-bit codec as Q2_0 (per-element indexing is block-size independent). + const int byte_index_0 = iqs / 4; + const int bit_offset_0 = (iqs % 4) * 2; + + const int byte_index_1 = (iqs + 1) / 4; + const int bit_offset_1 = ((iqs + 1) % 4) * 2; + + const int c0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 0x3; + const int c1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 0x3; + + v.x = (c0 - 1) * d; + v.y = (c1 - 1) * d; +} + +static __device__ __forceinline__ void dequantize_ptq1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_ptq1_0 * x = (const block_ptq1_0 *) vx; + + const float d = x[ib].d; + + v.x = ptq1_0_trit(&x[ib], iqs + 0) * d; + v.y = ptq1_0_trit(&x[ib], iqs + 1) * d; +} + static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_q4_0 * x = (const block_q4_0 *) vx; diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index b610ec455643..188b67ffc792 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -320,6 +320,14 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); break; + case GGML_TYPE_PQ2_0: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_PTQ1_0: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; case GGML_TYPE_Q2_0: get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 9eef5cfb06e9..576ac663914f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5351,6 +5351,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_F16: case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5410,6 +5412,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_I32: case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 9f9fd197382f..9261c67f65a3 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/ggml/src/ggml-cuda/mmq-config-ampere.cuh @@ -33,6 +33,35 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf CASE(GGML_TYPE_Q2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PQ2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_PTQ1_0, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); CASE(GGML_TYPE_Q4_0, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index a1dd040a0e5c..fa7fedd3020a 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -184,6 +184,186 @@ template static __device__ __forceinline_ } } +// Q2_0 group 128: identical layout to load_tiles_q2_0 with QK/QI/block for the 128-group format. +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_pq2_0( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + + constexpr int blocks_per_iter = MMQ_ITER_K / QK_PQ2_0; + constexpr int threads_per_row = blocks_per_iter * QI_PQ2_0; + constexpr int nrows = warp_size / threads_per_row; + constexpr int scale_entries_per_block = QK_PQ2_0 / QK8_1; + constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; + + const int txi = threadIdx.x % threads_per_row; + const int kbx = txi / QI_PQ2_0; + const int kqsx = txi % QI_PQ2_0; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { + int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; + + if (fallback) { + i = min(i, i_max); + } + + const block_pq2_0 * bxi = (const block_pq2_0 *) x + kbx0 + i*stride + kbx; + const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 4; + + const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; + +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int q = qxi[j]; + + const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0); + const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2); + const int qx = __byte_perm(qe, qo, 0x5140); + const int qy = __byte_perm(qe, qo, 0x7362); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_qs[i*sram_stride + dst_offset + j*2+0] = qx; + x_qs[i*sram_stride + dst_offset + j*2+1] = qy; +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+0] = qx; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*2+1] = qy; +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } + } + + const int ksx = threadIdx.x % scale_entries_per_row; + const int scale_block = ksx / scale_entries_per_block; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nwarps) { + int i = i0 + threadIdx.y; + + if (fallback) { + i = min(i, i_max); + } + + const block_pq2_0 * bxi = (const block_pq2_0 *) x + kbx0 + i*stride + scale_block; + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_df[i*sram_stride + ksx] = bxi->d; +#else + x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + ksx] = bxi->d; +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } +} + +#if !defined(GGML_USE_HIP) +static __device__ +__forceinline__ void ggml_cuda_mmq_decode_ptq1_0_qs4(uint32_t packed, int * __restrict__ dst, int stride) { + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + dst[t * stride] = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + } +} + +template +static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_ptq1_0(const char * __restrict__ x, + int * __restrict__ x_tile, + const int kbx0, + const int i_max, + const int stride) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + +# if defined(TURING_MMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2 * MMQ_TILE_NE_K); +# else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +# endif + + constexpr int blocks_per_iter = MMQ_ITER_K / QK_PTQ1_0; + constexpr int threads_per_block = 8; + constexpr int threads_per_row = blocks_per_iter * threads_per_block; + constexpr int nrows = warp_size / threads_per_row; + + const int txi = threadIdx.x % threads_per_row; + const int kbx = txi / threads_per_block; + const int lane = txi % threads_per_block; + +# pragma unroll + for (int i0 = 0; i0 < I; i0 += nrows * nwarps) { + int i = i0 + threadIdx.y * nrows + threadIdx.x / threads_per_row; + if (fallback) { + i = min(i, i_max); + } + + const block_ptq1_0 * bxi = (const block_ptq1_0 *) x + kbx0 + i * stride + kbx; +# if defined(TURING_MMA_AVAILABLE) + int * row = x_qs + i * sram_stride + kbx * (QK_PTQ1_0 / 4); +# else + int * row = x_qs + i * (2 * MMQ_TILE_NE_K + 1) + kbx * (QK_PTQ1_0 / 4); +# endif + + if (lane < 4) { + ggml_cuda_mmq_decode_ptq1_0_qs4(get_int_b4(bxi->qs, lane), row + lane, 4); + } else if (lane < 6) { + const int g = lane - 4; + ggml_cuda_mmq_decode_ptq1_0_qs4(get_int_b4(bxi->qs + 16, g), row + 20 + g, 2); + } else if (lane == 6) { + uint32_t v = (uint32_t) bxi->qh[0] | ((uint32_t) bxi->qh[1] << 16); +# pragma unroll + for (int t = 0; t < 4; t += 2) { + const uint32_t w0 = v * 3; + v = w0 & 0x00FF00FF; + const uint32_t w1 = v * 3; + v = w1 & 0x00FF00FF; + row[30 + t / 2] = __vsub4(__byte_perm(w0, w1, 0x7531), 0x01010101); + } + } + } + + constexpr int scale_entries_per_block = QK_PTQ1_0 / QK8_1; + constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; + constexpr int rows_per_warp = warp_size / scale_entries_per_row; + const int ksx = threadIdx.x % scale_entries_per_row; + const int scale_block = ksx / scale_entries_per_block; + +# pragma unroll + for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { + int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / scale_entries_per_row; + if (fallback) { + i = min(i, i_max); + } + + const block_ptq1_0 * bxi = (const block_ptq1_0 *) x + kbx0 + i * stride + scale_block; +# if defined(TURING_MMA_AVAILABLE) + x_df[i * sram_stride + ksx] = bxi->d; +# else + x_df[i * (2 * MMQ_TILE_NE_K / QI8_0) + i / (QI8_0 / 2) + ksx] = bxi->d; +# endif + } +} +#endif + template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index e93c2614b7d3..2e11a315b2b3 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -13,6 +13,14 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con case GGML_TYPE_Q2_0: mul_mat_q_case(ctx, args, stream); break; + case GGML_TYPE_PQ2_0: + mul_mat_q_case(ctx, args, stream); + break; +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: + mul_mat_q_case(ctx, args, stream); + break; +#endif case GGML_TYPE_Q4_0: mul_mat_q_case(ctx, args, stream); break; @@ -282,8 +290,14 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t bool mmq_supported; switch (type) { +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: + mmq_supported = turing_mma_available(cc); + break; +#endif case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -333,6 +347,19 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t } } +#if !defined(GGML_USE_HIP) + if (type == GGML_TYPE_PTQ1_0) { + // the fp16 dequantize + cuBLAS fallback is the source of PTQ1_0's extra error on CUDA, so + // the MMQ tile path runs at every batch by default; the env var is the A/B knob for + // deployments that prefer cuBLAS's ~7% at pp512 over the accuracy + static const int64_t max_batch = [] { + const char * s = getenv("GGML_CUDA_PTQ1_0_MMQ_MAX_BATCH"); + return s ? (int64_t) atoll(s) : (int64_t) MMQ_PTQ1_0_MAX_BATCH_SIZE; + }(); + return ne11 <= max_batch; + } +#endif + if (turing_mma_available(cc)) { return true; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 6367ec88b8ea..333e2450b90d 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -6,6 +6,7 @@ #include #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. +#define MMQ_PTQ1_0_MAX_BATCH_SIZE (1 << 30) // MMQ at every batch unless GGML_CUDA_PTQ1_0_MMQ_MAX_BATCH lowers it #define MMQ_ITER_K 256 #define MMQ_ITER_K_FP4 512 #define MMQ_NWARPS 8 @@ -61,6 +62,10 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { switch (type_x) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: + case GGML_TYPE_PQ2_0: +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: +#endif return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -425,6 +430,12 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml switch (type) { case GGML_TYPE_Q1_0: return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_Q2_0: return MMQ_DP4A_TXS_Q8_0; + case GGML_TYPE_PQ2_0: + return MMQ_DP4A_TXS_Q8_0; +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: + return MMQ_DP4A_TXS_Q8_0; +#endif case GGML_TYPE_Q4_0: return MMQ_DP4A_TXS_Q4_0; case GGML_TYPE_Q4_1: return MMQ_DP4A_TXS_Q4_1; case GGML_TYPE_Q5_0: return MMQ_DP4A_TXS_Q8_0; @@ -589,6 +600,18 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_q2_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); + case GGML_TYPE_PQ2_0: + return ggml_cuda_mmq_util_funcs( + VDR_PQ2_0_Q8_1_MMQ, + ggml_cuda_mmq_load_tiles_pq2_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, + ggml_cuda_mmq_write_back_dp4a); +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: + return ggml_cuda_mmq_util_funcs(VDR_PTQ1_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_ptq1_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, + ggml_cuda_mmq_write_back_dp4a); +#endif case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( VDR_Q4_0_Q8_1_MMQ, @@ -789,6 +812,19 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_q2_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, ggml_cuda_mmq_write_back_mma); + case GGML_TYPE_PQ2_0: + return ggml_cuda_mmq_util_funcs( + -1, + ggml_cuda_mmq_load_tiles_pq2_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_write_back_mma); +#if !defined(GGML_USE_HIP) + case GGML_TYPE_PTQ1_0: + return ggml_cuda_mmq_util_funcs( + -1, ggml_cuda_mmq_load_tiles_ptq1_0, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_write_back_mma); +#endif case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( -1, @@ -1676,6 +1712,10 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda extern DECL_MMQ_CASE(GGML_TYPE_Q1_0); extern DECL_MMQ_CASE(GGML_TYPE_Q2_0); +extern DECL_MMQ_CASE(GGML_TYPE_PQ2_0); +#if !defined(GGML_USE_HIP) +extern DECL_MMQ_CASE(GGML_TYPE_PTQ1_0); +#endif extern DECL_MMQ_CASE(GGML_TYPE_Q4_0); extern DECL_MMQ_CASE(GGML_TYPE_Q4_1); extern DECL_MMQ_CASE(GGML_TYPE_Q5_0); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 976251a1ba60..a5e7bfe9fc56 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -12,6 +12,8 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) switch (type) { case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1; case GGML_TYPE_Q2_0: return vec_dot_q2_0_q8_1; + case GGML_TYPE_PQ2_0: return vec_dot_pq2_0_q8_1; + case GGML_TYPE_PTQ1_0: return vec_dot_ptq1_0_q8_1; case GGML_TYPE_Q4_0: return vec_dot_q4_0_q8_1; case GGML_TYPE_Q4_1: return vec_dot_q4_1_q8_1; case GGML_TYPE_Q5_0: return vec_dot_q5_0_q8_1; @@ -70,6 +72,8 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) { case GGML_TYPE_IQ2_S: return VDR_IQ2_S_Q8_1_MMVQ; case GGML_TYPE_IQ3_XXS: return VDR_IQ3_XXS_Q8_1_MMVQ; case GGML_TYPE_IQ3_S: return VDR_IQ3_S_Q8_1_MMVQ; + case GGML_TYPE_PQ2_0: return VDR_PQ2_0_Q8_1_MMVQ; + case GGML_TYPE_PTQ1_0: return VDR_PTQ1_0_Q8_1_MMVQ; case GGML_TYPE_IQ4_NL: return VDR_IQ4_NL_Q8_1_MMVQ; case GGML_TYPE_IQ4_XS: return VDR_IQ4_XS_Q8_1_MMVQ; default: return 1; @@ -314,6 +318,11 @@ bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) { if (!ggml_is_quantized(type)) { return false; } +#if !defined(GGML_USE_HIP) + if (type == GGML_TYPE_PTQ1_0 && GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING) { + return ne11 <= 7; + } +#endif // k-quants cost more to decode and mvq redoes that per column, so MMQ wins sooner. // Only list quant-types MMQ supports, others would fall back to cuBLAS. if (GGML_CUDA_CC_IS_NVIDIA(cc) && cc == GGML_CUDA_CC_ADA_LOVELACE) { @@ -540,6 +549,8 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d // Only worth the wider block when it actually retires the K loop in half the trips (Observation) if (ncols_dst == 1 && !small_k && halve_iters) { switch (type) { + case GGML_TYPE_PQ2_0: + return generic + generic / 2; case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -701,6 +712,34 @@ static __global__ void mul_mat_vec_q( // x block quant index when casting the quants to int const int kqs = vdr * (tid % (qi/vdr)); +#if !defined(GGML_USE_HIP) + // PTQ1_0 decodes a whole 128-wide block of trits at once, so a few columns can share that + // decode instead of paying for it per column. + if constexpr (type == GGML_TYPE_PTQ1_0 && ncols_dst > 1 && ncols_dst <= 3) { +# pragma unroll + for (int i = 0; i < rows_per_cuda_block; ++i) { + float dots[ncols_dst]; + vec_dot_ptq1_0_q8_1_multi(vx, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp[j][i] += dots[j]; + } + + if constexpr (has_fusion) { + if (use_gate) { + vec_dot_ptq1_0_q8_1_multi(vgate, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp_gate[j][i] += dots[j]; + } + } + } + } + } else +#endif + { #pragma unroll for (int j = 0; j < ncols_dst; ++j) { #pragma unroll @@ -715,6 +754,7 @@ static __global__ void mul_mat_vec_q( } } } + } } __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; @@ -1071,6 +1111,10 @@ static void mul_mat_vec_q_switch_ncols_dst( // Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle. const int nwarps = calc_nwarps(type, c_ncols_dst, table_id); bool use = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp; + // on CDNA the small_k geometry (rows_per_block = nwarps) is several times slower than the default split-K geometry for these types, and with wave64 it triggers for every K < 4096 projection + if (GGML_CUDA_CC_IS_CDNA(cc) && (type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0)) { + use = false; + } constexpr std::array iq_slow_turing = { GGML_TYPE_IQ3_XXS, @@ -1108,6 +1152,10 @@ static void mul_mat_vec_q_switch_ncols_dst( if (table_id != MMVQ_PARAMETERS_GB10) { return false; } + if (type == GGML_TYPE_PQ2_0 && ncols_x == 6144 && nrows_x == 2048) { + return false; + } + // Expert rows are gathered per token, so a wider block adds reduction work without reuse. if (has_ids) { @@ -1249,6 +1297,18 @@ static void mul_mat_vec_q_switch_type( nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); break; + case GGML_TYPE_PQ2_0: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_PTQ1_0: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; case GGML_TYPE_Q4_0: mul_mat_vec_q_switch_ncols_dst (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index 6b05ed359997..2e0b9229ce62 100755 --- a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -106,6 +106,7 @@ TYPES_MMQ = [ "GGML_TYPE_Q1_0", "GGML_TYPE_Q2_0", + "GGML_TYPE_PQ2_0", "GGML_TYPE_PTQ1_0", "GGML_TYPE_Q2_0S", "GGML_TYPE_Q2_1", "GGML_TYPE_Q3_0", "GGML_TYPE_Q3_1", "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q6_0", "GGML_TYPE_Q6_1", "GGML_TYPE_Q8_0", "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K", diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-pq2_0.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-pq2_0.cu new file mode 100644 index 000000000000..ed0e2c4fef51 --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-pq2_0.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_PQ2_0); diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-ptq1_0.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-ptq1_0.cu new file mode 100644 index 000000000000..692e952986bd --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-ptq1_0.cu @@ -0,0 +1,7 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +#if !defined(GGML_USE_HIP) +DECL_MMQ_CASE(GGML_TYPE_PTQ1_0); +#endif diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index bdb42bc9a551..b9f44c4d09ae 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -110,7 +110,11 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) { #define VDR_Q1_0_Q8_1_MMQ 4 // Q1_0 has 128 bits (4 ints) per block #define VDR_Q2_0_Q8_1_MMVQ 1 // Process one 32-element chunk at a time for parallelism +#define VDR_PQ2_0_Q8_1_MMVQ 1 // one 32-element chunk at a time (same per-chunk codec as Q2_0) +#define VDR_PTQ1_0_Q8_1_MMVQ 4 // whole 128 block per call: keeps the byte walk uniform across lanes #define VDR_Q2_0_Q8_1_MMQ 2 // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks +#define VDR_PQ2_0_Q8_1_MMQ 2 // Q2_0 group 128: 4 32-element chunks per block +#define VDR_PTQ1_0_Q8_1_MMQ 2 // expanded to signed bytes in the MMQ tile loader #define VDR_Q4_0_Q8_1_MMVQ 2 #define VDR_Q4_0_Q8_1_MMQ 4 @@ -948,6 +952,192 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } +#if !defined(GGML_USE_HIP) +template +static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs, + const uint32_t stride_col_y, + float * result) { + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[ncols_dst][4] = {}; + + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes. +# pragma unroll + for (int g = 0; g < 4; ++g) { + const uint32_t packed = get_int_b4(bq->qs, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = t * 16 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + +# pragma unroll + for (int g = 0; g < 2; ++g) { + const uint32_t packed = get_int_b4(bq->qs + 16, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = 80 + t * 8 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + + uint32_t v = (uint32_t) bq->qh[0] | ((uint32_t) bq->qh[1] << 16); +# pragma unroll + for (int t = 0; t < 4; t += 2) { + const uint32_t w0 = v * 3; + v = w0 & 0x00FF00FF; + const uint32_t w1 = v * 3; + v = w1 & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w0, w1, 0x7531), 0x01010101); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + 3].qs, 6 + t / 2); + sumi[j][3] = ggml_cuda_dp4a(q, u, sumi[j][3]); + } + } + + const float d = (float) bq->d; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[j * stride_col_y + iqs + k].ds) * (float) sumi[j][k]; + } + result[j] = d * acc; + } +} +#endif + +// PTQ1_0 x Q8_1. One call consumes the full 128-weight block. +static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { +#if defined(GGML_USE_HIP) + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[4] = { 0, 0, 0, 0 }; + +# pragma unroll + for (int m = 0; m < 16; ++m) { + uint32_t v = bq->qs[m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = t * 16 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int m = 0; m < 8; ++m) { + uint32_t v = bq->qs[16 + m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 80 + t * 8 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int h = 0; h < 2; ++h) { + uint32_t v = bq->qh[h]; +# pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 120 + t * 2 + h; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[iqs + k].ds) * (float) sumi[k]; + } + return (float) bq->d * acc; +#else + float result; + vec_dot_ptq1_0_q8_1_multi<1>(vbq, bq8_1, kbx, iqs, 0, &result); + return result; +#endif +} + +static __device__ __forceinline__ float vec_dot_pq2_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { + const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq + kbx; + + // Q2_0 group 128: 128 elements, ONE scale, processed as four 32-element chunks + // (iqs selects the chunk, 0-3). Same per-chunk 2-bit codec as Q2_0. + const float d2 = bq2_0->d; + const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4; + + const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; + + int sumi = 0; +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int q = qs[j]; + const int u = get_int_b4(bq8_1_chunk->qs, j*2+0); + const int v = get_int_b4(bq8_1_chunk->qs, j*2+1); + +#if defined(GGML_USE_HIP) && defined(__HIP_DEVICE_COMPILE__) + const int qx = q2_0_symbols4_hip((uint32_t) q & 0xFFu); + const int qy = q2_0_symbols4_hip(((uint32_t) q >> 8) & 0xFFu); +#else + const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0); + const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2); + const int qx = __byte_perm(qe, qo, 0x5140); + const int qy = __byte_perm(qe, qo, 0x7362); +#endif + + sumi = ggml_cuda_dp4a(u, qx, sumi); + sumi = ggml_cuda_dp4a(v, qy, sumi); + } + + const float d8 = __low2float(bq8_1_chunk->ds); + return d2 * d8 * sumi; +} + static __device__ __forceinline__ float vec_dot_q4_0_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 8420035e5f5e..babe7226713d 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -36,6 +36,41 @@ static inline int best_index_int8(int n, const int8_t * val, float x) { return x - val[mu-1] < val[mu] - x ? mu-1 : mu; } +// PQ2_0: identical 2-bit codec to Q2_0, one fp16 scale per 128 weights. +void quantize_row_pq2_0_ref(const float * GGML_RESTRICT x, block_pq2_0 * GGML_RESTRICT y, int64_t k) { + static const int qk = QK_PQ2_0; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + float amax = 0.0f; + for (int j = 0; j < qk; j++) { + const float a = fabsf(x[i*qk + j]); + if (a > amax) amax = a; + } + const float d = amax; + const float id = d > 0.0f ? 1.0f / d : 0.0f; + + y[i].d = GGML_FP32_TO_FP16(d); + + for (int j = 0; j < qk / 4; ++j) { + y[i].qs[j] = 0; + } + + for (int j = 0; j < qk; ++j) { + const float w = x[i*qk + j]; + int q = (int)roundf(w * id) + 1; + if (q < 0) q = 0; + if (q > 3) q = 3; + const int byte_index = j / 4; + const int bit_offset = (j % 4) * 2; + y[i].qs[byte_index] |= ((uint8_t)q << bit_offset); + } + } +} + // reference implementation for deterministic creation of model files void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k) { static const int qk = QK1_0; @@ -712,6 +747,26 @@ void dequantize_row_q2_0(const block_q2_0 * GGML_RESTRICT x, float * GGML_RESTRI } } +void dequantize_row_pq2_0(const block_pq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + static const int qk = QK_PQ2_0; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + const float d = GGML_FP16_TO_FP32(x[i].d); + + for (int j = 0; j < qk; ++j) { + const int byte_index = j / 4; + const int bit_offset = (j % 4) * 2; + const uint8_t q = (x[i].qs[byte_index] >> bit_offset) & 0x03; + // 00=-1, 01=0, 10=+1, 11=+2 + y[i*qk + j] = ((int)q - 1) * d; + } + } +} + void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { static const int qk = QK4_0; @@ -2506,6 +2561,122 @@ size_t quantize_q2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, return nrow * row_size; } +size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + if (!quant_weights) { + quantize_row_pq2_0_ref(src, dst, (int64_t)nrow*n_per_row); + return nrow * ggml_row_size(GGML_TYPE_PQ2_0, n_per_row); + } + size_t row_size = ggml_row_size(GGML_TYPE_PQ2_0, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_pq2_0_ref(src, (block_pq2_0*)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + +// ====================== PTQ1_0 (Prism ternary, group 128) ====================== +// Base-3 trit packing identical to upstream TQ1_0, but at block 128 so one fp16 +// scale covers 128 weights. qs is 24 bytes, which TQ1_0's fixed 32-then-16 byte +// staging cannot cover, so the stages are generalised to 32/16/8; at TQ1_0's +// 48-byte qs this reduces to exactly its original 32-then-16 behaviour. +static const size_t ptq1_0_stages[3] = {32, 16, 8}; + +void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + for (int64_t i = 0; i < nb; i++) { + float amax = 0.0f; + for (int j = 0; j < QK_PTQ1_0; j++) { + amax = MAX(amax, fabsf(x[j])); + } + + const float d = amax; + const float id = d ? 1.0f/d : 0.0f; + + y[i].d = GGML_FP32_TO_FP16(d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(y->qs); j += c) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = 0; + for (size_t n = 0; n < 5; ++n) { + int xi = lroundf(x[m + n*c] * id) + 1; // -1, 0, 1 -> 0, 1, 2 + q *= 3; + q += xi; + } + // ceiling division (243 == pow(3, 5)) + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qs[j + m] = q; + } + x += 5*c; + } + } + // 4 elements per byte + for (size_t h = 0; h < sizeof(y->qh); ++h) { + uint8_t q = 0; + for (size_t m = 0; m < 4; ++m) { + int xi = lroundf(x[h + m*sizeof(y->qh)] * id) + 1; + q *= 3; + q += xi; + } + // shift the first value to the most significant trit + q *= 3; + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qh[h] = q; + } + x += 4*sizeof(y->qh); + } +} + +void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + + for (int64_t i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t n = 0; n < 5; ++n) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = x[i].qs[j + m] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } + } + for (size_t n = 0; n < 4; ++n) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + uint8_t q = x[i].qh[h] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } +} + +size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void)quant_weights; // ternary codes come from the weights themselves; an imatrix has no role + const size_t row_size = ggml_row_size(GGML_TYPE_PTQ1_0, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_ptq1_0_ref(src, (block_ptq1_0 *)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { if (!quant_weights) { quantize_row_q4_0_ref(src, dst, (int64_t)nrow*n_per_row); @@ -5960,6 +6131,14 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_q2_0, data, nb); } break; + case GGML_TYPE_PQ2_0: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_pq2_0, data, nb); + } break; + case GGML_TYPE_PTQ1_0: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_ptq1_0, data, nb); + } break; case GGML_TYPE_Q4_0: { VALIDATE_ROW_DATA_D_F16_IMPL(block_q4_0, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 4c6557bca5c3..50323ecb3971 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -16,6 +16,8 @@ extern "C" { // Quantization GGML_API void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q2_0_ref(const float * GGML_RESTRICT x, block_q2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_pq2_0_ref(const float * GGML_RESTRICT x, block_pq2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_q4_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_1_ref(const float * GGML_RESTRICT x, block_q4_1 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_RESTRICT y, int64_t k); @@ -45,6 +47,8 @@ GGML_API void quantize_row_iq3_xxs_ref(const float * GGML_RESTRICT x, block_iq3_ GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq4_xs_ref (const float * GGML_RESTRICT x, block_iq4_xs * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_s_ref (const float * GGML_RESTRICT x, block_iq3_s * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_pq2_0(const block_pq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_s * GGML_RESTRICT y, int64_t k); // Dequantization @@ -96,6 +100,8 @@ GGML_API size_t quantize_iq1_m (const float * GGML_RESTRICT src, void * GGML_RE GGML_API size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq3_s (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_tq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_tq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 14685463ce3c..7f74fe5dcf31 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -690,6 +690,22 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_q2_0, .from_float_ref = (ggml_from_float_t) quantize_row_q2_0_ref, }, + [GGML_TYPE_PQ2_0] = { + .type_name = "pq2_0", + .blck_size = QK_PQ2_0, + .type_size = sizeof(block_pq2_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_pq2_0, + .from_float_ref = (ggml_from_float_t) quantize_row_pq2_0_ref, + }, + [GGML_TYPE_PTQ1_0] = { + .type_name = "ptq1_0", + .blck_size = QK_PTQ1_0, + .type_size = sizeof(block_ptq1_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_ptq1_0, + .from_float_ref = (ggml_from_float_t) quantize_row_ptq1_0_ref, + }, [GGML_TYPE_Q6_0] = { .type_name = "q6_0", .blck_size = QK6_0, @@ -1491,6 +1507,8 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_Q4_1: wtype = GGML_TYPE_Q4_1; break; case GGML_FTYPE_MOSTLY_Q1_0: wtype = GGML_TYPE_Q1_0; break; case GGML_FTYPE_MOSTLY_Q2_0: wtype = GGML_TYPE_Q2_0; break; + case GGML_FTYPE_MOSTLY_PQ2_0: wtype = GGML_TYPE_PQ2_0; break; + case GGML_FTYPE_MOSTLY_PTQ1_0: wtype = GGML_TYPE_PTQ1_0; break; case GGML_FTYPE_MOSTLY_Q5_0: wtype = GGML_TYPE_Q5_0; break; case GGML_FTYPE_MOSTLY_Q5_1: wtype = GGML_TYPE_Q5_1; break; case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; @@ -8358,6 +8376,8 @@ size_t ggml_quantize_chunk( switch (type) { case GGML_TYPE_Q1_0: result = quantize_q1_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q2_0: result = quantize_q2_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_PQ2_0: result = quantize_pq2_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_PTQ1_0: result = quantize_ptq1_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_0: result = quantize_q4_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_1: result = quantize_q4_1 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_0: result = quantize_q5_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/ggml/src/gguf.cpp b/ggml/src/gguf.cpp index 6c7b5817812b..d97df4a8deea 100644 --- a/ggml/src/gguf.cpp +++ b/ggml/src/gguf.cpp @@ -774,11 +774,19 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr // compute the total size of the data section, taking into account the alignment { ctx->size = 0; + size_t size_if_legacy_q2 = 0; // see the legacy Q2_0 layout hint below + bool has_q2_0 = false; for (size_t i = 0; i < ctx->info.size(); ++i) { const gguf_tensor_info & ti = ctx->info[i]; if (ti.offset != ctx->size) { GGML_LOG_ERROR("%s: tensor '%s' has offset %" PRIu64 ", expected %zu\n", __func__, ti.t.name, ti.offset, ctx->size); + if (has_q2_0 && ti.offset == size_if_legacy_q2) { + GGML_LOG_ERROR("%s: this file matches the legacy Prism Q2_0 layout (group size 128 stored as ggml type id 42), " + "but this build reads Q2_0 as the official group-64 format\n", __func__); + GGML_LOG_ERROR("%s: you are probably using the wrong GGUF: use the PQ2_0 version of this model (ggml type id 142) " + "or download the group-64 Q2_0 file\n", __func__); + } GGML_LOG_ERROR("%s: failed to read tensor data\n", __func__); gguf_free(ctx); return nullptr; @@ -791,6 +799,17 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr return nullptr; } ctx->size += padded_size; + + // track the size the data section would have if the Q2_0 tensors used the + // legacy Prism group-128 layout, which is byte-identical to PQ2_0 + if (ti.t.type == GGML_TYPE_Q2_0) { + has_q2_0 = true; + const size_t nrows = ggml_nrows(&ti.t); + const size_t nbytes_legacy = ggml_row_size(GGML_TYPE_PQ2_0, ti.t.ne[0]) * nrows; + size_if_legacy_q2 += GGML_PAD(nbytes_legacy, ctx->alignment); + } else { + size_if_legacy_q2 += padded_size; + } } } diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 8c22ecb1c835..df2141e859c4 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -5671,6 +5671,8 @@ class GGMLQuantizationType(IntEnum): NVFP4 = 40 Q1_0 = 41 Q2_0 = 42 + PQ2_0 = 142 + PTQ1_0 = 143 class ExpertGatingFuncType(IntEnum): @@ -5727,6 +5729,8 @@ class LlamaFileType(IntEnum): MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors MOSTLY_Q2_0 = 41 # except 1d tensors + MOSTLY_PQ2_0 = 128 # except 1d tensors + MOSTLY_PTQ1_0 = 129 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -5864,6 +5868,8 @@ class VisionProjectorType: GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), GGMLQuantizationType.Q2_0: (64, 2 + 16), + GGMLQuantizationType.PQ2_0: (128, 2 + 32), + GGMLQuantizationType.PTQ1_0: (128, 2 + 24 + 2), } diff --git a/include/llama.h b/include/llama.h index 6989c853e5b2..0bb156ed68af 100644 --- a/include/llama.h +++ b/include/llama.h @@ -156,6 +156,9 @@ extern "C" { LLAMA_FTYPE_MOSTLY_NVFP4 = 39, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q1_0 = 40, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors + LLAMA_FTYPE_MOSTLY_PQ2_0 = 141, // except 1d tensors (Prism group-128 Q2_0) + LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY = 142, // pre-rename value for the same format, still found in published ggufs + LLAMA_FTYPE_MOSTLY_PTQ1_0 = 143, // except 1d tensors (Prism group-128 ternary, 1.75 bpw) LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 1a11cfc43827..424a42a933c3 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -39,6 +39,10 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_BF16: name = LLAMA_FTYPE_PREFIX "BF16"; break; case LLAMA_FTYPE_MOSTLY_Q1_0: name = LLAMA_FTYPE_PREFIX "Q1_0"; break; case LLAMA_FTYPE_MOSTLY_Q2_0: name = LLAMA_FTYPE_PREFIX "Q2_0"; break; + case LLAMA_FTYPE_MOSTLY_PQ2_0: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128)"; break; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: name = LLAMA_FTYPE_PREFIX "PTQ1_0 - 1.75 bpw (group 128)"; break; + // ggufs packed before the Q2_0_G128 -> PQ2_0 rename carry the old ftype value. + case LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128)"; break; case LLAMA_FTYPE_MOSTLY_Q4_0: name = LLAMA_FTYPE_PREFIX "Q4_0"; break; case LLAMA_FTYPE_MOSTLY_Q4_1: name = LLAMA_FTYPE_PREFIX "Q4_1"; break; case LLAMA_FTYPE_MOSTLY_Q5_0: name = LLAMA_FTYPE_PREFIX "Q5_0"; break; @@ -772,6 +776,8 @@ llama_model_loader::llama_model_loader( case GGML_TYPE_NVFP4: ftype = LLAMA_FTYPE_MOSTLY_NVFP4; break; case GGML_TYPE_Q1_0: ftype = LLAMA_FTYPE_MOSTLY_Q1_0; break; case GGML_TYPE_Q2_0: ftype = LLAMA_FTYPE_MOSTLY_Q2_0; break; + case GGML_TYPE_PQ2_0: ftype = LLAMA_FTYPE_MOSTLY_PQ2_0; break; + case GGML_TYPE_PTQ1_0: ftype = LLAMA_FTYPE_MOSTLY_PTQ1_0; break; default: { LLAMA_LOG_WARN("%s: unknown type %s\n", __func__, ggml_type_name(type_max)); diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 34ff25db57e6..edc031d51ef5 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -391,6 +391,8 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso case GGML_TYPE_IQ3_S: // types on the right: block size 32 case GGML_TYPE_IQ4_XS: return_type = GGML_TYPE_IQ4_NL; break; case GGML_TYPE_Q2_0: + case GGML_TYPE_PTQ1_0: + case GGML_TYPE_PQ2_0: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_TQ1_0: @@ -502,7 +504,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) { new_type = GGML_TYPE_IQ3_S; } - else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0) { + else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_PQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_PTQ1_0) { new_type = GGML_TYPE_Q4_K; } } @@ -856,6 +858,8 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_ALL_F32: return GGML_TYPE_F32; case LLAMA_FTYPE_MOSTLY_Q1_0: return GGML_TYPE_Q1_0; case LLAMA_FTYPE_MOSTLY_Q2_0: return GGML_TYPE_Q2_0; + case LLAMA_FTYPE_MOSTLY_PQ2_0: return GGML_TYPE_PQ2_0; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: return GGML_TYPE_PTQ1_0; case LLAMA_FTYPE_MOSTLY_MXFP4_MOE: return GGML_TYPE_MXFP4; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index e3db6816757b..ab8964a6935c 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9117,6 +9117,7 @@ static const ggml_type all_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_MXFP4, GGML_TYPE_NVFP4, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, @@ -9133,6 +9134,7 @@ static const ggml_type base_types[] = { GGML_TYPE_Q8_0, // for I8MM tests GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, // for I8MM tests GGML_TYPE_Q4_K, @@ -9146,6 +9148,7 @@ static const ggml_type other_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index 38950036cd82..2f1525e16e6f 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -34,6 +34,8 @@ struct quant_option { static const std::vector QUANT_OPTIONS = { { "Q1_0", LLAMA_FTYPE_MOSTLY_Q1_0, " 1.125 bpw quantization", }, { "Q2_0", LLAMA_FTYPE_MOSTLY_Q2_0, " 2.25 bpw quantization (group 64)", }, + { "PQ2_0", LLAMA_FTYPE_MOSTLY_PQ2_0, " 2.13 bpw quantization (group 128, Prism)", }, + { "PTQ1_0", LLAMA_FTYPE_MOSTLY_PTQ1_0, " 1.75 bpw ternarization (group 128, Prism)", }, { "Q4_0", LLAMA_FTYPE_MOSTLY_Q4_0, " 4.34G, +0.4685 ppl @ Llama-3-8B", }, { "Q4_1", LLAMA_FTYPE_MOSTLY_Q4_1, " 4.78G, +0.4511 ppl @ Llama-3-8B", }, { "MXFP4_MOE",LLAMA_FTYPE_MOSTLY_MXFP4_MOE," MXFP4 MoE", }, From 477afefd374807422779276feec6bc8243f82ab4 Mon Sep 17 00:00:00 2001 From: Andgihat Date: Sat, 19 Sep 2026 22:48:21 +0300 Subject: [PATCH 3/3] qwen35 : apply the Hadamard inverse to the MTP token-embedding lookup The MTP graph does its own row lookup on the embedding table and never restores the primal basis, while the trunk (build_inp_embd) does. On a prism.hadamard model that mismatch trips the graph check: Hadamard-latent table 'token_embd.weight' is read without the inverse transform so a Bonsai 2 gguf with an embedded MTP block refuses to load at all. Factor the inverse out of build_inp_embd into build_hadamard_inverse_embd and call it from both paths, against whichever table was actually read (layer.nextn.embed_tokens, else model.tok_embd). Models without folded weights have an empty inverse map, so this is a no-op for them. Same fix as PrismML-Eng/llama.cpp#205, adapted to this tree. Measured on Ternary-Bonsai-2-27B-PQ2_0-MTP-Q8_0 (RTX 5060 Ti, sm_120, CUDA 12.8, greedy, 4K context): the file now loads, and decode goes from 37.4 t/s with the drafter off to 71.1 t/s at --spec-draft-n-max 8, with draft acceptance 0.247 and mean draft length 2.98. The vendor default of 2 leaves ~13% on the table here (61.7 t/s); the curve peaks at 8 and falls off by 16 (52.9 t/s). Speculative decoding verifies every draft against the target, so output is unchanged. --- src/llama-graph.cpp | 35 ++++++++++++++++++++++++----------- src/llama-graph.h | 5 +++++ src/models/qwen35.cpp | 1 + src/models/qwen35moe.cpp | 1 + 4 files changed, 31 insertions(+), 11 deletions(-) diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 2680dcf24478..1c1002d1bf2e 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1904,6 +1904,29 @@ ggml_tensor * llm_graph_context::build_cvec( return cvec->apply_to(ctx0, cur, il); } +// A Hadamard-latent embedding table stores rotated rows; restore the primal basis +// right after the lookup: h = s * (H z). Both the trunk and the MTP graph do their +// own row lookup, so both have to undo it. +ggml_tensor * llm_graph_context::build_hadamard_inverse_embd( + ggml_tensor * w, + ggml_tensor * cur) const { + if (!hadamard_inverses) { + return cur; + } + + const auto it = hadamard_inverses->find(w); + if (it == hadamard_inverses->end()) { + return cur; + } + + cur = llama_mul_mat_hadamard(ctx0, cur, it->second.rot); + if (it->second.signs) { + cur = ggml_mul(ctx0, cur, it->second.signs); + } + + return cur; +} + ggml_tensor * llm_graph_context::build_hadamard_activation( ggml_tensor * w, ggml_tensor * cur) const { @@ -2822,17 +2845,7 @@ ggml_tensor * llm_graph_context::build_inp_embd(ggml_tensor * tok_embd) const { cur = ggml_get_rows(ctx0, tok_embd, inp->tokens); - // a Hadamard-latent embedding table stores rotated rows; restore the - // primal basis right after the lookup: h = s * (H z) - if (hadamard_inverses) { - const auto it = hadamard_inverses->find(tok_embd); - if (it != hadamard_inverses->end()) { - cur = llama_mul_mat_hadamard(ctx0, cur, it->second.rot); - if (it->second.signs) { - cur = ggml_mul(ctx0, cur, it->second.signs); - } - } - } + cur = build_hadamard_inverse_embd(tok_embd, cur); // apply lora for embedding tokens if needed for (const auto & lora : *loras) { diff --git a/src/llama-graph.h b/src/llama-graph.h index 152785f8ca50..56e4728dacf3 100644 --- a/src/llama-graph.h +++ b/src/llama-graph.h @@ -1156,6 +1156,11 @@ struct llm_graph_context { // do mat_mul, while optionally apply lora and per-tensor scale // if w is a Hadamard-folded weight, return the activation with its // transform applied (sign flip, then rotation); otherwise return it as is + // restore the primal basis of a Hadamard-latent embedding table after a row lookup + ggml_tensor * build_hadamard_inverse_embd( + ggml_tensor * w, + ggml_tensor * cur) const; + ggml_tensor * build_hadamard_activation( ggml_tensor * w, ggml_tensor * cur) const; diff --git a/src/models/qwen35.cpp b/src/models/qwen35.cpp index a1e263500ee8..0358d086fe7f 100644 --- a/src/models/qwen35.cpp +++ b/src/models/qwen35.cpp @@ -519,6 +519,7 @@ llama_model_qwen35::graph_mtp::graph_mtp(const llama_model & model, const llm_gr ggml_tensor * tok_embd_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd; tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens); + tok_embd = build_hadamard_inverse_embd(tok_embd_w, tok_embd); } else { tok_embd = inp->embd; } diff --git a/src/models/qwen35moe.cpp b/src/models/qwen35moe.cpp index bdf772625093..b9c9a98c7e61 100644 --- a/src/models/qwen35moe.cpp +++ b/src/models/qwen35moe.cpp @@ -583,6 +583,7 @@ llama_model_qwen35moe::graph_mtp::graph_mtp(const llama_model & model, const llm ggml_tensor * tok_embd_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd; tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens); + tok_embd = build_hadamard_inverse_embd(tok_embd_w, tok_embd); } else { tok_embd = inp->embd; }