Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-17 16:25:01 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'upstream' into concedo_experimental

Browse Source 
# Conflicts:
#	.github/workflows/build-ibm.yml
#	.github/workflows/build-self-hosted.yml
#	.github/workflows/code-style.yml
#	.github/workflows/editorconfig.yml
#	.github/workflows/fusion.yml
#	.github/workflows/release.yml
#	.github/workflows/server-sanitize.yml
#	.pi/gg/SYSTEM.md
#	CMakeLists.txt
#	ci/run.sh
#	common/CMakeLists.txt
#	docs/backend/SYCL.md
#	ggml/CMakeLists.txt
#	ggml/src/ggml-cpu/CMakeLists.txt
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-sycl/CMakeLists.txt
#	ggml/src/ggml-sycl/backend.hpp
#	ggml/src/ggml-sycl/base.hpp
#	ggml/src/ggml-sycl/common.hpp
#	ggml/src/ggml-sycl/gemm.hpp
#	ggml/src/ggml-sycl/ggml-sycl.cpp
#	ggml/src/ggml-sycl/mem.cpp
#	scripts/sync-ggml.last
#	src/CMakeLists.txt
#	tests/CMakeLists.txt
#	tests/snapshots/nvidia-nemotron-3-nano-30b-a3b.schema
#	tests/test-backend-ops.cpp
#	tests/test-chat.cpp
#	tests/test-llama-archs.cpp
#	tests/test-quant-type-selection.cpp
#	tools/mtmd/CMakeLists.txt
#	tools/server/CMakeLists.txt
This commit is contained in:
Concedo
2026-09-14 22:44:04 +08:00
parent a4f7e70b9f 41abbfd599
commit c43b36bcb5
39 changed files with 503 additions and 1530 deletions
Download Patch File Download Diff File Expand all files Collapse all files
cmake/arm64-windows-msvc-cuda.cmake
-26
View File
@@ -1,26 +0,0 @@
# Used to cross-compile ggml-cuda for Windows ARM64 on an x64 Windows host.
set( CMAKE_SYSTEM_NAME Windows )
set( CMAKE_SYSTEM_PROCESSOR arm64 )
if ( DEFINED CUDAToolkit_ROOT )
file( TO_CMAKE_PATH "${CUDAToolkit_ROOT}" CUDA_ROOT )
elseif ( DEFINED ENV{CUDA_PATH} )
file( TO_CMAKE_PATH "$ENV{CUDA_PATH}" CUDA_ROOT )
else()
message( FATAL_ERROR "Set CUDAToolkit_ROOT or CUDA_PATH to a Windows CUDA Toolkit with ARM64 target libraries" )
endif()
if ( DEFINED ENV{VCToolsInstallDir} )
file( TO_CMAKE_PATH "$ENV{VCToolsInstallDir}" MSVC_TOOLS_ROOT )
set( CMAKE_CUDA_HOST_COMPILER "${MSVC_TOOLS_ROOT}/bin/Hostx64/arm64/cl.exe" CACHE FILEPATH "" )
endif()
set( CMAKE_CUDA_COMPILER "${CUDA_ROOT}/bin/nvcc.exe" CACHE FILEPATH "" )
set( CMAKE_CUDA_FLAGS_INIT "-target-dir=arm64" )
# FindCUDAToolkit selects lib/x64 from the host architecture on Windows.
set( CUDA_CUDART "${CUDA_ROOT}/lib/arm64/cudart.lib" CACHE FILEPATH "" )
set( CUDA_cudart_LIBRARY "${CUDA_ROOT}/lib/arm64/cudart.lib" CACHE FILEPATH "" )
set( CUDA_cublas_LIBRARY "${CUDA_ROOT}/lib/arm64/cublas.lib" CACHE FILEPATH "" )
set( CUDA_cublasLt_LIBRARY "${CUDA_ROOT}/lib/arm64/cublasLt.lib" CACHE FILEPATH "" )
set( CUDA_cuda_driver_LIBRARY "${CUDA_ROOT}/lib/arm64/cuda.lib" CACHE FILEPATH "" )
common/arg.cpp
+1 -1
View File
@@ -3876,7 +3876,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
{"--no-log-jsonl"},
"Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)",
[](common_params &, bool value) {
common_log_set_jsonl(common_log_main(), value);
common_log_set_jsonl(value);
}
).set_env("LLAMA_ARG_LOG_JSONL"));
add_opt(common_arg(
common/common.cpp
+5 -6
View File
@@ -1593,6 +1593,11 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
return COMMON_CONTEXT_SEQ_RM_TYPE_NO;
}
if (llama_n_rs_seq(ctx) > 0) {
COM_TRC("%s", "the context supports bounded partial sequence removal\n");
return COMMON_CONTEXT_SEQ_RM_TYPE_RS;
}
common_context_seq_rm_type res = COMMON_CONTEXT_SEQ_RM_TYPE_PART;
llama_memory_clear(mem, true);
@@ -1609,12 +1614,6 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
goto done;
}
if (llama_n_rs_seq(ctx) > 0) {
COM_TRC("%s", "the context supports bounded partial sequence removal\n");
res = COMMON_CONTEXT_SEQ_RM_TYPE_RS;
goto done;
}
// try to remove the last tokens
if (!llama_memory_seq_rm(mem, 0, 1, -1)) {
COM_TRC("%s", "the context does not support partial sequence removal\n");
common/fit.cpp
+41
View File
@@ -1,5 +1,6 @@
#include "fit.h"
#include "json.h"
#include "log.h"
#include "../src/llama-ext.h"
@@ -915,6 +916,9 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
std::vector<std::array<std::string, 9>> table_data;
table_data.reserve(devices.size());
// same data as the table below, for --log-jsonl consumers
common_json rows = common_json::array();
const std::string template_header = "%s: | %s | %s %s %s %s %s %s %s |\n";
const std::string template_gpu = "%s: | %s | %s = %s + (%s = %s + %s + %s) + %s |\n";
const std::string template_other = "%s: | %s | %s %s %s = %s + %s + %s %s |\n";
@@ -989,6 +993,19 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
std::to_string(mb.context / MiB),
std::to_string(mb.compute / MiB),
std::to_string(unaccounted / static_cast<int64_t>(MiB))});
rows.push_back({
{"kind", "device"},
{"name", name},
{"description", desc},
{"total", total / MiB},
{"free", free / MiB},
{"self", self / MiB},
{"model", mb.model / MiB},
{"context", mb.context / MiB},
{"compute", mb.compute / MiB},
{"unaccounted", unaccounted / static_cast<int64_t>(MiB)},
});
}
// print memory breakdown for host:
@@ -1004,6 +1021,15 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
std::to_string(mb_host.context / MiB),
std::to_string(mb_host.compute / MiB),
""}); // unaccounted
rows.push_back({
{"kind", "host"},
{"name", "Host"},
{"self", self / MiB},
{"model", mb_host.model / MiB},
{"context", mb_host.context / MiB},
{"compute", mb_host.compute / MiB},
});
}
// print memory breakdown for all remaining buffer types:
@@ -1025,6 +1051,16 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
std::to_string(mb.context / MiB),
std::to_string(mb.compute / MiB),
""}); // unaccounted
rows.push_back({
{"kind", "buffer_type"},
{"name", name},
{"self", self / MiB},
{"model", mb.model / MiB},
{"context", mb.context / MiB},
{"compute", mb.compute / MiB},
});
seen_buffer_types.insert(buft);
}
@@ -1042,6 +1078,11 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
__func__, td[1].c_str(), td[2].c_str(), td[3].c_str(), td[4].c_str(), td[5].c_str(),
td[6].c_str(), td[7].c_str(), td[8].c_str());
}
LOG_JSON("fit_memory_breakdown", common_json({
{"unit", "MiB"},
{"rows", rows},
}));
}
void common_fit_print(
common/log.cpp
+63 -13
View File
@@ -37,6 +37,16 @@ void common_log_set_verbosity_thold(int verbosity) {
common_log_verbosity_thold = verbosity;
}
static bool common_log_jsonl = false;
bool common_log_get_jsonl(void) {
return common_log_jsonl;
}
void common_log_set_jsonl(bool jsonl) {
common_log_jsonl = jsonl;
}
static int64_t t_us() {
return std::chrono::duration_cast<std::chrono::microseconds>(std::chrono::system_clock::now().time_since_epoch()).count();
}
@@ -87,6 +97,7 @@ struct common_log_entry {
bool is_end { false }; // signals the worker thread to stop
bool prefix { false };
bool jsonl { false };
bool is_json { false }; // msg already holds a serialized JSON object
common_log_entry(size_t size = 256) : msg(size) { }
@@ -107,6 +118,12 @@ struct common_log_entry {
}
if (jsonl) {
if (is_json) {
fprintf(fcur, "%s\n", msg.data());
fflush(fcur);
return;
}
common_json obj = {
{"type", "log"},
{"time", timestamp},
@@ -156,7 +173,6 @@ struct common_log {
file = nullptr;
prefix = false;
timestamps = false;
jsonl = false;
running = false;
t_start = t_us();
@@ -184,7 +200,6 @@ private:
bool prefix;
bool timestamps;
bool jsonl;
bool running;
int64_t t_start;
@@ -273,7 +288,8 @@ public:
entry.is_end = false;
entry.level = level;
entry.prefix = prefix;
entry.jsonl = jsonl;
entry.jsonl = common_log_jsonl;
entry.is_json = false;
entry.timestamp = 0;
if (timestamps) {
entry.timestamp = t_us() - t_start;
@@ -283,6 +299,42 @@ public:
cv_new.notify_one();
}
void add_json(const char * type, const common_json & obj) {
const common_json full = {
{"type", type},
{"data", obj},
};
const std::string text = full.dump_safe();
std::unique_lock<std::mutex> lock(mtx);
// block if the queue is full
cv_full.wait(lock, [this]() { return !running || !is_full(); });
if (!running) {
// discard messages while the worker thread is paused
return;
}
auto & entry = queue[tail];
if (entry.msg.size() < text.size() + 1) {
entry.msg.resize(text.size() + 1);
}
memcpy(entry.msg.data(), text.c_str(), text.size() + 1);
entry.is_end = false;
entry.level = GGML_LOG_LEVEL_NONE;
entry.prefix = false;
entry.jsonl = true;
entry.is_json = true;
entry.timestamp = 0;
tail = (tail + 1) % queue.size();
cv_new.notify_one();
}
void resume() {
std::lock_guard<std::mutex> lock(mtx);
@@ -388,12 +440,6 @@ public:
this->timestamps = timestamps;
}
void set_jsonl(bool jsonl) {
std::lock_guard<std::mutex> lock(mtx);
this->jsonl = jsonl;
}
};
//
@@ -440,6 +486,14 @@ void common_log_add(struct common_log * log, enum ggml_log_level level, const ch
va_end(args);
}
void common_log_add_json(struct common_log * log, const char * type, const common_json & obj) {
if (!common_log_jsonl) {
return;
}
log->add_json(type, obj);
}
void common_log_set_file(struct common_log * log, const char * file) {
log->set_file(file);
}
@@ -467,10 +521,6 @@ void common_log_set_timestamps(struct common_log * log, bool timestamps) {
log->set_timestamps(timestamps);
}
void common_log_set_jsonl(struct common_log * log, bool jsonl) {
log->set_jsonl(jsonl);
}
void common_log_flush(struct common_log * log) {
log->pause();
log->resume();
common/log.h
+18 -1
View File
@@ -43,6 +43,10 @@ int common_log_get_verbosity_thold(void);
void common_log_set_verbosity_thold(int verbosity); // not thread-safe
bool common_log_get_jsonl(void);
void common_log_set_jsonl(bool jsonl); // not thread-safe
int common_log_get_verbosity(enum ggml_log_level level);
void common_log_default_callback(enum ggml_log_level level, const char * text, void * user_data);
@@ -91,7 +95,6 @@ void common_log_set_file (struct common_log * log, const char * file); // n
void common_log_set_colors (struct common_log * log, log_colors colors); // not thread-safe
void common_log_set_prefix (struct common_log * log, bool prefix); // whether to output prefix to each log
void common_log_set_timestamps(struct common_log * log, bool timestamps); // whether to output timestamps in the prefix
void common_log_set_jsonl (struct common_log * log, bool jsonl); // print each log as a JSON object on one line, not thread-safe
void common_log_flush (struct common_log * log); // flush all pending log messages
// helper macros for logging
@@ -127,3 +130,17 @@ void common_log_flush (struct common_log * log); // f
#define LOG_WRNV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_WARN, verbosity, __VA_ARGS__)
#define LOG_ERRV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_ERROR, verbosity, __VA_ARGS__)
#define LOG_CNTV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_CONT, verbosity, __VA_ARGS__)
class common_json; // defined in common/json.h
// helper allows different types of json output
// no-op if --log-jsonl is not set
void common_log_add_json(struct common_log * log, const char * type, const common_json & data);
// will only print if --log-jsonl is set
#define LOG_JSON(type, data) \
do { \
if (common_log_get_jsonl()) { \
common_log_add_json(common_log_main(), type, data); \
} \
} while (0)
common/parsers/qwen3-coder.cpp
+28 -3
View File
@@ -104,9 +104,34 @@ common_chat_params common_chat_params_init_qwen3_coder(const common_chat_templat
auto arg_open = p.tool_arg_open("<parameter=" + p.tool_arg_name(p.literal(param.name)) + ">\n");
auto arg_value = param.schema->may_be_string() ?
arg_string :
p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
auto types = param.schema->value_types();
auto arg_value = p.eps();
if (!types.has(common_chat_schema::TYPE_STRING)) {
arg_value = p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
} else if (types.is_only(common_chat_schema::TYPE_STRING)) {
arg_value = arg_string;
} else {
// The string alternative accepts any text, so the grammar only keeps the raw string
// rule. The parser still tries the JSON alternatives first to type the value.
auto json_value = p.choice();
if (types.has(common_chat_schema::TYPE_OBJECT)) {
json_value |= p.json_object();
}
if (types.has(common_chat_schema::TYPE_ARRAY)) {
json_value |= p.json_array();
}
if (types.has(common_chat_schema::TYPE_NUMBER) || types.has(common_chat_schema::TYPE_INTEGER)) {
json_value |= p.json_number();
}
if (types.has(common_chat_schema::TYPE_BOOLEAN)) {
json_value |= p.json_bool();
}
if (types.has(common_chat_schema::TYPE_NULL)) {
json_value |= p.json_null();
}
arg_value = p.gbnf(p.atomic(p.tool_arg_json_value(json_value) + arg_close) | arg_string, "xml-arg-string");
}
auto arg_rule = p.rule(rule_name, p.tool_arg(arg_open + arg_value));
conversion/__init__.py
+1
View File
@@ -168,6 +168,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Mamba2ForCausalLM": "mamba",
"MambaForCausalLM": "mamba",
"MambaLMHeadModel": "mamba",
"MapleForCausalLM": "maple",
"MellumForCausalLM": "mellum",
"MiMoV2FlashForCausalLM": "mimo",
"MiMoV2ForCausalLM": "mimo",
conversion/maple.py
+87
View File
@@ -0,0 +1,87 @@
from __future__ import annotations
from typing import Iterable, TYPE_CHECKING, cast
import torch
if TYPE_CHECKING:
from torch import Tensor
from .base import LazyTorchTensor, ModelBase, TextModel, gguf
@ModelBase.register("MapleForCausalLM")
@ModelBase.example("deepgrove/maple-preview")
class MapleModel(TextModel):
model_arch = gguf.MODEL_ARCH.MAPLE
def set_gguf_parameters(self):
super().set_gguf_parameters()
hparams = self.hparams
assert hparams["hidden_act"] == "silu"
assert hparams.get("num_shared_experts", 0) == 0
assert hparams.get("norm_topk_prob", True)
assert hparams.get("nope_on_global_attention", False)
head_dim = hparams.get("head_dim", hparams["hidden_size"] // hparams["num_attention_heads"])
partial_rotary_factor = self.rope_parameters.get("partial_rotary_factor", 1.0)
self.gguf_writer.add_vocab_size(hparams["vocab_size"])
self.gguf_writer.add_rope_dimension_count(int(head_dim * partial_rotary_factor))
self.gguf_writer.add_sliding_window(hparams["sliding_window"])
self.gguf_writer.add_sliding_window_pattern([layer_type == "sliding_attention" for layer_type in hparams["layer_types"]])
self.gguf_writer.add_expert_feed_forward_length(hparams["moe_intermediate_size"])
# the reference clamps the MoE SwiGLU gate/up at 7.0 (modeling_maple.py)
self.gguf_writer.add_swiglu_clamp_exp([7.0] * self.block_count)
_experts: list[dict[str, Tensor]] | None = None
@staticmethod
def _stack_experts(tensors: list[Tensor]) -> Tensor:
shape = (len(tensors), *tensors[0].shape)
dtype = tensors[0].dtype
meta = LazyTorchTensor.meta_with_dtype_and_shape(dtype, shape)
# tensors goes through args, not the closure, so that `func` matches
# LazyBase's single-argument shape
def stack(ts: list[Tensor]) -> Tensor:
result = torch.empty(shape, dtype=dtype)
for expert_id, tensor in enumerate(ts):
result[expert_id].copy_(LazyTorchTensor.to_eager(tensor))
ts.clear()
return result
return cast(torch.Tensor, LazyTorchTensor(meta=meta, args=(tensors,), func=stack))
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
if "mlp.experts" in name:
n_experts = self.hparams["num_experts"]
assert bid is not None
if self._experts is None:
self._experts = [{} for _ in range(self.block_count)]
self._experts[bid][name] = data_torch
if len(self._experts[bid]) >= n_experts * 3:
for weight_name in ("down_proj", "gate_proj", "up_proj"):
tensors = []
for expert_id in range(n_experts):
expert_name = f"model.layers.{bid}.mlp.experts.{expert_id}.{weight_name}.weight"
tensors.append(self._experts[bid].pop(expert_name))
merged_name = f"model.layers.{bid}.mlp.experts.{weight_name}.weight"
yield from super().modify_tensors(self._stack_experts(tensors), merged_name, bid)
return
yield from super().modify_tensors(data_torch, name, bid)
def prepare_tensors(self):
super().prepare_tensors()
if self._experts is not None:
experts = [name for layer in self._experts for name in layer]
if experts:
raise ValueError(f"Unprocessed experts: {experts}")
docs/release.md
-55
View File
@@ -1,55 +0,0 @@
# Release process
llama.cpp uses [semantic versioning](https://semver.org) (`MAJOR.MINOR.PATCH`).
## Version bump guidelines
| Change type | Version component |
|---|---|
| Breaking change to the public C API (`include/llama.h`) | `MAJOR` |
| Backward-compatible features, model support, or API addition | `MINOR` |
| Bug fix with no API change | `PATCH` |
The version is set in the three variables at the top of the root `CMakeLists.txt`:
```cmake
set(LLAMA_VERSION_MAJOR 0)
set(LLAMA_VERSION_MINOR 1)
set(LLAMA_VERSION_PATCH 0)
```
_A version bump should be included in the PR that introduces the change, or in a
dedicated bump commit merged before the release is cut._
_TODO: add PR labels (`semver: patch`, `semver: minor`, `semver: major`) to help
identify which PRs require a version bump before cutting a release._
## Making a release
Releases are created by running the [make-release](.github/workflows/make-release.yml)
which is a manual workflow.
The workflow runs against the branch selected in the "Run workflow" dialog
(default `master`) and takes an optional `commit` SHA. When a commit is given,
the workflow validates that the commit belongs to the branch and is not older
than 3 days from the branch HEAD, then releases that commit instead of the
branch HEAD.
The workflow creates an annotated git tag (e.g. `v0.1.0`) and pushes it to the
remote. No GitHub Release object is created, the tag is the release artifact.
## Building a release
By default, `LLAMA_BUILD_IS_DEV=ON` which appends a `-dev` suffix to `LLAMA_VERSION`,
marking the build as a nightly/development build. Distributors building from a
release tag must pass `-DLLAMA_BUILD_IS_DEV=OFF` to produce a clean version string
(e.g. `0.1.0` instead of `0.1.0-dev`).
## How releases reach users
Currently releases are not published to github releases, only nightly/development
builds are available there. The way users can access releases are using the following
channels:
- **llama-install.sh** — downloads pre-built binaries built from the release tag.
- **Package managers** — consume the git tag directly.
- **Build from source** — users clone the repo and check out the tag.
ggml/src/ggml-cpu/arch/s390/quants.c
+1
View File
@@ -417,6 +417,7 @@ void ggml_vec_dot_mxfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo
sumf = vec_hsum_f32x4(v_acc);
*s = sumf;
#else
UNUSED(nb);
UNUSED(x);
UNUSED(y);
UNUSED(ib);
ggml/src/ggml-cpu/arch/s390/repack.cpp
+2
View File
@@ -70,6 +70,7 @@ void ggml_quantize_mat_q8_0_4x4(const float * GGML_RESTRICT x, void * GGML_RESTR
#endif
}
#if defined(__VXE__) || defined(__VXE2__)
static inline int16x8_t vxe_dot_acc(const int8x16_t v_x, const int8x16_t v_y, const int16x8_t v_acc) {
return vec_meadd(v_x, v_y, vec_moadd(v_x, v_y, v_acc));
}
@@ -84,6 +85,7 @@ static inline int32x4_t vxe_fold(const int16x8_t v_sumi) {
const int16x8_t v_ones = vec_splats((int16_t)1);
return vec_add(vec_mule(v_sumi, v_ones), vec_mulo(v_sumi, v_ones));
}
#endif
void ggml_gemv_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
const int qk = QK8_0;
ggml/src/ggml-cpu/ops.h
+4 -13
View File
@@ -5,10 +5,10 @@
//
// cache line
//
#if defined(__cpp_lib_hardware_interference_size)
#define CACHE_LINE_SIZE std::hardware_destructive_interference_size
#else
// TODO: rework CACHE_LINE_SIZE so std::hardware_destructive_interference_size
// can be used consistently between C and C++ TUs; the previous macro form
// diverged based on include order and undersized the work buffer.
// ref: https://github.com/ggml-org/llama.cpp/pull/28882
#if defined(__POWER9_VECTOR__)
#define CACHE_LINE_SIZE 128
#elif defined(__VXE__) || defined(__VXE2__)
@@ -16,17 +16,8 @@
#else
#define CACHE_LINE_SIZE 64
#endif
#endif
// -Winterference-size was introduced in GCC 12
#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Winterference-size"
#endif
static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
#pragma GCC diagnostic pop
#endif
// Work buffer size for im2col operations in CONV2D
#define GGML_IM2COL_WORK_SIZE (16 * 1024 * 1024)
ggml/src/ggml-cuda/common.cuh
+6
View File
@@ -334,6 +334,12 @@ static bool fp16_mma_hardware_available(const int cc) {
(GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
}
// To be used for feature selection of external libraries, e.g. cuBLAS.
static bool fast_bf16_hardware_available(const int cc) {
return (GGML_CUDA_CC_IS_AMD(cc) && (cc >= GGML_CUDA_CC_RDNA3 || GGML_CUDA_CC_IS_CDNA(cc)))
|| (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE);
}
static bool bf16_mma_hardware_available(const int cc) {
return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) ||
GGML_CUDA_CC_IS_CDNA(cc) || cc >= GGML_CUDA_CC_RDNA3 ||
ggml/src/ggml-cuda/ggml-cuda.cu
+10 -2
View File
@@ -1620,11 +1620,19 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
}
static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
const int cc = ggml_cuda_info().devices[ctx.device].cc;
ggml_type compute_type = src0->type;
if (ggml_is_quantized(compute_type)) {
compute_type = fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc)) {
compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
compute_type = GGML_TYPE_F32;
} else if (compute_type == GGML_TYPE_BF16 && !fast_bf16_hardware_available(cc)) {
if (GGML_CUDA_CC_IS_AMD(cc) && src1->ne[1] > 32) {
compute_type = GGML_TYPE_F32;
}
if (GGML_CUDA_CC_IS_NVIDIA(cc) && src1->ne[1] > (cc >= GGML_CUDA_CC_VOLTA ? 8 : 128)) {
compute_type = GGML_TYPE_F32;
}
}
if (dst->op_params[0] == GGML_PREC_F32) {
compute_type = GGML_TYPE_F32;
ggml/src/ggml-opencl/cl-program-cache.cpp
-453
View File
@@ -1,453 +0,0 @@
// Match the version setup ggml-opencl.cpp uses, so any cl.h declarations we
// touch are consistent across this backend's translation units.
#define CL_TARGET_OPENCL_VERSION GGML_OPENCL_TARGET_VERSION
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
#include "cl-program-cache.h"
#include "ggml-impl.h" // GGML_LOG_INFO / WARN
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <filesystem>
#include <fstream>
#include <system_error>
#include <vector>
#if defined(_WIN32)
# ifndef WIN32_LEAN_AND_MEAN
# define WIN32_LEAN_AND_MEAN
# endif
# ifndef NOMINMAX
# define NOMINMAX
# endif
# include <windows.h>
# include <process.h>
# define ggml_getpid() ((int) GetCurrentProcessId())
#else
# include <unistd.h>
# define ggml_getpid() ((int) getpid())
#endif
namespace fs = std::filesystem;
// ----------------------------------------------------------------------------
// SHA-256 (FIPS 180-4). Self-contained, ~80 lines, public-domain reference.
// Hot path is a few KB of source per kernel ⇒ <1 ms total per process init.
// ----------------------------------------------------------------------------
namespace {
struct sha256_ctx {
uint32_t state[8];
uint64_t bitlen;
uint8_t buf[64];
size_t buf_len;
};
const uint32_t K256[64] = {
0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174,
0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da,
0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967,
0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85,
0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070,
0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3,
0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2,
};
inline uint32_t rotr32(uint32_t x, unsigned n) { return (x >> n) | (x << (32 - n)); }
void sha256_compress(uint32_t state[8], const uint8_t block[64]) {
uint32_t w[64];
for (int i = 0; i < 16; ++i) {
w[i] = ((uint32_t)block[i*4 ] << 24) |
((uint32_t)block[i*4 + 1] << 16) |
((uint32_t)block[i*4 + 2] << 8) |
((uint32_t)block[i*4 + 3] );
}
for (int i = 16; i < 64; ++i) {
uint32_t s0 = rotr32(w[i-15], 7) ^ rotr32(w[i-15], 18) ^ (w[i-15] >> 3);
uint32_t s1 = rotr32(w[i-2], 17) ^ rotr32(w[i-2], 19) ^ (w[i-2] >> 10);
w[i] = w[i-16] + s0 + w[i-7] + s1;
}
uint32_t a = state[0],b = state[1],c = state[2],d = state[3],e = state[4],f = state[5],g = state[6],h = state[7];
for (int i = 0; i < 64; ++i) {
uint32_t S1 = rotr32(e, 6) ^ rotr32(e, 11) ^ rotr32(e, 25);
uint32_t ch = (e & f) ^ ((~e) & g);
uint32_t t1 = h + S1 + ch + K256[i] + w[i];
uint32_t S0 = rotr32(a, 2) ^ rotr32(a, 13) ^ rotr32(a, 22);
uint32_t maj = (a & b) ^ (a & c) ^ (b & c);
uint32_t t2 = S0 + maj;
h = g; g = f; f = e; e = d + t1;
d = c; c = b; b = a; a = t1 + t2;
}
state[0]+=a; state[1]+=b; state[2]+=c; state[3]+=d;
state[4]+=e; state[5]+=f; state[6]+=g; state[7]+=h;
}
void sha256_init(sha256_ctx & c) {
c.state[0]=0x6a09e667; c.state[1]=0xbb67ae85; c.state[2]=0x3c6ef372; c.state[3]=0xa54ff53a;
c.state[4]=0x510e527f; c.state[5]=0x9b05688c; c.state[6]=0x1f83d9ab; c.state[7]=0x5be0cd19;
c.bitlen = 0;
c.buf_len = 0;
}
void sha256_update(sha256_ctx & c, const void * data, size_t len) {
const uint8_t * p = (const uint8_t *) data;
c.bitlen += (uint64_t) len * 8;
if (c.buf_len > 0) {
size_t n = 64 - c.buf_len;
if (n > len) { n = len; }
memcpy(c.buf + c.buf_len, p, n);
c.buf_len += n;
p += n;
len -= n;
if (c.buf_len == 64) {
sha256_compress(c.state, c.buf);
c.buf_len = 0;
}
}
while (len >= 64) {
sha256_compress(c.state, p);
p += 64;
len -= 64;
}
if (len > 0) {
memcpy(c.buf, p, len);
c.buf_len = len;
}
}
void sha256_final(sha256_ctx & c, uint8_t out[32]) {
uint64_t bitlen = c.bitlen;
c.buf[c.buf_len++] = 0x80;
if (c.buf_len > 56) {
while (c.buf_len < 64) { c.buf[c.buf_len++] = 0; }
sha256_compress(c.state, c.buf);
c.buf_len = 0;
}
while (c.buf_len < 56) { c.buf[c.buf_len++] = 0; }
for (int i = 7; i >= 0; --i) { c.buf[c.buf_len++] = (uint8_t) (bitlen >> (i * 8)); }
sha256_compress(c.state, c.buf);
for (int i = 0; i < 8; ++i) {
out[i*4 ] = (uint8_t) (c.state[i] >> 24);
out[i*4 + 1] = (uint8_t) (c.state[i] >> 16);
out[i*4 + 2] = (uint8_t) (c.state[i] >> 8);
out[i*4 + 3] = (uint8_t) (c.state[i] );
}
}
std::string sha256_hex(const uint8_t digest[32]) {
static const char hex[] = "0123456789abcdef";
std::string s(64, '0');
for (int i = 0; i < 32; ++i) {
s[i*2 ] = hex[digest[i] >> 4];
s[i*2 + 1] = hex[digest[i] & 0xf];
}
return s;
}
std::string compute_key(const std::string & key_suffix,
const char * source,
const std::string & compile_opts) {
sha256_ctx c;
sha256_init(c);
static const uint8_t sep = 0;
sha256_update(c, source, strlen(source));
sha256_update(c, &sep, 1);
sha256_update(c, compile_opts.data(), compile_opts.size());
sha256_update(c, &sep, 1);
sha256_update(c, key_suffix.data(), key_suffix.size());
uint8_t digest[32];
sha256_final(c, digest);
return sha256_hex(digest);
}
bool make_dir_recursive(const std::string & path) {
if (path.empty()) { return false; }
// create_directories() already creates missing parents. It returns false
// (with ec clear) when the directory is already there, so re-check.
const fs::path p = fs::u8path(path);
std::error_code ec;
if (fs::create_directories(p, ec)) { return true; }
std::error_code ec_stat;
return fs::is_directory(p, ec_stat);
}
std::string default_cache_dir() {
#if defined(_WIN32)
const char * base = std::getenv("LOCALAPPDATA");
if (!base || !*base) { base = std::getenv("APPDATA"); }
if (!base || !*base) { base = std::getenv("TEMP"); }
if (!base || !*base) { base = "."; }
return std::string(base) + "\\llama.cpp\\cl-cache";
#elif defined(__APPLE__)
const char * home = std::getenv("HOME");
if (!home || !*home) { home = "."; }
return std::string(home) + "/Library/Caches/llama.cpp/cl-cache";
#else
// The throwing overload aborts the process when no usable temp directory
// exists (e.g. Android app contexts with TMPDIR unset); an empty return
// here just disables the cache instead.
std::error_code ec;
const fs::path tmp_path = fs::temp_directory_path(ec);
if (ec || tmp_path.empty()) { return {}; }
return tmp_path.string() + "/llama.cpp/cl-cache";
#endif
}
// Query a NUL-terminated string from clGetDeviceInfo / clGetPlatformInfo.
template <typename GetInfoFn, typename Object>
std::string query_string(GetInfoFn fn, Object obj, cl_uint name) {
size_t sz = 0;
if (fn(obj, name, 0, nullptr, &sz) != CL_SUCCESS || sz == 0) {
return {};
}
std::string s(sz, '\0');
if (fn(obj, name, sz, &s[0], nullptr) != CL_SUCCESS) {
return {};
}
if (!s.empty() && s.back() == '\0') {
s.pop_back();
}
return s;
}
std::string compute_key_suffix(cl_device_id device) {
cl_platform_id platform = nullptr;
clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(platform), &platform, nullptr);
std::string s;
s.reserve(512);
s += query_string(clGetDeviceInfo, device, CL_DEVICE_NAME); s.push_back('\0');
s += query_string(clGetDeviceInfo, device, CL_DRIVER_VERSION); s.push_back('\0');
s += query_string(clGetDeviceInfo, device, CL_DEVICE_VERSION); s.push_back('\0');
if (platform) {
s += query_string(clGetPlatformInfo, platform, CL_PLATFORM_VERSION); s.push_back('\0');
}
s += "fmt=" + std::to_string(CL_PROGRAM_CACHE_FORMAT_VERSION);
return s;
}
const uint8_t MAGIC[8] = { 'G','G','M','L','C','L','B','C' };
bool read_all(const std::string & path, std::vector<uint8_t> & out) {
std::ifstream f(fs::u8path(path), std::ios::binary);
if (!f) { return false; }
f.seekg(0, std::ios::end);
std::streamsize sz = f.tellg();
if (sz < 0) { return false; }
f.seekg(0, std::ios::beg);
out.resize((size_t) sz);
if (sz > 0) { f.read((char *) out.data(), sz); }
return f.good() || f.eof();
}
bool write_atomic(const std::string & path, const uint8_t * data, size_t len) {
const fs::path dst = fs::u8path(path);
const fs::path tmp = fs::u8path(path + ".tmp." + std::to_string(ggml_getpid()));
{
std::ofstream f(tmp, std::ios::binary | std::ios::trunc);
if (!f) { return false; }
f.write((const char *) data, (std::streamsize) len);
if (!f.good()) {
std::error_code ec_rm;
fs::remove(tmp, ec_rm);
return false;
}
}
std::error_code ec;
fs::rename(tmp, dst, ec);
if (ec) {
std::error_code ec_rm;
fs::remove(tmp, ec_rm);
return false;
}
return true;
}
} // namespace
static bool cache_debug_enabled() {
static int cached = -1;
if (cached < 0) {
const char * e = std::getenv("GGML_OPENCL_KERNEL_CACHE_DEBUG");
cached = (e && *e) ? 1 : 0;
}
return cached != 0;
}
static std::string opts_preview(const std::string & opts, size_t n = 120) {
if (opts.size() <= n) { return opts; }
return opts.substr(0, n) + "...";
}
// Running cache tally (diagnostic; plain ints — a benign race in the rare
// multi-threaded lazy-compile case at worst miscounts by one).
static int g_cache_hits = 0, g_cache_misses = 0, g_cache_saves = 0;
// Debug trace directly to stderr
static void cache_debug_line(const char * kind, const std::string & key,
const char * source, const std::string & opts) {
if (!cache_debug_enabled()) { return; }
fprintf(stderr, "ggml_opencl: cache %-4s [h=%d m=%d s=%d] key=%s src=%zuB opts='%s'\n",
kind, g_cache_hits, g_cache_misses, g_cache_saves,
key.substr(0, 16).c_str(), strlen(source), opts_preview(opts).c_str());
fflush(stderr);
}
cl_program_cache_state cl_program_cache_init(cl_device_id device) {
cl_program_cache_state st;
const char * env = std::getenv("GGML_OPENCL_KERNEL_CACHE_DIR");
if (env && (!std::strcmp(env, "0") || !std::strcmp(env, "off") ||
!std::strcmp(env, "none") || !std::strcmp(env, "disable") ||
!std::strcmp(env, "disabled"))) {
if (cache_debug_enabled()) {
fprintf(stderr, "ggml_opencl: kernel cache disabled by GGML_OPENCL_KERNEL_CACHE_DIR=%s\n", env);
fflush(stderr);
}
return st;
}
std::string dir;
if (!env || !*env || !std::strcmp(env, "1") || !std::strcmp(env, "default")) {
dir = default_cache_dir();
if (dir.empty()) {
GGML_LOG_INFO("ggml_opencl: kernel cache disabled (no usable default cache directory)\n");
return st;
}
} else {
dir = env;
}
if (!make_dir_recursive(dir)) {
GGML_LOG_INFO("ggml_opencl: kernel cache disabled (cannot create directory '%s')\n", dir.c_str());
return st;
}
st.dir = dir;
st.key_suffix = compute_key_suffix(device);
GGML_LOG_INFO("ggml_opencl: kernel cache enabled at '%s'\n", st.dir.c_str());
if (cache_debug_enabled()) {
fprintf(stderr, "ggml_opencl: kernel cache enabled at '%s' "
"(GGML_OPENCL_KERNEL_CACHE_DIR=off to disable)\n", st.dir.c_str());
fflush(stderr);
}
return st;
}
cl_program cl_program_cache_try_load(
const cl_program_cache_state & state,
cl_context context,
cl_device_id device,
const char * source,
const std::string & compile_opts) {
if (state.dir.empty() || !source) { return nullptr; }
const std::string key = compute_key(state.key_suffix, source, compile_opts);
const std::string path = state.dir + "/" + key + ".clbin";
std::vector<uint8_t> file;
if (!read_all(path, file)) {
++g_cache_misses;
cache_debug_line("MISS", key, source, compile_opts);
return nullptr;
}
if (file.size() < 16 || std::memcmp(file.data(), MAGIC, 8) != 0) { return nullptr; }
uint32_t fmt =
((uint32_t) file[ 8]) | ((uint32_t) file[ 9] << 8) |
((uint32_t) file[10] << 16) | ((uint32_t) file[11] << 24);
if (fmt != CL_PROGRAM_CACHE_FORMAT_VERSION) { return nullptr; }
const size_t hdr_len = 16;
const unsigned char * bin = file.data() + hdr_len;
const size_t bin_len = file.size() - hdr_len;
cl_int err = CL_SUCCESS;
cl_int bin_err = CL_SUCCESS;
cl_program p = clCreateProgramWithBinary(context, 1, &device, &bin_len, &bin, &bin_err, &err);
if (err != CL_SUCCESS || bin_err != CL_SUCCESS || p == nullptr) {
if (p) { clReleaseProgram(p); }
return nullptr;
}
err = clBuildProgram(p, 0, nullptr, compile_opts.c_str(), nullptr, nullptr);
if (err != CL_SUCCESS) {
clReleaseProgram(p);
return nullptr;
}
++g_cache_hits;
cache_debug_line("HIT", key, source, compile_opts);
return p;
}
void cl_program_cache_try_save(
const cl_program_cache_state & state,
cl_program program,
cl_device_id /*device*/,
const char * source,
const std::string & compile_opts) {
if (state.dir.empty() || !program || !source) {
return;
}
cl_uint n_dev = 0;
if (clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(n_dev), &n_dev, nullptr) != CL_SUCCESS || n_dev == 0) {
return;
}
std::vector<size_t> sizes(n_dev);
if (clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * n_dev, sizes.data(), nullptr) != CL_SUCCESS) {
return;
}
if (sizes.empty() || sizes[0] == 0) {
return;
}
std::vector<std::vector<uint8_t>> binaries(n_dev);
std::vector<unsigned char *> bin_ptrs(n_dev);
for (cl_uint i = 0; i < n_dev; ++i) {
binaries[i].resize(sizes[i]);
bin_ptrs[i] = binaries[i].data();
}
if (clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char *) * n_dev, bin_ptrs.data(), nullptr) != CL_SUCCESS) {
return;
}
// We only care about the first device's binary — that's the one we'd
// re-load with on a future cache hit. Multi-device contexts aren't a
// pattern this backend uses today.
const std::vector<uint8_t> & bin = binaries[0];
std::vector<uint8_t> file;
file.reserve(16 + bin.size());
file.insert(file.end(), MAGIC, MAGIC + 8);
uint32_t fmt = CL_PROGRAM_CACHE_FORMAT_VERSION;
file.push_back((uint8_t) (fmt & 0xff));
file.push_back((uint8_t) ((fmt >> 8) & 0xff));
file.push_back((uint8_t) ((fmt >> 16) & 0xff));
file.push_back((uint8_t) ((fmt >> 24) & 0xff));
file.push_back(0); file.push_back(0); file.push_back(0); file.push_back(0); // reserved
file.insert(file.end(), bin.begin(), bin.end());
const std::string key = compute_key(state.key_suffix, source, compile_opts);
const std::string path = state.dir + "/" + key + ".clbin";
if (!write_atomic(path, file.data(), file.size())) {
GGML_LOG_INFO("ggml_opencl: kernel cache: failed to write '%s'\n", path.c_str());
} else {
++g_cache_saves;
cache_debug_line("SAVE", key, source, compile_opts);
}
}
ggml/src/ggml-opencl/cl-program-cache.h
-75
View File
@@ -1,75 +0,0 @@
// On-disk cache for OpenCL cl_program binaries. Lets a fresh process skip the
// expensive clBuildProgram-from-source step when a binary for the exact same
// (source, compile options, device, driver, platform) was previously saved.
//
// Activation: default on via GGML_OPENCL_KERNEL_CACHE_DIR:
// unset / empty / "1" / "default" : platform default cache dir
// (%LOCALAPPDATA%\llama.cpp\cl-cache,
// ~/Library/Caches/llama.cpp/cl-cache,
// <temp dir>/llama.cpp/cl-cache elsewhere)
// "0" / "off" / "none" / "disable(d)" : disabled (all functions no-op)
// any other value : used verbatim as the cache path
// If the chosen directory cannot be created/used, the cache silently disables
// itself for the process and falls back to source compile.
// GGML_OPENCL_KERNEL_CACHE_DEBUG=1 prints a HIT/MISS/SAVE trace (with a running
// tally) straight to stderr — visible even in tools that filter INFO/WARN logs;
// redirect stderr to record it.
//
// Cache key (SHA-256 hex):
// sha256(source_bytes || '\x00' ||
// compile_opts || '\x00' ||
// CL_DEVICE_NAME || '\x00' ||
// CL_DRIVER_VERSION || '\x00' ||
// CL_PLATFORM_VERSION || '\x00' ||
// CL_PROGRAM_CACHE_FORMAT_VERSION)
//
// The key fully captures everything that can affect the produced binary,
// without needing the host source revision (a kernel source change shows up
// in source_bytes; a compile-option change shows up in compile_opts).
//
// File layout per cache entry: <cache_dir>/<sha256-hex>.clbin
// bytes [0..7] : magic "GGMLCLBC"
// bytes [8..11] : uint32_t format version (CL_PROGRAM_CACHE_FORMAT_VERSION)
// bytes [12..15] : uint32_t reserved (0)
// bytes [16..] : raw cl_program binary as returned by
// clGetProgramInfo(CL_PROGRAM_BINARIES)
//
// Concurrency: writes go to <name>.tmp.<pid> then atomic rename. On race,
// last-writer-wins. No locks.
#pragma once
#include <CL/cl.h>
#include <string>
// Bumped manually if host-side OpenCL API usage changes in a way that
// affects compile semantics but does not show up in source_bytes /
// compile_opts (e.g. switching from clCreateProgramWithSource to
// clCompileProgram + clLinkProgram, or changing how multiple sources
// are concatenated). Most commits — including kernel changes — do NOT
// require bumping this; the source bytes already capture those.
#define CL_PROGRAM_CACHE_FORMAT_VERSION 1u
struct cl_program_cache_state {
// Empty string means cache is disabled.
std::string dir;
// Concatenated device/driver/platform identity + cache format version,
// computed once at init and folded into every key.
std::string key_suffix;
};
cl_program_cache_state cl_program_cache_init(cl_device_id device);
cl_program cl_program_cache_try_load(
const cl_program_cache_state & state,
cl_context context,
cl_device_id device,
const char * source,
const std::string & compile_opts);
void cl_program_cache_try_save(
const cl_program_cache_state & state,
cl_program program,
cl_device_id device,
const char * source,
const std::string & compile_opts);
ggml/src/ggml-sycl/dsv4-hc.cpp
-280
View File
@@ -1,280 +0,0 @@
#include "ggml-impl.h"
#include "dsv4-hc.hpp"
#include <cmath>
static constexpr int DSV4_HC = 4;
static void dsv4_hc_pre_f32_sycl(
const float * x, const float * weights, float * dst,
int64_t n_embd, int64_t hc, int64_t n_tokens,
int64_t sx0, int64_t sx1, int64_t sx2,
int64_t sw0, int64_t sw1,
int64_t sd0, int64_t sd1,
queue_ptr stream) {
const int64_t nr = n_embd * n_tokens;
const int64_t block_size = 256;
const int64_t num_blocks = (nr + block_size - 1) / block_size;
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item) {
const int64_t ir = item.get_global_id(0);
if (ir >= nr) {
return;
}
const int64_t i0 = ir % n_embd;
const int64_t it = ir / n_embd;
float sum = x[i0*sx0 + it*sx2] * weights[it*sw1];
for (int64_t ih = 1; ih < hc; ++ih) {
const float xv = x[i0*sx0 + ih*sx1 + it*sx2];
const float wv = weights[ih*sw0 + it*sw1];
sum += xv * wv;
}
dst[i0*sd0 + it*sd1] = sum;
});
}
static void dsv4_hc_comb_norm_cols(float * comb, float eps) {
for (int idst = 0; idst < DSV4_HC; ++idst) {
float sum = eps;
for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
sum += comb[idst + DSV4_HC*isrc];
}
const float inv_sum = 1.0f / sum;
for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
comb[idst + DSV4_HC*isrc] *= inv_sum;
}
}
}
static void dsv4_hc_comb_norm_rows(float * comb, float eps) {
for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
float sum = eps;
for (int idst = 0; idst < DSV4_HC; ++idst) {
sum += comb[idst + DSV4_HC*isrc];
}
const float inv_sum = 1.0f / sum;
for (int idst = 0; idst < DSV4_HC; ++idst) {
comb[idst + DSV4_HC*isrc] *= inv_sum;
}
}
}
static void dsv4_hc_comb_f32_sycl(
const float * mixes,
const float * scale,
const float * base,
float * dst,
int64_t n_tokens,
int64_t sm0,
int64_t sm1,
int64_t ss0,
int64_t sb0,
int64_t sd0,
int64_t sd1,
int64_t sd2,
float eps,
int32_t n_iter,
queue_ptr stream) {
constexpr int comb_offset = 2*DSV4_HC;
const int64_t block_size = 256;
const int64_t num_blocks = (n_tokens + block_size - 1) / block_size;
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item_ct1) {
const int64_t it = item_ct1.get_global_id(0);
if (it >= n_tokens) {
return;
}
const float scale_comb = scale[2*ss0];
float comb[DSV4_HC*DSV4_HC];
for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
float max = -INFINITY;
for (int idst = 0; idst < DSV4_HC; ++idst) {
const int idx = idst + DSV4_HC*isrc;
const float v = mixes[(comb_offset + idx)*sm0 + it*sm1] * scale_comb + base[(comb_offset + idx)*sb0];
comb[idx] = v;
max = fmaxf(max, v);
}
float sum = 0.0f;
for (int idst = 0; idst < DSV4_HC; ++idst) {
const int idx = idst + DSV4_HC*isrc;
const float v = expf(comb[idx] - max);
comb[idx] = v;
sum += v;
}
const float inv_sum = 1.0f / sum;
for (int idst = 0; idst < DSV4_HC; ++idst) {
const int idx = idst + DSV4_HC*isrc;
comb[idx] = comb[idx] * inv_sum + eps;
}
}
dsv4_hc_comb_norm_cols(comb, eps);
for (int32_t i = 1; i < n_iter; ++i) {
dsv4_hc_comb_norm_rows(comb, eps);
dsv4_hc_comb_norm_cols(comb, eps);
}
for (int isrc = 0; isrc < DSV4_HC; ++isrc) {
for (int idst = 0; idst < DSV4_HC; ++idst) {
const int idx = idst + DSV4_HC*isrc;
dst[idst*sd0 + isrc*sd1 + it*sd2] = comb[idx];
}
}
});
}
static void dsv4_hc_post_f32_sycl(
const float * x, const float * residual, const float * post, const float * comb, float * dst,
int64_t n_embd, int64_t hc, int64_t n_tokens,
int64_t sx0, int64_t sx1,
int64_t sr0, int64_t sr1, int64_t sr2,
int64_t sp0, int64_t sp1,
int64_t sc0, int64_t sc1, int64_t sc2,
int64_t sd0, int64_t sd1, int64_t sd2,
queue_ptr stream) {
const int64_t nr = n_embd * hc * n_tokens;
const int64_t block_size = 256;
const int64_t num_blocks = (nr + block_size - 1) / block_size;
stream->parallel_for(
sycl::nd_range<1>(sycl::range<1>(num_blocks * block_size), sycl::range<1>(block_size)),
[=](sycl::nd_item<1> item) {
const int64_t ir = item.get_global_id(0);
if (ir >= nr) {
return;
}
const int64_t i0 = ir % n_embd;
const int64_t idst = (ir / n_embd) % hc;
const int64_t it = ir / (n_embd * hc);
float sum = x[i0*sx0 + it*sx1] * post[idst*sp0 + it*sp1];
for (int64_t isrc = 0; isrc < hc; ++isrc) {
sum += residual[i0*sr0 + isrc*sr1 + it*sr2] * comb[idst*sc0 + isrc*sc1 + it*sc2];
}
dst[i0*sd0 + idst*sd1 + it*sd2] = sum;
});
}
void ggml_sycl_op_dsv4_hc_pre(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
const ggml_tensor * x = dst->src[0];
const ggml_tensor * weights = dst->src[1];
GGML_ASSERT(x->type == GGML_TYPE_F32);
GGML_ASSERT(weights->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_TENSOR_LOCALS(size_t, nbx, x, nb);
GGML_TENSOR_LOCALS(size_t, nbw, weights, nb);
GGML_TENSOR_LOCALS(size_t, nbd, dst, nb);
const int64_t n_embd = x->ne[0];
const int64_t hc = x->ne[1];
const int64_t n_tokens = x->ne[2];
queue_ptr stream = ctx.stream();
dsv4_hc_pre_f32_sycl(
(const float *) x->data, (const float *) weights->data, (float *) dst->data,
n_embd, hc, n_tokens,
nbx0 / sizeof(float), nbx1 / sizeof(float), nbx2 / sizeof(float),
nbw0 / sizeof(float), nbw1 / sizeof(float),
nbd0 / sizeof(float), nbd1 / sizeof(float),
stream);
}
void ggml_sycl_op_dsv4_hc_comb(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/3);
const ggml_tensor * mixes = dst->src[0];
const ggml_tensor * scale = dst->src[1];
const ggml_tensor * base = dst->src[2];
GGML_ASSERT(mixes->type == GGML_TYPE_F32);
GGML_ASSERT(scale->type == GGML_TYPE_F32);
GGML_ASSERT(base->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
constexpr int64_t hc_mix_dim = (2 + DSV4_HC)*DSV4_HC;
GGML_ASSERT(mixes->ne[0] == hc_mix_dim);
GGML_ASSERT(dst->ne[0] == DSV4_HC);
GGML_ASSERT(dst->ne[1] == DSV4_HC);
GGML_ASSERT(dst->ne[2] == mixes->ne[1]);
GGML_ASSERT(scale->ne[0] >= 3);
GGML_ASSERT(base->ne[0] == hc_mix_dim);
GGML_TENSOR_LOCALS(size_t, nbm, mixes, nb);
GGML_TENSOR_LOCALS(size_t, nbs, scale, nb);
GGML_TENSOR_LOCALS(size_t, nbb, base, nb);
GGML_TENSOR_LOCALS(size_t, nbd, dst, nb);
const int64_t n_tokens = mixes->ne[1];
const float eps = ggml_get_op_params_f32(dst, 0);
const int32_t n_iter = ggml_get_op_params_i32(dst, 1);
queue_ptr stream = ctx.stream();
dsv4_hc_comb_f32_sycl(
(const float *) mixes->data, (const float *) scale->data, (const float *) base->data, (float *) dst->data,
n_tokens,
nbm0 / sizeof(float), nbm1 / sizeof(float),
nbs0 / sizeof(float),
nbb0 / sizeof(float),
nbd0 / sizeof(float), nbd1 / sizeof(float), nbd2 / sizeof(float),
eps, n_iter, stream);
}
void ggml_sycl_op_dsv4_hc_post(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/4);
const ggml_tensor * x = dst->src[0];
const ggml_tensor * residual = dst->src[1];
const ggml_tensor * post = dst->src[2];
const ggml_tensor * comb = dst->src[3];
GGML_ASSERT(x->type == GGML_TYPE_F32);
GGML_ASSERT(residual->type == GGML_TYPE_F32);
GGML_ASSERT(post->type == GGML_TYPE_F32);
GGML_ASSERT(comb->type == GGML_TYPE_F32);
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_TENSOR_LOCALS(size_t, nbx, x, nb);
GGML_TENSOR_LOCALS(size_t, nbr, residual, nb);
GGML_TENSOR_LOCALS(size_t, nbp, post, nb);
GGML_TENSOR_LOCALS(size_t, nbc, comb, nb);
GGML_TENSOR_LOCALS(size_t, nbd, dst, nb);
const int64_t n_embd = x->ne[0];
const int64_t n_tokens = x->ne[1];
const int64_t hc = residual->ne[1];
queue_ptr stream = ctx.stream();
dsv4_hc_post_f32_sycl(
(const float *) x->data, (const float *) residual->data,
(const float *) post->data, (const float *) comb->data, (float *) dst->data,
n_embd, hc, n_tokens,
nbx0 / sizeof(float), nbx1 / sizeof(float),
nbr0 / sizeof(float), nbr1 / sizeof(float), nbr2 / sizeof(float),
nbp0 / sizeof(float), nbp1 / sizeof(float),
nbc0 / sizeof(float), nbc1 / sizeof(float), nbc2 / sizeof(float),
nbd0 / sizeof(float), nbd1 / sizeof(float), nbd2 / sizeof(float),
stream);
}
ggml/src/ggml-sycl/dsv4-hc.hpp
-10
View File
@@ -1,10 +0,0 @@
#ifndef GGML_SYCL_DSV4_HC_HPP
#define GGML_SYCL_DSV4_HC_HPP
#include "common.hpp"
void ggml_sycl_op_dsv4_hc_pre(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_op_dsv4_hc_comb(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
void ggml_sycl_op_dsv4_hc_post(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
#endif // GGML_SYCL_DSV4_HC_HPP
ggml/src/ggml-sycl/lightning-indexer.cpp
-197
View File
@@ -1,197 +0,0 @@
#include "lightning-indexer.hpp"
#include "dequantize.hpp"
static void lightning_indexer_f32_sycl(
const char * q, const char * k, const char * w, const char * m, float * dst,
int64_t n_embd, int64_t n_head, int64_t n_batch, int64_t n_stream, int64_t n_kv,
int64_t nem3,
int64_t nbq1, int64_t nbq2, int64_t nbq3,
int64_t nbk2, int64_t nbk3,
int64_t nbw1, int64_t nbw3,
int64_t nbm1, int64_t nbm3,
int64_t nb1, int64_t nb3,
ggml_type k_type,
queue_ptr stream) {
constexpr int64_t LANES = WARP_SIZE;
constexpr int64_t ELEMS_PER_LANE = 8;
constexpr int64_t ROWS_PER_BLOCK = 4;
constexpr int64_t BLOCK_SIZE = ROWS_PER_BLOCK * LANES;
const int64_t n_rows = n_batch * n_stream * n_kv;
const int64_t n_blocks = (n_rows + ROWS_PER_BLOCK - 1) / ROWS_PER_BLOCK;
stream->parallel_for(
sycl::nd_range<1>(
sycl::range<1>(n_blocks * BLOCK_SIZE),
sycl::range<1>(BLOCK_SIZE)),
[=](sycl::nd_item<1> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
const int64_t ir = item.get_global_id(0);
const int64_t lane = ir % LANES;
const int64_t row = ir / LANES;
if (row >= n_rows) {
return;
}
const int64_t i_bs = row / n_kv;
const int64_t i_kv = row % n_kv;
const int64_t i_batch = i_bs / n_stream;
const int64_t i_stream = i_bs % n_stream;
// load K row slice into registers (row is contiguous, nbk0 == type size)
const char * k_base = k + i_kv*nbk2 + i_stream*nbk3;
float k_local[ELEMS_PER_LANE];
if (k_type == GGML_TYPE_F16) {
const sycl::half * k_row = (const sycl::half *) k_base;
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
k_local[j] = static_cast<float>(k_row[lane*ELEMS_PER_LANE + j]);
}
} else if (k_type == GGML_TYPE_F32) {
const float * k_row = (const float *) k_base;
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
k_local[j] = k_row[lane*ELEMS_PER_LANE + j];
}
} else {
const int64_t lane_base = lane * ELEMS_PER_LANE;
switch (k_type) {
case GGML_TYPE_BF16: {
const sycl::ext::oneapi::bfloat16 * k_row = (const sycl::ext::oneapi::bfloat16 *) k_base;
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
k_local[j] = static_cast<float>(k_row[lane_base + j]);
}
} break;
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
case GGML_TYPE_Q5_1: {
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
const int64_t idx = lane_base + j;
const int64_t ib = idx / QK4_0;
const int iqs = idx % (QK4_0/2);
dfloat2 kv;
if (k_type == GGML_TYPE_Q4_0) {
dequantize_q4_0(k_base, ib, iqs, kv);
} else if (k_type == GGML_TYPE_Q4_1) {
dequantize_q4_1(k_base, ib, iqs, kv);
} else if (k_type == GGML_TYPE_Q5_0) {
dequantize_q5_0(k_base, ib, iqs, kv);
} else {
dequantize_q5_1(k_base, ib, iqs, kv);
}
k_local[j] = (idx % QK4_0) < (QK4_0/2) ? static_cast<float>(kv.x()) : static_cast<float>(kv.y());
}
} break;
case GGML_TYPE_Q8_0: {
#pragma unroll
for (int64_t pair = 0; pair < ELEMS_PER_LANE / 2; ++pair) {
const int64_t elem0 = lane_base + 2 * pair;
dfloat2 kv;
dequantize_q8_0(k_base, elem0 / QK8_0, elem0 % QK8_0, kv);
k_local[2 * pair + 0] = static_cast<float>(kv.x());
k_local[2 * pair + 1] = static_cast<float>(kv.y());
}
} break;
case GGML_TYPE_IQ4_NL: {
#pragma unroll
for (int64_t pair = 0; pair < ELEMS_PER_LANE / 2; ++pair) {
const int64_t elem0 = lane_base + 2 * pair;
dfloat2 kv;
dequantize_iq4_nl(k_base, elem0 / QK4_NL, elem0 % QK4_NL, kv);
k_local[2 * pair + 0] = static_cast<float>(kv.x());
k_local[2 * pair + 1] = static_cast<float>(kv.y());
}
} break;
default:
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
k_local[j] = 0.0f;
}
break;
}
}
const char * q_base = q + i_batch*nbq2 + i_stream*nbq3;
const float * w_base = (const float *) (w + i_batch*nbw1 + i_stream*nbw3);
float score = 0.0f;
for (int64_t h = 0; h < n_head; ++h) {
const float * q_row = (const float *) (q_base + h*nbq1);
float dot = 0.0f;
#pragma unroll
for (int64_t j = 0; j < ELEMS_PER_LANE; ++j) {
const int64_t i = lane*ELEMS_PER_LANE + j;
if (i < n_embd) {
dot += q_row[i] * k_local[j];
}
}
dot = sycl::reduce_over_group(item.get_sub_group(), dot, sycl::plus<float>());
if (lane == 0) {
score += sycl::max(dot, 0.0f) * w_base[h];
}
}
if (lane == 0) {
const sycl::half * m_base = (const sycl::half *) (m + i_batch*nbm1 + (i_stream % nem3)*nbm3);
// flat-index store: storing through a strided base pointer
// hangs/misroutes writes on this stack when n_batch*n_stream > 1
const int64_t dst_idx = i_kv + i_batch*(nb1/sizeof(float)) + i_stream*(nb3/sizeof(float));
dst[dst_idx] = score + static_cast<float>(m_base[i_kv]);
}
});
}
void ggml_sycl_op_lightning_indexer(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/4);
const ggml_tensor * q = dst->src[0];
const ggml_tensor * k = dst->src[1];
const ggml_tensor * w = dst->src[2]; // weights
const ggml_tensor * m = dst->src[3]; // mask
GGML_ASSERT(dst->type == GGML_TYPE_F32);
GGML_ASSERT( q->type == GGML_TYPE_F32);
GGML_ASSERT( w->type == GGML_TYPE_F32);
GGML_ASSERT( m->type == GGML_TYPE_F16);
GGML_ASSERT(k->type == GGML_TYPE_F16 || k->type == GGML_TYPE_F32 || k->type == GGML_TYPE_BF16 ||
k->type == GGML_TYPE_Q8_0 || k->type == GGML_TYPE_Q5_1 || k->type == GGML_TYPE_Q5_0 ||
k->type == GGML_TYPE_Q4_1 || k->type == GGML_TYPE_Q4_0 || k->type == GGML_TYPE_IQ4_NL);
GGML_TENSOR_LOCALS(int64_t, neq, q, ne);
GGML_TENSOR_LOCALS(size_t, nbq, q, nb);
GGML_TENSOR_LOCALS(int64_t, nek, k, ne);
GGML_TENSOR_LOCALS(size_t, nbk, k, nb);
GGML_TENSOR_LOCALS(size_t, nbw, w, nb);
GGML_TENSOR_LOCALS(int64_t, nem, m, ne);
GGML_TENSOR_LOCALS(size_t, nbm, m, nb);
GGML_TENSOR_LOCALS(int64_t, ne, dst, ne);
GGML_TENSOR_LOCALS(size_t, nb, dst, nb);
// input rows must be contiguous
GGML_ASSERT(nbq0 == ggml_type_size(q->type));
GGML_ASSERT(nbk0 == ggml_type_size(k->type));
GGML_ASSERT(nbm0 == ggml_type_size(m->type));
GGML_ASSERT(nb0 == ggml_type_size(dst->type));
const int64_t n_embd = neq0;
const int64_t n_head = neq1;
const int64_t n_batch = neq2;
const int64_t n_stream = neq3;
const int64_t n_kv = nek2;
GGML_ASSERT(n_embd == WARP_SIZE * 8);
lightning_indexer_f32_sycl(
(const char *) q->data, (const char *) k->data,
(const char *) w->data, (const char *) m->data, (float *) dst->data,
n_embd, n_head, n_batch, n_stream, n_kv, nem3,
nbq1, nbq2, nbq3,
nbk2, nbk3,
nbw1, nbw3,
nbm1, nbm3,
nb1, nb3,
k->type,
ctx.stream());
}
ggml/src/ggml-sycl/lightning-indexer.hpp
-8
View File
@@ -1,8 +0,0 @@
#ifndef GGML_SYCL_LIGHTNING_INDEXER_HPP
#define GGML_SYCL_LIGHTNING_INDEXER_HPP
#include "common.hpp"
void ggml_sycl_op_lightning_indexer(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
#endif // GGML_SYCL_LIGHTNING_INDEXER_HPP
ggml/src/ggml-vulkan/ggml-vulkan.cpp
+17 -1
View File
@@ -339,6 +339,7 @@ static void ggml_vk_print_device_lost_info(const vk_device& device);
struct vk_queue_handle {
vk::Queue queue;
vk_device_ref device;
std::mutex * device_submit_mutex = nullptr;
virtual void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) = 0;
virtual void lock() {} // no-op by default (internally synchronized case)
virtual void unlock() {}
@@ -348,6 +349,11 @@ struct vk_queue_handle {
struct vk_queue_handle_synchronized : vk_queue_handle {
std::mutex mutex;
void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
// Workaround for NVIDIA driver bug
std::unique_lock<std::mutex> device_guard;
if (device_submit_mutex) {
device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
}
std::lock_guard<std::mutex> guard(mutex);
try {
queue.submit(submits, fence);
@@ -362,9 +368,14 @@ struct vk_queue_handle_synchronized : vk_queue_handle {
void unlock() override { mutex.unlock(); }
};
// Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
struct vk_queue_handle_unsynchronized : vk_queue_handle {
void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
// Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
// Workaround for NVIDIA driver bug
std::unique_lock<std::mutex> device_guard;
if (device_submit_mutex) {
device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
}
try {
queue.submit(submits, fence);
} catch (vk::DeviceLostError &) {
@@ -841,6 +852,7 @@ static bool ggml_vk_lightning_indexer_k_type_supported(ggml_type type) {
struct vk_device_struct {
std::recursive_mutex mutex;
std::mutex queue_submit_mutex;
mutable std::shared_mutex pinned_memory_mutex;
// Guards compile_pending, all_pipelines, and the dynamic pipeline maps
@@ -3526,6 +3538,10 @@ static std::unique_ptr<vk_queue> ggml_vk_create_queue(vk_device& device, uint32_
h->queue = device->device.getQueue2(queue_info2);
h->device = device;
// Avoid concurrent submissions on NVIDIA due to driver bug.
if (device->vendor_id == VK_VENDOR_ID_NVIDIA) {
h->device_submit_mutex = &device->queue_submit_mutex;
}
q->handle = h;
q->cmd_pool.init(device, q.get());
gguf-py/gguf/constants.py
+19
View File
@@ -541,6 +541,7 @@ class MODEL_ARCH(IntEnum):
ARWKV7 = auto()
MAMBA = auto()
MAMBA2 = auto()
MAPLE = auto()
JAMBA = auto()
XVERSE = auto()
COMMAND_R = auto()
@@ -1295,6 +1296,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
MODEL_ARCH.ARWKV7: "arwkv7",
MODEL_ARCH.MAMBA: "mamba",
MODEL_ARCH.MAMBA2: "mamba2",
MODEL_ARCH.MAPLE: "maple",
MODEL_ARCH.JAMBA: "jamba",
MODEL_ARCH.XVERSE: "xverse",
MODEL_ARCH.COMMAND_R: "command-r",
@@ -3487,6 +3489,23 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.SSM_NORM,
MODEL_TENSOR.SSM_OUT,
],
MODEL_ARCH.MAPLE: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT_NORM,
MODEL_TENSOR.OUTPUT,
MODEL_TENSOR.ATTN_NORM,
MODEL_TENSOR.ATTN_Q,
MODEL_TENSOR.ATTN_Q_NORM,
MODEL_TENSOR.ATTN_K,
MODEL_TENSOR.ATTN_K_NORM,
MODEL_TENSOR.ATTN_V,
MODEL_TENSOR.ATTN_OUT,
MODEL_TENSOR.FFN_NORM,
MODEL_TENSOR.FFN_GATE_INP,
MODEL_TENSOR.FFN_GATE_EXP,
MODEL_TENSOR.FFN_DOWN_EXP,
MODEL_TENSOR.FFN_UP_EXP,
],
MODEL_ARCH.JAMBA: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT_NORM,
models/templates/muse-glimmer.jinja
-211
View File
@@ -1,211 +0,0 @@
{#
Template: Muse Glimmer ATEM Chat Template
Renders the ATEM tool-calling protocol: reasoning channel (to=self), tool
channels (to=<tool>), and the user channel, plus tool definitions and the
valid-recipient list in the system block.
Whitespace note: every tag uses the {%- -%} / {{- -}} stripping markers, so
the indentation below is purely for readability and contributes nothing to
the rendered output.
#}
{%- macro render_content(content) -%}
{%- if content is string -%}
{{- content -}}
{%- elif content is not none -%}
{%- for part in content -%}
{%- if part['type'] == 'image' -%}
{{- '<|patch|>' -}}
{%- elif part['type'] == 'video' -%}
{{- '<|video|>' -}}
{%- elif part['type'] == 'text' -%}
{{- part['text'] -}}
{%- endif -%}
{%- endfor -%}
{%- endif -%}
{%- endmacro -%}
{%- macro render_atem(tc) -%}
{%- set args = tc.function.arguments -%}
{%- if args is not mapping -%}
{{- raise_exception('Muse Glimmer ATEM chat template requires tool_call.function.arguments to be a dict (mapping); a JSON string cannot be parsed in the HF jinja sandbox.') -}}
{%- endif -%}
{{- '<atem:function_calls>\n<atem:invoke name="' + tc.function.name + '">\n' -}}
{%- for k, v in args.items() -%}
{{- '<atem:parameter name="' + k + '">' -}}
{%- if v is boolean -%}
{%- if v -%}
true
{%- else -%}
false
{%- endif -%}
{%- elif v is none -%}
null
{%- elif v is mapping or (v is iterable and v is not string) -%}
{{- v | tojson -}}
{%- else -%}
{{- v -}}
{%- endif -%}
{{- '</atem:parameter>\n' -}}
{%- endfor -%}
{{- '</atem:invoke>\n</atem:function_calls>' -}}
{%- endmacro -%}
{%- macro render_tool_defs(tools) -%}
{{- 'In this environment you have access to a set of tools you can use to answer the user\'s question.\n\n' -}}
{{- 'You can invoke a function by writing a "<atem:function_calls>" block like the following:\n' -}}
{{- '<atem:function_calls>\n<atem:invoke name="$FUNCTION_NAME">\n<atem:parameter name="$PARAMETER_NAME">$PARAMETER_VALUE</atem:parameter>\n...\n</atem:invoke>\n</atem:function_calls>\n\n' -}}
{{- 'String and scalar parameters should be specified as is, while lists and objects should use JSON format. Note that spaces for string values are not stripped. The output is not expected to be valid XML and is parsed with regular expressions.\n' -}}
{{- 'Here are the functions available in JSONSchema format:\n' -}}
{{- '// Tool metadata\n' -}}
{%- set nsns = namespace(seen=[]) -%}
{%- for tool in tools -%}
{%- set fn = tool.function if tool.function is defined else tool -%}
{%- set tns = fn.name.split('.')[0] -%}
{%- if tns not in nsns.seen -%}
{%- set nsns.seen = nsns.seen + [tns] -%}
{%- endif -%}
{%- endfor -%}
{%- set nd = tool_namespace_descriptions if tool_namespace_descriptions is defined else {} -%}
{%- for tns in nsns.seen -%}
{{- '{"name": ' + (tns | tojson) + ', "description": ' + ((nd[tns] if tns in nd else '') | tojson) + '}\n' -}}
{%- endfor -%}
{{- '// Function schemas' -}}
{%- for tool in tools -%}
{%- set fn = tool.function if tool.function is defined else tool -%}
{{- '\n{"name": ' + (fn.name | tojson) + ', "description": ' + (fn.description | tojson) + ', "parameters": ' + (fn.parameters | tojson) + '}' -}}
{%- endfor -%}
{{- '\n\nHere\'s an example of how to call a function in the tool set:\n' -}}
{{- '(If the tool namespace is not specified, invoke the function directly as `example_function_name` rather than `example_tool_name.example_function_name`)\n\n' -}}
{{- 'to=example_tool_name.example_function_name\n\n' -}}
{{- '<atem:function_calls>\n<atem:invoke name="example_tool_name.example_function_name">\n' -}}
{{- '<atem:parameter name="example_parameter_1">value_1</atem:parameter>\n' -}}
{{- '<atem:parameter name="example_parameter_2">This is the value for the second parameter\nthat can span\n"multiple" lines\n</atem:parameter>\n' -}}
{{- '</atem:invoke>\n</atem:function_calls>' -}}
{%- endmacro -%}
{%- macro render_reasoning() -%}
{%- set rs = reasoning_strength if reasoning_strength is defined and reasoning_strength else 'high' -%}
{{- 'Reasoning strength: ' + rs + '.' -}}
{%- endmacro -%}
{%- macro render_system_meta(tools) -%}
{%- set rns = namespace(recipients=['"self"'], nslist=[]) -%}
{%- if tools -%}
{%- for tool in tools -%}
{%- set fn = tool.function if tool.function is defined else tool -%}
{%- set tns = fn.name.split('.')[0] -%}
{%- if tns not in rns.nslist -%}
{%- set rns.nslist = rns.nslist + [tns] -%}
{%- endif -%}
{%- endfor -%}
{%- for tns in rns.nslist -%}
{%- set rns.recipients = rns.recipients + ['"' + tns + '.*"'] -%}
{%- endfor -%}
{%- endif -%}
{%- set rns.recipients = rns.recipients + ['"user"'] -%}
{{- '# Valid recipients: ' + rns.recipients | join(', ') + '.' -}}
{%- endmacro -%}
{{- bos_token -}}
{%- set ns = namespace(has_system=false) -%}
{%- for m in messages -%}
{%- if m['role'] == 'system' -%}
{%- set ns.has_system = true -%}
{%- endif -%}
{%- endfor -%}
{%- if not ns.has_system -%}
{{- '<|start|>system<|message|>You are a helpful AI assistant.' -}}
{%- set kc = knowledge_cutoff if knowledge_cutoff is defined and knowledge_cutoff else '2026-01-04' -%}
{{- '\nKnowledge cutoff: ' + kc + '.' -}}
{%- if current_date is defined and current_date -%}
{{- '\nCurrent date: ' + current_date + '.' -}}
{%- elif strftime_now is defined -%}
{{- '\nCurrent date: ' + strftime_now('%Y-%m-%d') + '.' -}}
{%- endif -%}
{{- '\n\n' -}}
{{- render_reasoning() -}}
{%- if tools -%}
{{- '\n\n' -}}
{{- render_tool_defs(tools) -}}
{%- endif -%}
{{- '\n\n' -}}
{{- render_system_meta(tools) -}}
{{- '<|eot|>' -}}
{%- endif -%}
{%- for message in messages -%}
{%- set role = message['role'] -%}
{%- set end_token = '<|eom|>' if (not loop.last and messages[loop.index0 + 1]['role'] == role) else '<|eot|>' -%}
{%- if role == 'system' -%}
{#- Callers sometimes write the directive into the system prompt themselves.
Normalise "Reasoning effort" to "Reasoning strength" (jinja has no
case-insensitive replace, hence the four realistic casings), then skip
the kwarg-driven line below if the prompt already carries one. -#}
{%- set sys_text = render_content(message['content'])
| replace('Reasoning effort', 'Reasoning strength')
| replace('Reasoning Effort', 'Reasoning Strength')
| replace('reasoning effort', 'reasoning strength')
| replace('REASONING EFFORT', 'REASONING STRENGTH') -%}
{{- '<|start|>system<|message|>' -}}
{{- sys_text -}}
{%- if 'reasoning strength' not in (sys_text | lower) -%}
{{- '\n\n' -}}
{{- render_reasoning() -}}
{%- endif -%}
{%- if tools -%}
{{- '\n\n' -}}
{{- render_tool_defs(tools) -}}
{%- endif -%}
{{- '\n\n' -}}
{{- render_system_meta(tools) -}}
{{- '<|eot|>' -}}
{%- elif role == 'user' -%}
{{- '<|start|>user<|message|>' -}}
{{- render_content(message['content']) -}}
{{- '<|eot|>' -}}
{%- elif role == 'tool' -%}
{%- set tname = message.get('name') -%}
{%- if not tname -%}
{%- set tcid = message.get('tool_call_id') -%}
{%- set rns = namespace(name=tcid if tcid else '') -%}
{%- for m in messages -%}
{%- if m.get('tool_calls') -%}
{%- for tc in m['tool_calls'] -%}
{%- if tcid is not none and tc.id is defined and tc.id == tcid -%}
{%- set rns.name = tc.function.name -%}
{%- endif -%}
{%- endfor -%}
{%- endif -%}
{%- endfor -%}
{%- set tname = rns.name -%}
{%- endif -%}
{{- '<|start|>tool ' + tname + '<|message|><tool_output name="' + tname + '">\n' -}}
{{- render_content(message['content']) -}}
{{- '\n</tool_output><|eot|>' -}}
{%- elif role == 'assistant' -%}
{%- if message.get('reasoning_content') -%}
{{- '<|start|>assistant to=self<|message|>' + message['reasoning_content'] + '<|eom|>' -}}
{%- endif -%}
{%- if message.get('tool_calls') -%}
{%- for tc in message['tool_calls'] -%}
{{- '<|start|>assistant to=' + tc.function.name + '<|message|>' -}}
{{- render_atem(tc) -}}
{%- if loop.last -%}
{{- end_token -}}
{%- else -%}
{{- '<|eom|>' -}}
{%- endif -%}
{%- endfor -%}
{%- else -%}
{%- set recipient = message.get('recipient') or 'user' -%}
{%- set end_turn = message.get('end_turn') -%}
{%- if end_turn is none -%}
{%- set end_turn = not (recipient and recipient != 'user') -%}
{%- endif -%}
{{- '<|start|>assistant' -}}
{%- if recipient -%}
{{- ' to=' + recipient -}}
{%- endif -%}
{{- '<|message|>' -}}
{{- render_content(message['content']) -}}
{{- ('<|eot|>' if end_turn else '<|eom|>') -}}
{%- endif -%}
{%- endif -%}
{%- endfor -%}
{%- if add_generation_prompt -%}
{{- '<|start|>assistant' -}}
{%- endif -%}
skills/code-review/SKILL.md
-149
View File
@@ -1,149 +0,0 @@
---
name: code-review
description: Review llama.cpp changes against project conventions and common reviewer pitfalls before a PR. Use when the user wants to review a diff, branch, or PR.
---
# Review llama.cpp changes
This skill reviews changes against llama.cpp's conventions and the pitfalls that reviewers flag most often, so the contributor can fix them before a maintainer has to. It has two modes:
- **Self-review (default):** review the contributor's own local changes (uncommitted work, or a branch vs `master`) as a pre-PR pass. Ask which if it's ambiguous; default to `git diff master...HEAD` plus any uncommitted changes.
- **Read-only review of a PR/file:** if the user points at a PR number or specific files (including code they didn't write), review those and report findings.
In both modes the output is **private review notes for the user to read and act on** - it is never something to post. This is a hard rule from `AGENTS.md`: an agent must NEVER write, or help write, a PR comment, a review comment, or a reply to a reviewer, by any means including `gh`. Do not offer to. If the user asks you to post the notes, refuse and point them at that rule. Present findings in the conversation only.
Before starting, read `AGENTS.md` and `CONTRIBUTING.md` if not already in context - the "Coding guidelines", "Naming guidelines", and AI usage sections are the baseline this review enforces. For a diff that adds a new model architecture, also read `docs/development/HOWTO-add-model.md` and consider the dedicated `add-new-model` skill.
## Step 0 - Scope the diff and pick the checklists
Identify what actually changed and which area checklists below apply. Run `git diff --stat` (or `gh pr view <n> --json files` for PR mode) and bucket the touched paths:
- `conversion/`, `gguf-py/`, `src/models/`, `src/llama-arch.*` -> **New model / architecture**
- `ggml/` (any backend, op, or `ggml.h`) -> **ggml / backend**
- `include/llama.h` and other public headers -> **Public API**
- `tools/server/` -> **Server**
- anything else, plus all of the above -> **General** (always runs)
Always run the **Scope and quick-reject gate**, the **Security review**, and the **General** checklist. Run each area checklist whose paths were touched. Additionally, if the diff introduces a new component, subsystem, or piece of infrastructure (a new file/class/module, a new abstraction, or hand-rolled machinery), run the **Approach and design** review. Tell the user which checklists you're running and why.
## Scope and quick-reject gate (always)
These are the patterns that get PRs closed without a full review. Check them first - a finding here is more important than any code nit, because it can mean the change shouldn't be a PR in its current form at all.
- Is there a prior issue/discussion for this? Features are supposed to start as an issue, not a PR (`CONTRIBUTING.md`). If this is a nontrivial feature with no linked issue, flag it and suggest opening one first.
- Is it a duplicate of existing/in-flight work? Suggest `gh search prs` / `gh search issues` for the feature. Many closed PRs were duplicates of something already queued.
- Is it self-contained and single-purpose? Multiple unrelated changes/optimizations bundled together get sent back to be split. Flag unrelated changes and suggest separate PRs.
- Does it touch multiple ggml backends at once? Initial support should be CPU-only, other backends as follow-ups (`CONTRIBUTING.md`). Flag CUDA/Metal/Vulkan/etc. changes bundled into a feature's first PR.
- Does it add a new `ggml_type` / quantization type? That carries a disproportionate maintenance burden and needs the full justification package (GGUF sample upload, perplexity vs FP16/BF16 and similar sizes, KL-divergence data, CPU perf numbers). Absent that, it will be rejected regardless of code quality.
- Is it invasive - new subsystem, core-API reshaping, changes to shared graph/sampler code that other models don't need? Flag it and suggest a discussion with maintainers before investing further.
- Is it niche/vendor-specific in a way that adds a maintenance burden nobody will own long-term? Flag the maintenance-ownership question.
- Is the change semantically correct, or a plausible-looking "fix" that misunderstands the code? Sanity-check the actual behavior, not just that it compiles.
- AI-disclosure: if AI meaningfully contributed, is the PR template's disclosure section filled in? Remind the user. Never suggest writing the PR description or commit message for them.
## Security review (mandatory)
Mandatory on every review; any finding here is **blocking**. Rule of thumb: GGUF metadata, tensor shapes, tokenizer/grammar input, and all server/RPC fields are attacker-controlled - bound them before use.
- **Sizes/counts from tensor dims:** validate before allocating. Products like `ne[i]*nb[i]`/nbytes can overflow on crafted dims into an undersized alloc then heap overflow. Overflow checks must run BEFORE the arithmetic they guard - padding/alignment macros wrap to 0 near `SIZE_MAX`, so a guard after the pad passes.
- **GGUF strings/arrays:** cap declared lengths and element counts before using them to size a loop or buffer; validate element type and length before casting an array to a pointer or reading fixed indices (`[i+1]`, `[0..2]`).
- **Element-type confusion:** casting `gguf_get_arr_data()` or `tensor->data` to `float *`/`int32_t *` needs an element-type check first (`gguf_get_kv_type() == GGUF_TYPE_ARRAY` then `gguf_get_arr_type()`; `type == GGML_TYPE_F32` for tensors). A `UINT8` array or `I8` tensor passes every length check, then gets read 4 bytes per element - a nearby length check is not a type check.
- **Loaders:** `GGML_ASSERT` on a file-derived value aborts the process; throw instead where the caller already catches (vocab, model loader, clip).
- **File-supplied counts indexing fixed arrays:** bound any count (e.g. layer/block count into a `LLAMA_MAX_*` array) before indexing; watch checks that only fire when an optional key is present.
- **Declared vs actual array length:** check the declared length of a GGUF array against the count actually read, not just against a buffer size.
- **Bounds comparisons:** flag narrowing casts (`size_t`->`int32_t`) and signed/unsigned mixing that can bypass a length check and copy past a buffer.
- **Parsed/derived indices:** range-check `stoi`/`atoi` results and catch parse throws; never use a default or derived token id (EOS/BOS/...) as an index without a bounds check.
- **Reused/reserved buffers:** recheck bounds after a buffer is shrunk or reused; watch `reserve()` then index-by-assumed-size, and header fields read before their length is checked.
- **Server JSON ints:** clamp client-supplied integers (token/discard counts, offsets) to non-negative and an upper bound before they reach index/pointer arithmetic.
- **RPC-deserialized fields:** treat every field (type/buffer/data/ne/nb/op_params) as hostile - validate before use. Null/zero buffers skipping validation, attacker data pointers, out-of-range type indices, and negative strides sign-extending past a corner-only assert all give arbitrary read/write.
- **Lifetime/UAF:** flag stored raw pointers to caller/temporary storage, cached pointers to buffers a later free releases, async ops whose source may drop before completion, and structures not invalidated on free/realloc. Null-check conditionally-built or "not required" tensors before dereferencing.
## Approach and design (when a new component/infra is introduced)
Run this whenever the diff adds a new component, subsystem, or piece of infrastructure. Reviews too often stop at "does it work" - a diff can be correct and still be the wrong approach, and a messy design costs more long-term than a bug. Evaluate the *approach*, not just the behavior; raising a cleaner one is a high-value finding, not a nit. If you see a better design, describe it concretely rather than just calling the current one bad.
- **Simpler approach upstream:** the biggest win is often a different data model or design that removes whole subsystems, not tweaks to the code as written. Complexity must be justified by the problem, not by the first thing that worked.
- **Reuse over reinvention:** grep for an existing helper, library, object, or mechanism before adding a new one. Reimplementing what the codebase already has reintroduces solved bugs and adds maintenance surface.
- **Clear ownership/lifetime:** prefer RAII and obvious ownership over manual liveness flags, hand-tracked pointers, and "is it still alive?" checks - manual lifetime tracking is a recurring source of subtle bugs.
- **Right-sized machinery:** flag redundant, overkill, or heavier-than-needed primitives and abstractions; use the minimum the design actually needs.
- **Right structure and fit:** a new type should earn its place (split it if it serves two roles); follow existing patterns, idioms, and naming, and avoid constructs the project shuns.
- **Root cause vs symptom:** fixes layered on fixes signal a design to correct, not guard around.
## New model / architecture
See the `add-new-model` skill and `docs/development/HOWTO-add-model.md` for the full workflow; this is the review-time subset that reviewers most often catch:
- Don't branch on `model.arch` when the real dependency is a config/capability value - gate on the hparam/capability, not the architecture enum.
- If the model is a close variant of an existing arch, is the delta justified? Prefer reusing or subclassing the existing arch/model class over duplicating it. A near-duplicate class or `src/models/<name>.cpp` will be asked to merge with its sibling.
- New tensor names go through `tensor_mapping.py`, not ad-hoc name matching.
- For QKV, split the *activation* with `ggml_view`, not the *weight* tensor; rely on ggml broadcasting instead of manually duplicating tensors.
- New graph inputs are declared at the top of the graph-build function, not inline where first used.
- Hparams that the model can't run correctly without must be mandatory (hard-error if missing), not read with a silent default fallback. Only genuinely-optional-across-configs values get a fallback accessor.
- New/optional weight tensors (scales, etc.) must route through `build_lora_mm` and the existing helpers, matching convention - don't leave raw matmuls copied from another arch.
- Don't hack RoPE with a custom sin/cos implementation. If `ggml_rope_ext` genuinely can't express it, that's an issue for discussion, not a PR.
- Test the quantized-KV path (`-ctk`/`-ctv q8_0`), not just default f16 - new speculative/attention features silently break there.
- Preserve existing explanatory comments about model-specific quirks when copying code; note the provenance ("copied from X, with Y added").
- Remove dead code/branches left over from adapting a reference implementation.
## ggml / backend
- `supports_op` (and any dispatch/gating condition) must be scoped exactly to the cases being changed - a condition meant for a few quant types must not silently disable or enable everything else.
- No hardcoded warp/lane size - use `ggml_cuda_get_physical_warp_size()` (32 on CUDA, 64 on HIP/ROCm) and the portable helpers.
- Strip leftover debug/profiling/logging code before review.
- New or changed op? Update `docs/ops.md` and the relevant `docs/ops/*.csv` for the touched backend.
- New op or operator change needs corresponding `test-backend-ops` cases, and (per `CONTRIBUTING.md`) consistency across at least two backends.
- New kernels are expected to come with concrete perf data (throughput across realistic tensor shapes), not just correctness.
- Don't have a backend mutate the cgraph as a shortcut - that's an unresolved architectural question, not something to slip in.
- Expect this to need two maintainer approvals; that's normal for `ggml/` changes, not a sign something is wrong.
- For CUDA: Avoid excessively templating kernels, only add this where it shows visible performance gain.
## Public API (`include/llama.h`)
Public API changes carry a higher bar than internal ones (`CONTRIBUTING.md`). Review for:
- Justification: why doesn't an existing mechanism (e.g. `cb_eval`, existing batch/sampler knobs) suffice? If it does, the change likely shouldn't add public surface. This is the single most common reason these PRs are rejected.
- Experimental or stop-gap surface belongs in a side header (`llama-ext.h`), not in `llama.h`.
- Keep it minimal and general: prefer one general call over several narrow convenience wrappers; make new calls forward-compatible (e.g. mixed-modality batches) rather than assuming today's shape.
- The C API is the first-class, stable, ABI-defining surface - don't propose a parallel C++ API as a replacement. `llama-cpp.h` stays a thin convenience layer.
- Types and naming: sized integer types (`int32_t`, `size_t` for sizes/offsets); `snake_case`; `<class>_<method>` = `<class>_<action>_<noun>`; enum values upper-case and prefixed with the enum name; `_t` suffix for opaque types. Avoid gratuitous signature/ABI changes to existing exported functions.
- Every new API needs a working example/tool exercising it in the same PR - reviewers find real bugs by requiring it to be wired into `server`, `embedding`, `perplexity`, etc.
## Server (`tools/server/`)
- Is the feature within server's defined scope? Check `tools/server/README-dev.md` - out-of-scope features get declined.
- Security: don't trust client-supplied headers (e.g. `X-Forwarded-For`) or add footguns; things like IP allowlisting belong at a reverse proxy unless there's a trusted-proxy design.
- Wire new behavior into the existing request/response and checkpoint paths correctly; watch for resource leaks across requests.
## Multimodal (`tools/mtmd/`)
- Tensor names must be prefixed by `v.`, `a.`, `mm.` or `a.mm.` (legacy naming doesn't follow this convention - this is expected, but new code should follow it).
- Do not use explicit sin/cos for RoPE; use `ggml_rope_ext` instead, see `HOWTO-add-model.md`. If it can't express the needed behavior, that's a design discussion, not a PR.
- New GGML ops must not be introduced in the same PR, you must push it as a separate PR.
- In most cases, `build_vit` should be enough to build the transformer graph for vision models. Do not add a loop to build the transformer graph manually, unless you have a very good reason to do so. If you do, please explain why in the PR description.
- If you need a dedicated preprocessor, there is a high chance that it can be a derived class from one of the existing preprocessors. Check carefully before adding a new preprocessor class.
- If the model need a new public API in `mtmd.h`, open a discussion first.
- For audio generation models, see `tools/mtmd/README-dev.md`
## General (always)
Enforce the `AGENTS.md` / `CONTRIBUTING.md` coding and naming guidelines on every changed line - this is a distinct pass from checking that the code works, and matters just as much for review speed:
- ASCII only in code and comments - no emdash, unicode arrows, `x`, `...` used as unicode; use `-`, `->`, `x`, `...` ASCII equivalents.
- Comments are concise and explain non-obvious *why*, not *what*. Flag verbose comments, comments that restate the code, comments that reference the current task/PR, and comments hard-wrapped to a fixed column width.
- Do not force-wrap prose/comments to a fixed character count or split a sentence across lines.
- `snake_case` names; `kebab-case` (lowercase-with-dashes) file names for C/C++, `.h` headers; Python files lowercase-with-underscores. Naming optimizes for longest common prefix (`number_small`, not `small_number`).
- 4-space indentation, brackets on the same line, `void * ptr`, `int & a`, no trailing whitespace; match the surrounding style.
- Reuse existing infrastructure over introducing new components; no new third-party dependencies, extra headers, or files unless clearly justified.
- Keep it simple: a simpler change doing 90% is often preferable to a complex one doing 100%. Flag unnecessary templates/fancy STL; basic `for` loops are fine here.
- Every added line should be something the contributor can explain and defend to a reviewer without AI help - flag anything that looks copied-in without understanding.
- `Co-authored-by:` must be reserved for human co-authors; AI contributions (claude, cursor, codex, etc.) must use `Assisted-by:`; if this point is violated, it's a blocking finding.
- Any mentions of Minja must be treated as blocking; see `AGENTS.md` for why.
## Reporting
Group findings by severity so the user knows what actually blocks a merge:
1. **Blocking** - quick-reject/scope issues and correctness bugs; these can sink the PR regardless of everything else.
2. **Will slow the review** - convention/naming/comment violations, missing tests/docs/perf data, missing API justification or example.
3. **Nits** - minor style, optional cleanups.
For each finding, point to the file and line and say concretely what to change and why. Do not rewrite the whole diff unprompted; let the contributor make the fixes so they own and understand them. And do not draft any PR text, commit message, or reviewer reply - that is the contributor's to write.
src/llama-arch.cpp
+1
View File
@@ -62,6 +62,7 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
{ LLM_ARCH_STARCODER2, "starcoder2" },
{ LLM_ARCH_MAMBA, "mamba" },
{ LLM_ARCH_MAMBA2, "mamba2" },
{ LLM_ARCH_MAPLE, "maple" },
{ LLM_ARCH_JAMBA, "jamba" },
{ LLM_ARCH_FALCON_H1, "falcon-h1" },
{ LLM_ARCH_XVERSE, "xverse" },
src/llama-arch.h
+1
View File
@@ -67,6 +67,7 @@ enum llm_arch {
LLM_ARCH_STARCODER2,
LLM_ARCH_MAMBA,
LLM_ARCH_MAMBA2,
LLM_ARCH_MAPLE,
LLM_ARCH_JAMBA,
LLM_ARCH_FALCON_H1,
LLM_ARCH_XVERSE,
src/llama-grammar.cpp
+6 -5
View File
@@ -896,17 +896,18 @@ static void llama_grammar_advance_stack(
std::set<llama_grammar_stack, decltype(stack_cmp)> seen(stack_cmp);
while (!todo.empty()) {
llama_grammar_stack curr_stack = std::move(todo.back());
llama_grammar_stack curr_stack_candidate = std::move(todo.back());
todo.pop_back();
if (seen.find( curr_stack) != seen.end()) {
auto [curr_stack_it, inserted] = seen.insert(std::move(curr_stack_candidate));
if (!inserted) {
continue;
}
seen.insert(curr_stack);
const llama_grammar_stack & curr_stack = *curr_stack_it;
if (curr_stack.empty()) {
if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
new_stacks.emplace_back(std::move(curr_stack));
new_stacks.emplace_back(curr_stack);
}
continue;
}
@@ -949,7 +950,7 @@ static void llama_grammar_advance_stack(
case LLAMA_GRETYPE_TOKEN_NOT:
if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
// only add the stack if it's not a duplicate of one we already have
new_stacks.emplace_back(std::move(curr_stack));
new_stacks.emplace_back(curr_stack);
}
break;
default:
src/llama-graph.cpp
+1 -1
View File
@@ -2226,7 +2226,7 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
const float limit = hparams.swiglu_clamp_exp[il];
constexpr float eps = 1e-6f;
if (limit > eps) {
if (arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
if (arch == LLM_ARCH_MAPLE || arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
cur = ggml_swiglu_clamp(ctx0, cur, up, limit);
} else {
up = ggml_clamp(ctx0, up, -limit, limit);
src/llama-model-saver.cpp
+1
View File
@@ -33,6 +33,7 @@ bool llama_model_saver_supports_arch(llm_arch arch) {
case LLM_ARCH_LAGUNA:
case LLM_ARCH_GRANITE_SWA:
case LLM_ARCH_DOTS3NOTE: // TODO: need to handle SWA pattern and MLA+SWA config
case LLM_ARCH_MAPLE:
return false;
default:
return true;
src/llama-model.cpp
+4
View File
@@ -125,6 +125,7 @@
#include "models/mamba-base.cpp"
#include "models/mamba.cpp"
#include "models/mamba2.cpp"
#include "models/maple.cpp"
#include "models/mellum.cpp"
#include "models/mimo2.cpp"
#include "models/minicpm.cpp"
@@ -316,6 +317,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
return new llama_model_mamba(params);
case LLM_ARCH_MAMBA2:
return new llama_model_mamba2(params);
case LLM_ARCH_MAPLE:
return new llama_model_maple(params);
case LLM_ARCH_JAMBA:
return new llama_model_jamba(params);
case LLM_ARCH_XVERSE:
@@ -3173,6 +3176,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
case LLM_ARCH_SPARK2_5:
case LLM_ARCH_TALKIE:
case LLM_ARCH_MELLUM:
case LLM_ARCH_MAPLE:
return LLAMA_ROPE_TYPE_NEOX;
case LLM_ARCH_DFLASH:
src/models/gemma4-assistant.cpp
+1 -1
View File
@@ -4,7 +4,7 @@ void llama_model_gemma4_assistant::load_arch_hparams(llama_model_loader & ml) {
hparams.n_embd_inp_impl = hparams.n_embd_out();
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
uint32_t n_kv_shared_layers = 0;
ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
src/models/gemma4.cpp
+1 -1
View File
@@ -2,7 +2,7 @@
void llama_model_gemma4::load_arch_hparams(llama_model_loader & ml) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
uint32_t n_kv_shared_layers = 0;
ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
src/models/maple.cpp
+150
View File
@@ -0,0 +1,150 @@
#include "models.h"
void llama_model_maple::load_arch_hparams(llama_model_loader & ml) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
hparams.rope_freq_base_train_swa = hparams.rope_freq_base_train;
hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train;
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer_all);
switch (hparams.n_layer()) {
case 24: type = LLM_TYPE_20B; break;
default: type = LLM_TYPE_UNKNOWN;
}
}
void llama_model_maple::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t head_dim = hparams.n_embd_head_k();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0);
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, 0);
if (n_expert == 0) {
throw std::runtime_error("n_expert must be > 0 for Maple");
}
if (n_expert_used == 0) {
throw std::runtime_error("n_expert_used must be > 0 for Maple");
}
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
create_tensor_qkv(layer, i, n_embd, n_head * head_dim, n_head_kv * head_dim, n_head_kv * head_dim, 0);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head * head_dim, n_embd}, 0);
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {head_dim}, 0);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {head_dim}, 0);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0);
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
}
}
std::unique_ptr<llm_graph_context> llama_model_maple::build_arch_graph(const llm_graph_params & params) const {
return std::make_unique<graph>(*this, params);
}
llama_model_maple::graph::graph(const llama_model & model, const llm_graph_params & params) :
llm_graph_context(params) {
const int64_t n_embd_head = hparams.n_embd_head_k();
GGML_ASSERT(n_embd_head == hparams.n_embd_head_v());
ggml_tensor * inpL = build_inp_embd(model.tok_embd);
ggml_tensor * inp_pos = build_inp_pos();
auto * inp_attn = build_attn_inp_kv_iswa();
ggml_tensor * inp_out_ids = build_inp_out_ids();
for (int il = 0; il < n_layer; ++il) {
ggml_tensor * inpSA = inpL;
ggml_tensor * cur = build_norm(inpL, model.layers[il].attn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "attn_norm", il);
{
auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur, n_embd_head, n_head, n_head_kv, il);
Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, nullptr, LLM_NORM_RMS, il);
Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, nullptr, LLM_NORM_RMS, il);
cb(Qcur, "Qcur_normed", il);
cb(Kcur, "Kcur_normed", il);
if (hparams.is_swa(il)) {
const int64_t n_rot_l = hparams.n_rot(il);
const float freq_base_l = model.get_rope_freq_base(cparams, il);
const float freq_scale_l = model.get_rope_freq_scale(cparams, il);
Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
}
cb(Qcur, "Qcur", il);
cb(Kcur, "Kcur", il);
cb(Vcur, "Vcur", il);
cur = build_attn(inp_attn,
model.layers[il].wo, nullptr, model.layers[il].wo_s,
Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f / sqrtf(float(n_embd_head)), il);
cb(cur, "attn_out", il);
}
if (il == n_layer - 1 && inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
}
ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
cb(ffn_inp, "ffn_inp", il);
cur = build_norm(ffn_inp, model.layers[il].ffn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "ffn_norm", il);
cur = build_moe_ffn(cur,
model.layers[il].ffn_gate_inp,
model.layers[il].ffn_up_exps,
model.layers[il].ffn_gate_exps,
model.layers[il].ffn_down_exps,
nullptr,
n_expert, n_expert_used,
LLM_FFN_SILU, true,
1.0f,
LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
il);
cb(cur, "ffn_moe_out", il);
cur = ggml_add(ctx0, cur, ffn_inp);
cur = build_cvec(cur, il);
cb(cur, "l_out", il);
inpL = cur;
}
ggml_tensor * cur = build_norm(inpL, model.output_norm, nullptr, LLM_NORM_RMS, -1);
cb(cur, "result_norm", -1);
res->t_embd = cur;
cur = build_lora_mm(model.output, cur, model.output_s);
cb(cur, "result_output", -1);
res->t_logits = cur;
ggml_build_forward_expand(gf, cur);
}
src/models/mimo2.cpp
+1 -1
View File
@@ -9,7 +9,7 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
float value_scale = 0.0f;
if (ml.get_key(LLM_KV_ATTENTION_VALUE_SCALE, value_scale, false) && value_scale != 1.0f) {
src/models/models.h
+13
View File
@@ -945,6 +945,19 @@ struct llama_model_mamba2 : public llama_model_base {
};
struct llama_model_maple : public llama_model_base {
llama_model_maple(const struct llama_model_params & params) : llama_model_base(params) {}
void load_arch_hparams(llama_model_loader & ml) override;
void load_arch_tensors(llama_model_loader & ml) override;
struct graph : public llm_graph_context {
graph(const llama_model & model, const llm_graph_params & params);
};
std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
};
struct llama_model_jamba : public llama_model_base {
llama_model_jamba(const struct llama_model_params & params) : llama_model_base(params) {}
void load_arch_hparams(llama_model_loader & ml) override;
src/models/nemotron-h.cpp
+8 -2
View File
@@ -145,8 +145,14 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
const int64_t n_head_i = hparams.n_head(i);
const int64_t n_embd_k_gqa_i = hparams.n_embd_k_gqa(i);
const int64_t n_embd_v_gqa_i = hparams.n_embd_v_gqa(i);
const int64_t n_ff_exp = hparams.n_ff_exp(i) ? (int64_t)hparams.n_ff_exp(i) : n_ff / (int64_t)hparams.n_expert_used(i);
const int64_t n_ff_shexp = hparams.n_ff_shexp;
const int64_t n_expert_used_i = hparams.n_expert_used(i);
const int64_t n_ff_exp_i = hparams.n_ff_exp(i);
if (n_ff_exp_i == 0 && n_expert_used_i == 0) {
throw std::runtime_error(format("%s: layer %d declares neither expert_feed_forward_length nor expert_used_count, "
"cannot determine the expert FFN size", __func__, i));
}
const int64_t n_ff_exp = n_ff_exp_i ? n_ff_exp_i : n_ff / n_expert_used_i;
const int64_t n_ff_shexp = hparams.n_ff_shexp;
// NextN input-fusion tensors
layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), {n_embd}, mtp_flags);
src/models/qwen4exp.cpp
+11 -14
View File
@@ -157,7 +157,8 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);
// there is no output_norm: the final hyper-connection mixer carries it
hc_head_norm = create_tensor(tn(LLM_TENSOR_HC_HEAD_NORM, "weight"), { hc_dim }, 0);
// the gammas load as [n_embd, hc] so the grouped norm multiplies them without a graph reshape
hc_head_norm = create_tensor(tn(LLM_TENSOR_HC_HEAD_NORM, "weight"), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
hc_head_down = create_tensor(tn(LLM_TENSOR_HC_HEAD_DOWN, "weight"), { hc_dim, hc_lr }, 0);
hc_head_up = create_tensor(tn(LLM_TENSOR_HC_HEAD_UP, "weight"), { hc_lr, hc_dim }, 0);
@@ -203,11 +204,11 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
const int64_t conv_dim = key_dim * 2 + value_dim;
// two HC modules per layer: before the token mixer, before the MoE
layer.hc_attn_norm = create_tensor(tn(LLM_TENSOR_HC_ATTN_NORM, "weight", il), { hc_dim }, 0);
layer.hc_attn_norm = create_tensor(tn(LLM_TENSOR_HC_ATTN_NORM, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
layer.hc_attn_down = create_tensor(tn(LLM_TENSOR_HC_ATTN_DOWN, "weight", il), { hc_dim, hc_lr }, 0);
layer.hc_attn_up = create_tensor(tn(LLM_TENSOR_HC_ATTN_UP, "weight", il), { hc_lr, hc_dim }, 0);
layer.hc_attn_inject = create_tensor(tn(LLM_TENSOR_HC_ATTN_INJECT, "weight", il), { hc_dim, hc }, 0);
layer.hc_ffn_norm = create_tensor(tn(LLM_TENSOR_HC_FFN_NORM, "weight", il), { hc_dim }, 0);
layer.hc_ffn_norm = create_tensor(tn(LLM_TENSOR_HC_FFN_NORM, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
layer.hc_ffn_down = create_tensor(tn(LLM_TENSOR_HC_FFN_DOWN, "weight", il), { hc_dim, hc_lr }, 0);
layer.hc_ffn_up = create_tensor(tn(LLM_TENSOR_HC_FFN_UP, "weight", il), { hc_lr, hc_dim }, 0);
layer.hc_ffn_inject = create_tensor(tn(LLM_TENSOR_HC_FFN_INJECT, "weight", il), { hc_dim, hc }, 0);
@@ -240,9 +241,9 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
if (hparams.is_ple(il)) {
layer.ple_key = create_tensor(tn(LLM_TENSOR_PLE_KEY, "weight", il), { n_embd, hc_dim }, 0);
layer.ple_value = create_tensor(tn(LLM_TENSOR_PLE_VALUE, "weight", il), { n_embd, n_embd }, 0);
layer.ple_norm_key = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY, "weight", il), { hc_dim }, 0);
layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { hc_dim }, 0);
layer.ple_norm_conv = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV, "weight", il), { hc_dim }, 0);
layer.ple_norm_key = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
layer.ple_norm_conv = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
layer.ple_conv1d = create_tensor(tn(LLM_TENSOR_PLE_CONV1D, "weight", il), { hparams.ple_conv_kernel, hc_dim }, 0);
}
@@ -275,11 +276,10 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
const int64_t hc_dim = hc * n_embd;
const int64_t nt = x->ne[2];
// grouped RMSNorm: reduce over one stream, then scale all streams with the [hc_dim] gamma
// grouped RMSNorm: reduce over one stream, then scale all streams with the [n_embd, hc] gamma
// the converter folded each gamma to (1 + w)
ggml_tensor * xn = ggml_rms_norm(ctx0, x, hparams.f_norm_rms_eps);
ggml_tensor * xn = ggml_mul(ctx0, ggml_rms_norm(ctx0, x, hparams.f_norm_rms_eps), w_norm);
xn = ggml_reshape_2d(ctx0, xn, hc_dim, nt);
xn = ggml_mul(ctx0, xn, w_norm);
cb(xn, "hc_norm", il);
ggml_tensor * lo = build_lora_mm(w_down, xn);
@@ -1200,13 +1200,10 @@ ggml_tensor * llama_model_qwen4exp::graph::build_ple(
ggml_tensor * key = build_lora_mm(model.layers[il].ple_key, emb);
ggml_tensor * value = build_lora_mm(model.layers[il].ple_value, emb);
// both norms group over one hc stream, with a weight over the whole hc*n_embd layout
// both norms group over one hc stream, with a [n_embd, hc] weight
auto grouped_norm = [&](ggml_tensor * x, ggml_tensor * w) {
ggml_tensor * t = ggml_reshape_3d(ctx0, x, n_embd, hc, n_tokens);
t = ggml_rms_norm(ctx0, t, hparams.f_norm_rms_eps);
t = ggml_reshape_2d(ctx0, t, hc_dim, n_tokens);
t = ggml_mul(ctx0, t, w);
return ggml_reshape_3d(ctx0, t, n_embd, hc, n_tokens);
return ggml_mul(ctx0, ggml_rms_norm(ctx0, t, hparams.f_norm_rms_eps), w);
};
key = grouped_norm(key, model.layers[il].ple_norm_key);
src/models/step35.cpp
+1 -1
View File
@@ -23,7 +23,7 @@ void llama_model_step35::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer_all, false);
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer_all, false);