From a52dc60ba3ae0ef1e941ce9a4585672cc335a175 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sat, 27 Dec 2025 09:59:19 +0100 Subject: [PATCH 01/12] llama_fit_params: return enum for fail vs. error (#18374) --- include/llama.h | 8 +++++++- src/llama.cpp | 31 +++++++++++++++++++------------ tools/fit-params/fit-params.cpp | 4 ++-- 3 files changed, 28 insertions(+), 15 deletions(-) diff --git a/include/llama.h b/include/llama.h index f86293009..2d4e9a94e 100644 --- a/include/llama.h +++ b/include/llama.h @@ -467,10 +467,16 @@ extern "C" { // Frees all allocated memory LLAMA_API void llama_free(struct llama_context * ctx); + enum llama_params_fit_status { + LLAMA_PARAMS_FIT_STATUS_SUCCESS = 0, // found allocations that are projected to fit + LLAMA_PARAMS_FIT_STATUS_FAILURE = 1, // could not find allocations that are projected to fit + LLAMA_PARAMS_FIT_STATUS_ERROR = 2, // a hard error occured, e.g. because no model could be found at the specified path + }; + // fits mparams and cparams to free device memory (assumes system memory is unlimited) // returns true if the parameters could be successfully modified to fit device memory // this function is NOT thread safe because it modifies the global llama logger state - LLAMA_API bool llama_params_fit( + LLAMA_API enum llama_params_fit_status llama_params_fit( const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams, diff --git a/src/llama.cpp b/src/llama.cpp index 3428b1bd3..c53f2472b 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -140,6 +140,10 @@ enum layer_fraction_t { }; // this enum is only used in llama_params_fit_impl but needs to be defined outside of it to fix a Windows compilation issue +class llama_params_fit_exception : public std::runtime_error { + using std::runtime_error::runtime_error; +}; + static void llama_params_fit_impl( const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams, float * tensor_split, struct llama_model_tensor_buft_override * tensor_buft_overrides, @@ -281,28 +285,28 @@ static void llama_params_fit_impl( } if (mparams->n_gpu_layers != default_mparams.n_gpu_layers) { - throw std::runtime_error("n_gpu_layers already set by user to " + std::to_string(mparams->n_gpu_layers) + ", abort"); + throw llama_params_fit_exception("n_gpu_layers already set by user to " + std::to_string(mparams->n_gpu_layers) + ", abort"); } if (nd > 1) { if (!tensor_split) { - throw std::runtime_error("did not provide a buffer to write the tensor_split to, abort"); + throw llama_params_fit_exception("did not provide a buffer to write the tensor_split to, abort"); } if (mparams->tensor_split) { for (size_t id = 0; id < nd; id++) { if (mparams->tensor_split[id] != 0.0f) { - throw std::runtime_error("model_params::tensor_split already set by user, abort"); + throw llama_params_fit_exception("model_params::tensor_split already set by user, abort"); } } } if (mparams->split_mode == LLAMA_SPLIT_MODE_ROW) { - throw std::runtime_error("changing weight allocation for LLAMA_SPLIT_MODE_ROW not implemented, abort"); + throw llama_params_fit_exception("changing weight allocation for LLAMA_SPLIT_MODE_ROW not implemented, abort"); } } if (!tensor_buft_overrides) { - throw std::runtime_error("did not provide buffer to set tensor_buft_overrides, abort"); + throw llama_params_fit_exception("did not provide buffer to set tensor_buft_overrides, abort"); } if (mparams->tensor_buft_overrides && (mparams->tensor_buft_overrides->pattern || mparams->tensor_buft_overrides->buft)) { - throw std::runtime_error("model_params::tensor_buft_overrides already set by user, abort"); + throw llama_params_fit_exception("model_params::tensor_buft_overrides already set by user, abort"); } // step 3: iteratively fill the back to front with "dense" layers @@ -385,7 +389,7 @@ static void llama_params_fit_impl( tensor_buft_overrides[itbo].buft = nullptr; itbo++; mparams.tensor_buft_overrides = tensor_buft_overrides; - throw std::runtime_error("llama_params_fit_n_tensor_buft_overrides() == " + throw llama_params_fit_exception("llama_params_fit_n_tensor_buft_overrides() == " + std::to_string(ntbo) + " is insufficient for model\n"); } tensor_buft_overrides[itbo].pattern = get_overflow_pattern(il, il == il0 ? ngl_per_device[id].overflow_type : LAYER_FRACTION_MOE); @@ -683,22 +687,25 @@ static void llama_params_fit_impl( set_ngl_tensor_split_tbo(ngl_per_device, overflow_bufts, *mparams); } -bool llama_params_fit( +enum llama_params_fit_status llama_params_fit( const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams, float * tensor_split, struct llama_model_tensor_buft_override * tensor_buft_overrides, size_t margin_s, uint32_t n_ctx_min, enum ggml_log_level log_level) { const int64_t t0_us = llama_time_us(); - bool ok = true; + llama_params_fit_status status = LLAMA_PARAMS_FIT_STATUS_SUCCESS; try { llama_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margin_s, n_ctx_min, log_level); LLAMA_LOG_INFO("%s: successfully fit params to free device memory\n", __func__); - } catch (const std::runtime_error & e) { + } catch (const llama_params_fit_exception & e) { LLAMA_LOG_WARN("%s: failed to fit params to free device memory: %s\n", __func__, e.what()); - ok = false; + status = LLAMA_PARAMS_FIT_STATUS_FAILURE; + } catch (const std::runtime_error & e) { + LLAMA_LOG_ERROR("%s: encountered an error while trying to fit params to free device memory: %s\n", __func__, e.what()); + status = LLAMA_PARAMS_FIT_STATUS_ERROR; } const int64_t t1_us = llama_time_us(); LLAMA_LOG_INFO("%s: fitting params to free memory took %.2f seconds\n", __func__, (t1_us - t0_us) * 1e-6); - return ok; + return status; } struct llama_sampler_chain_params llama_sampler_chain_default_params() { diff --git a/tools/fit-params/fit-params.cpp b/tools/fit-params/fit-params.cpp index de47763d3..c7e7748ca 100644 --- a/tools/fit-params/fit-params.cpp +++ b/tools/fit-params/fit-params.cpp @@ -26,10 +26,10 @@ int main(int argc, char ** argv) { llama_numa_init(params.numa); auto mparams = common_model_params_to_llama(params); auto cparams = common_context_params_to_llama(params); - const bool success = llama_params_fit(params.model.path.c_str(), &mparams, &cparams, + const llama_params_fit_status status = llama_params_fit(params.model.path.c_str(), &mparams, &cparams, params.tensor_split, params.tensor_buft_overrides.data(), params.fit_params_target, params.fit_params_min_ctx, params.verbosity >= 4 ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR); - if (!success) { + if (status != LLAMA_PARAMS_FIT_STATUS_SUCCESS) { LOG_ERR("%s: failed to fit CLI arguments to free memory, exiting...\n", __func__); exit(1); } From 06705fdcb3ef199d2c2c95a5e3cbb7cf9cc5256e Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Sat, 27 Dec 2025 19:56:27 +0800 Subject: [PATCH 02/12] ggml-cuda: Use same regex for GGML_NATIVE=OFF (#18407) --- ggml/src/ggml-cuda/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt index c0f8bcaa3..3b438c30c 100644 --- a/ggml/src/ggml-cuda/CMakeLists.txt +++ b/ggml/src/ggml-cuda/CMakeLists.txt @@ -61,7 +61,7 @@ if (CUDAToolkit_FOUND) set(CMAKE_CUDA_ARCHITECTURES ${PROCESSED_ARCHITECTURES}) else() foreach(ARCH ${CMAKE_CUDA_ARCHITECTURES}) - if(ARCH MATCHES "^12[0-9]$") + if(ARCH MATCHES "^12[0-9](-real|-virtual)?$") message(FATAL_ERROR "Compute capability ${ARCH} used, use ${ARCH}a or ${ARCH}f for Blackwell specific optimizations") endif() endforeach() From 026d2ad47247c006bd9ef6e583d76b6ab69875b6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sat, 27 Dec 2025 20:18:35 +0100 Subject: [PATCH 03/12] llama: fix magic number of 999 for GPU layers (#18266) * llama: fix magic number of 999 for GPU layers * use strings for -ngl, -ngld * enacapsulate n_gpu_layers, split_mode --- common/arg.cpp | 27 +++++++++++++++++++++------ common/common.cpp | 5 +---- common/common.h | 2 +- include/llama.h | 7 ++++--- src/llama-context.cpp | 6 +++--- src/llama-model.cpp | 14 +++++++++++--- src/llama-model.h | 7 +++++-- 7 files changed, 46 insertions(+), 22 deletions(-) diff --git a/common/arg.cpp b/common/arg.cpp index 774f8731a..189470182 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2137,11 +2137,18 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } } ).set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_N_CPU_MOE_DRAFT")); + GGML_ASSERT(params.n_gpu_layers < 0); // string_format would need to be extended for a default >= 0 add_opt(common_arg( {"-ngl", "--gpu-layers", "--n-gpu-layers"}, "N", - string_format("max. number of layers to store in VRAM (default: %d)", params.n_gpu_layers), - [](common_params & params, int value) { - params.n_gpu_layers = value; + string_format("max. number of layers to store in VRAM, either an exact number, 'auto', or 'all' (default: %s)", params.n_gpu_layers == -1 ? "auto" : "all"), + [](common_params & params, const std::string & value) { + if (value == "auto") { + params.n_gpu_layers = -1; + } else if (value == "all") { + params.n_gpu_layers = -2; + } else { + params.n_gpu_layers = std::stoi(value); + } if (!llama_supports_gpu_offload()) { fprintf(stderr, "warning: no usable GPU found, --gpu-layers option will be ignored\n"); fprintf(stderr, "warning: one possible reason is that llama.cpp was compiled without GPU support\n"); @@ -3175,11 +3182,19 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.speculative.devices = parse_device_list(value); } ).set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI})); + GGML_ASSERT(params.speculative.n_gpu_layers < 0); // string_format would need to be extended for a default >= 0 add_opt(common_arg( {"-ngld", "--gpu-layers-draft", "--n-gpu-layers-draft"}, "N", - "number of layers to store in VRAM for the draft model", - [](common_params & params, int value) { - params.speculative.n_gpu_layers = value; + string_format("max. number of draft model layers to store in VRAM, either an exact number, 'auto', or 'all' (default: %s)", + params.speculative.n_gpu_layers == -1 ? "auto" : "all"), + [](common_params & params, const std::string & value) { + if (value == "auto") { + params.speculative.n_gpu_layers = -1; + } else if (value == "all") { + params.speculative.n_gpu_layers = -2; + } else { + params.speculative.n_gpu_layers = std::stoi(value); + } if (!llama_supports_gpu_offload()) { fprintf(stderr, "warning: no usable GPU found, --gpu-layers-draft option will be ignored\n"); fprintf(stderr, "warning: one possible reason is that llama.cpp was compiled without GPU support\n"); diff --git a/common/common.cpp b/common/common.cpp index acf2ec841..8d6289337 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1341,10 +1341,7 @@ struct llama_model_params common_model_params_to_llama(common_params & params) { mparams.devices = params.devices.data(); } - if (params.n_gpu_layers != -1) { - mparams.n_gpu_layers = params.n_gpu_layers; - } - + mparams.n_gpu_layers = params.n_gpu_layers; mparams.main_gpu = params.main_gpu; mparams.split_mode = params.split_mode; mparams.tensor_split = params.tensor_split; diff --git a/common/common.h b/common/common.h index 334372073..f8bc686b6 100644 --- a/common/common.h +++ b/common/common.h @@ -329,7 +329,7 @@ struct common_params { // offload params std::vector devices; // devices to use for offloading - int32_t n_gpu_layers = -1; // number of layers to store in VRAM (-1 - use default) + int32_t n_gpu_layers = -1; // number of layers to store in VRAM, -1 is auto, <= -2 is all int32_t main_gpu = 0; // the GPU that is used for scratch and small tensors float tensor_split[128] = {0}; // how split tensors should be distributed across GPUs bool fit_params = true; // whether to fit unset model/context parameters to free device memory diff --git a/include/llama.h b/include/llama.h index 2d4e9a94e..4f0124fdc 100644 --- a/include/llama.h +++ b/include/llama.h @@ -286,7 +286,7 @@ extern "C" { // NULL-terminated list of buffer types to use for tensors that match a pattern const struct llama_model_tensor_buft_override * tensor_buft_overrides; - int32_t n_gpu_layers; // number of layers to store in VRAM + int32_t n_gpu_layers; // number of layers to store in VRAM, a negative value means all layers enum llama_split_mode split_mode; // how to split the model across multiple GPUs // the GPU that is used for the entire model when split_mode is LLAMA_SPLIT_MODE_NONE @@ -474,8 +474,9 @@ extern "C" { }; // fits mparams and cparams to free device memory (assumes system memory is unlimited) - // returns true if the parameters could be successfully modified to fit device memory - // this function is NOT thread safe because it modifies the global llama logger state + // - returns true if the parameters could be successfully modified to fit device memory + // - this function is NOT thread safe because it modifies the global llama logger state + // - only parameters that have the same value as in llama_default_model_params are modified LLAMA_API enum llama_params_fit_status llama_params_fit( const char * path_model, struct llama_model_params * mparams, diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 015ebae71..1c530fdc9 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -294,8 +294,8 @@ llama_context::llama_context( // enabling pipeline parallelism in the scheduler increases memory usage, so it is only done when necessary bool pipeline_parallel = model.n_devices() > 1 && - model.params.n_gpu_layers > (int) model.hparams.n_layer && - model.params.split_mode == LLAMA_SPLIT_MODE_LAYER && + model.n_gpu_layers() > model.hparams.n_layer && + model.split_mode() == LLAMA_SPLIT_MODE_LAYER && cparams.offload_kqv && !model.has_tensor_overrides(); @@ -1570,7 +1570,7 @@ llm_graph_cb llama_context::graph_get_cb() const { // norm may be automatically assigned to the backend of the previous layer, increasing data transfer between backends // FIXME: fix in ggml_backend_sched - const bool full_offload = model.params.n_gpu_layers > (int) model.hparams.n_layer; + const bool full_offload = model.n_gpu_layers() > model.hparams.n_layer; if (ubatch.n_tokens < 32 || full_offload) { if (il != -1 && strcmp(name, "norm") == 0) { const auto & dev_layer = model.dev_layer(il); diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 69075742c..1d6134ec0 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -2378,11 +2378,11 @@ void llama_model::load_vocab(llama_model_loader & ml) { bool llama_model::load_tensors(llama_model_loader & ml) { const auto & split_mode = params.split_mode; - const auto & n_gpu_layers = params.n_gpu_layers; const auto & use_mlock = params.use_mlock; const auto & tensor_split = params.tensor_split; - const int n_layer = hparams.n_layer; + const int n_layer = hparams.n_layer; + const int n_gpu_layers = this->n_gpu_layers(); const bool use_mmap_buffer = true; @@ -6884,6 +6884,14 @@ size_t llama_model::n_devices() const { return devices.size(); } +uint32_t llama_model::n_gpu_layers() const { + return params.n_gpu_layers >= 0 ? params.n_gpu_layers : hparams.n_layer + 1; +} + +llama_split_mode llama_model::split_mode() const { + return params.split_mode; +} + std::map llama_model::memory_breakdown() const { std::map ret; for (const auto & [ctx, bufs] : pimpl->ctxs_bufs) { @@ -7794,7 +7802,7 @@ llama_model_params llama_model_default_params() { llama_model_params result = { /*.devices =*/ nullptr, /*.tensor_buft_overrides =*/ nullptr, - /*.n_gpu_layers =*/ 999, + /*.n_gpu_layers =*/ -1, /*.split_mode =*/ LLAMA_SPLIT_MODE_LAYER, /*.main_gpu =*/ 0, /*.tensor_split =*/ nullptr, diff --git a/src/llama-model.h b/src/llama-model.h index 9c00eec75..dbe5edc15 100644 --- a/src/llama-model.h +++ b/src/llama-model.h @@ -466,8 +466,6 @@ struct llama_model { struct ggml_tensor * dense_2_out_layers = nullptr; struct ggml_tensor * dense_3_out_layers = nullptr; - llama_model_params params; - // gguf metadata std::unordered_map gguf_kv; @@ -498,6 +496,9 @@ struct llama_model { size_t n_tensors() const; size_t n_devices() const; + uint32_t n_gpu_layers() const; + llama_split_mode split_mode() const; + std::map memory_breakdown() const; // total number of parameters in the model @@ -526,6 +527,8 @@ struct llama_model { ggml_cgraph * build_graph(const llm_graph_params & params) const; private: + llama_model_params params; + struct impl; std::unique_ptr pimpl; }; From a4bf35889eda36d3597cd0f8f333f5b8a2fcaefc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sat, 27 Dec 2025 20:20:45 +0100 Subject: [PATCH 04/12] llama-fit-params: fix overflow check (#18354) --- src/llama.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/src/llama.cpp b/src/llama.cpp index c53f2472b..93a9c408b 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -389,8 +389,8 @@ static void llama_params_fit_impl( tensor_buft_overrides[itbo].buft = nullptr; itbo++; mparams.tensor_buft_overrides = tensor_buft_overrides; - throw llama_params_fit_exception("llama_params_fit_n_tensor_buft_overrides() == " - + std::to_string(ntbo) + " is insufficient for model\n"); + throw llama_params_fit_exception("llama_max_tensor_buft_overrides() == " + + std::to_string(ntbo) + " is insufficient for model"); } tensor_buft_overrides[itbo].pattern = get_overflow_pattern(il, il == il0 ? ngl_per_device[id].overflow_type : LAYER_FRACTION_MOE); tensor_buft_overrides[itbo].buft = overflow_bufts[id]; @@ -647,7 +647,7 @@ static void llama_params_fit_impl( ngl_per_device_test[id].overflow_type = LAYER_FRACTION_UP; LLAMA_LOG_DEBUG("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_UP\n", __func__); std::vector mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts); - if (mem_test[id] < targets[id]) { + if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) { ngl_per_device = ngl_per_device_test; mem = mem_test; id_dense_start = id_dense_start_test; @@ -657,7 +657,7 @@ static void llama_params_fit_impl( ngl_per_device_test[id].overflow_type = LAYER_FRACTION_GATE; LLAMA_LOG_DEBUG("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_GATE\n", __func__); mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts); - if (mem_test[id] < targets[id]) { + if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) { ngl_per_device = ngl_per_device_test; mem = mem_test; id_dense_start = id_dense_start_test; @@ -668,7 +668,7 @@ static void llama_params_fit_impl( ngl_per_device_test[id].overflow_type = LAYER_FRACTION_ATTN; LLAMA_LOG_DEBUG("%s: trying to fit one extra layer with overflow_type=LAYER_FRACTION_ATTN\n", __func__); mem_test = get_memory_for_layers(__func__, ngl_per_device_test, overflow_bufts); - if (mem_test[id] < targets[id]) { + if (mem_test[id] < targets[id] && (id + 1 == nd || mem_test[id + 1] < targets[id + 1])) { ngl_per_device = ngl_per_device_test; mem = mem_test; id_dense_start = id_dense_start_test; From 08566977a74a09ee5fbe8d35846e50bb507b8a97 Mon Sep 17 00:00:00 2001 From: lhez Date: Sat, 27 Dec 2025 15:51:14 -0800 Subject: [PATCH 05/12] opencl: allow resizing transpose buffers (#18384) * opencl: allow resizing transpose buffers instead of using fixed sizes * opencl: remove commented code --- ggml/src/ggml-opencl/ggml-opencl.cpp | 56 ++++++++++++++++++++++------ 1 file changed, 44 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 639715537..353f6a4b4 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -263,6 +263,32 @@ static ggml_cl_compiler_version get_adreno_cl_compiler_version(const char *drive return { type, major, minor, patch }; } +// cl buffer wrapper +struct ggml_cl_buffer { + cl_mem buffer; + size_t size; + + ggml_cl_buffer() + : buffer(nullptr), size(0) {} + + ~ggml_cl_buffer() { + if (buffer) { + CL_CHECK(clReleaseMemObject(buffer)); + } + } + + void allocate(cl_context context, size_t new_size) { + if (new_size > size) { + size = new_size; + if (buffer) { + CL_CHECK(clReleaseMemObject(buffer)); + } + cl_int err; + CL_CHECK((buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &err), err)); + } + } +}; + // Profiling struct ProfilingInfo { std::string op_name; @@ -376,6 +402,11 @@ struct ggml_backend_opencl_context { cl_context context; cl_command_queue queue; + // prealloc buffers for transposing weights and activations + ggml_cl_buffer prealloc_quant_trans; + ggml_cl_buffer prealloc_scales_trans; + ggml_cl_buffer prealloc_act_trans; + cl_program program_add; cl_program program_add_id; cl_program program_clamp; @@ -638,10 +669,6 @@ struct ggml_backend_opencl_context { cl_kernel kernel_transpose_16_buf; cl_kernel kernel_transpose_16_4x1; - cl_mem A_s_d_max; // max scale buffer size for transpose - cl_mem A_q_d_max; // max weight buffer size for transpose - cl_mem B_d_max; // max activation buffer size for transpose - // Gemm and Gemv related programs, kernels, etc cl_program program_CL_gemm; cl_program program_CL_gemv_general; @@ -2600,9 +2627,9 @@ static ggml_backend_opencl_context * ggml_cl2_init(ggml_backend_dev_t dev) { required_B_d_bytes, max_B_d_bytes); } - CL_CHECK((backend_ctx->A_q_d_max = clCreateBuffer(context, 0, max_A_q_d_bytes, NULL, &err), err)); - CL_CHECK((backend_ctx->A_s_d_max = clCreateBuffer(context, 0, max_A_s_d_bytes, NULL, &err), err)); - CL_CHECK((backend_ctx->B_d_max = clCreateBuffer(context, 0, max_B_d_bytes, NULL, &err), err)); + backend_ctx->prealloc_quant_trans.allocate(context, max_A_q_d_bytes); + backend_ctx->prealloc_scales_trans.allocate(context, max_A_s_d_bytes); + backend_ctx->prealloc_act_trans.allocate(context, max_B_d_bytes); #endif // GGML_OPENCL_USE_ADRENO_KERNELS backend_ctx->disable_fusion = getenv("GGML_OPENCL_DISABLE_FUSION") != nullptr; @@ -3607,32 +3634,35 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, // use sub_buffer of max buffer size instead size_t q_size_bytes = K * M / 8 * sizeof(float); + backend_ctx->prealloc_quant_trans.allocate(context, q_size_bytes); + cl_buffer_region region; region.origin = 0; region.size = q_size_bytes; cl_mem qT_d = clCreateSubBuffer( - backend_ctx->A_q_d_max, + backend_ctx->prealloc_quant_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); - // cl_mem qT_d = clCreateBuffer(context, CL_MEM_READ_WRITE, q_size_bytes, NULL, &err); CL_CHECK(err); bool K_tile_trans = true; if ((K / 32) % 4 != 0){ K_tile_trans =false; } + size_t d_size_bytes = M * (K / 32) * 2; + backend_ctx->prealloc_scales_trans.allocate(context, d_size_bytes); + region.origin = 0; region.size = d_size_bytes; cl_mem dT_d = clCreateSubBuffer( - backend_ctx->A_s_d_max, + backend_ctx->prealloc_scales_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err); - // cl_mem dT_d = clCreateBuffer(context, CL_MEM_READ_WRITE, d_size_bytes, NULL, &err); CL_CHECK(err); // <----------------------------------------------------------------------------------> // @@ -7395,8 +7425,10 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co region.origin = 0; // Specify the size of the sub-buffer (divide by 2 for FP16) region.size = K * (N + padding) * sizeof(float)/2; + backend_ctx->prealloc_act_trans.allocate(context, region.size); + B_d = clCreateSubBuffer( - backend_ctx->B_d_max, + backend_ctx->prealloc_act_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, From 4fd59e8427e78241b01b1db352ef28be6bf304fd Mon Sep 17 00:00:00 2001 From: QDelta <60222316+QDelta@users.noreply.github.com> Date: Sat, 27 Dec 2025 20:33:14 -0500 Subject: [PATCH 06/12] ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON (#18413) --- ggml/src/ggml-cuda/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt index 3b438c30c..f49121754 100644 --- a/ggml/src/ggml-cuda/CMakeLists.txt +++ b/ggml/src/ggml-cuda/CMakeLists.txt @@ -37,14 +37,13 @@ if (CUDAToolkit_FOUND) endif() endif() endif() - message(STATUS "Using CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}") enable_language(CUDA) # Replace any 12x-real architectures with 12x{a}-real. FP4 ptx instructions are not available in just 12x if (GGML_NATIVE) set(PROCESSED_ARCHITECTURES "") - if (CMAKE_CUDA_ARCHITECTURES_NATIVE) + if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND CMAKE_CUDA_ARCHITECTURES_NATIVE) set(ARCH_LIST ${CMAKE_CUDA_ARCHITECTURES_NATIVE}) else() set(ARCH_LIST ${CMAKE_CUDA_ARCHITECTURES}) @@ -66,6 +65,7 @@ if (CUDAToolkit_FOUND) endif() endforeach() endif() + message(STATUS "Using CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}") file(GLOB GGML_HEADERS_CUDA "*.cuh") list(APPEND GGML_HEADERS_CUDA "../../include/ggml-cuda.h") From 94de74e7b1b9ee6404b54ebb0df273a3ef35a555 Mon Sep 17 00:00:00 2001 From: Boian Berberov <7432115+bberberov@users.noreply.github.com> Date: Sun, 28 Dec 2025 07:33:29 +0000 Subject: [PATCH 07/12] cmake: Added more x86_64 CPU backends when building with `GGML_CPU_ALL_VARIANTS=On` (#18186) * minor: Consolidated `#include ` under `ggml-cpu-impl.h` * cmake: Added more x86-64 CPU backends when building with `GGML_CPU_ALL_VARIANTS=On` - `ivybridge` - `piledriver` - `cannonlake` - `cascadelake` - `cooperlake` - `zen4` Resolves: #17966 --- ggml/CMakeLists.txt | 12 ++++++++++++ ggml/src/CMakeLists.txt | 28 +++++++++++++++++++++------- ggml/src/ggml-cpu/ggml-cpu-impl.h | 2 +- ggml/src/ggml-cpu/simd-mappings.h | 4 ---- ggml/src/ggml-impl.h | 4 ---- 5 files changed, 34 insertions(+), 16 deletions(-) diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index 18d117f7c..cb46c3210 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -430,10 +430,22 @@ if (MSVC) configure_msvc_target(ggml-cpu-x64) configure_msvc_target(ggml-cpu-sse42) configure_msvc_target(ggml-cpu-sandybridge) + # __FMA__ and __F16C__ are not defined in MSVC, however they are implied with AVX2/AVX512 + # skipping ggml-cpu-ivybridge + # skipping ggml-cpu-piledriver configure_msvc_target(ggml-cpu-haswell) configure_msvc_target(ggml-cpu-skylakex) + configure_msvc_target(ggml-cpu-cannonlake) + configure_msvc_target(ggml-cpu-cascadelake) configure_msvc_target(ggml-cpu-icelake) + # MSVC 2022 doesn't support BF16 intrinsics without `/arch:AVX10.1` ?! + # https://learn.microsoft.com/en-us/cpp/intrinsics/x64-amd64-intrinsics-list?view=msvc-170 + # https://learn.microsoft.com/en-us/cpp/build/reference/arch-x64?view=msvc-170 + # skipping ggml-cpu-cooperlake + # skipping ggml-cpu-zen4 configure_msvc_target(ggml-cpu-alderlake) + # MSVC doesn't support AMX + # skipping ggml-cpu-sapphirerapids if (GGML_BUILD_EXAMPLES) configure_msvc_target(common-ggml) diff --git a/ggml/src/CMakeLists.txt b/ggml/src/CMakeLists.txt index 262d78a4c..25f25c423 100644 --- a/ggml/src/CMakeLists.txt +++ b/ggml/src/CMakeLists.txt @@ -357,15 +357,29 @@ if (GGML_CPU_ALL_VARIANTS) endif() if (GGML_SYSTEM_ARCH STREQUAL "x86") ggml_add_cpu_backend_variant(x64) - ggml_add_cpu_backend_variant(sse42 SSE42) - ggml_add_cpu_backend_variant(sandybridge SSE42 AVX) - ggml_add_cpu_backend_variant(haswell SSE42 AVX F16C AVX2 BMI2 FMA) - ggml_add_cpu_backend_variant(skylakex SSE42 AVX F16C AVX2 BMI2 FMA AVX512) - ggml_add_cpu_backend_variant(icelake SSE42 AVX F16C AVX2 BMI2 FMA AVX512 AVX512_VBMI AVX512_VNNI) - ggml_add_cpu_backend_variant(alderlake SSE42 AVX F16C AVX2 BMI2 FMA AVX_VNNI) + ggml_add_cpu_backend_variant(sse42 SSE42) + ggml_add_cpu_backend_variant(sandybridge SSE42 AVX) + if (NOT MSVC) + # __FMA__ and __F16C__ are not defined in MSVC, however they are implied with AVX2/AVX512 + ggml_add_cpu_backend_variant(ivybridge SSE42 AVX F16C) + ggml_add_cpu_backend_variant(piledriver SSE42 AVX F16C FMA) + endif() + ggml_add_cpu_backend_variant(haswell SSE42 AVX F16C FMA AVX2 BMI2) + ggml_add_cpu_backend_variant(skylakex SSE42 AVX F16C FMA AVX2 BMI2 AVX512) + ggml_add_cpu_backend_variant(cannonlake SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VBMI) + ggml_add_cpu_backend_variant(cascadelake SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VNNI) + ggml_add_cpu_backend_variant(icelake SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VBMI AVX512_VNNI) + if (NOT MSVC) + # MSVC 2022 doesn't support BF16 intrinsics without `/arch:AVX10.1` ?! + # https://learn.microsoft.com/en-us/cpp/intrinsics/x64-amd64-intrinsics-list?view=msvc-170 + # https://learn.microsoft.com/en-us/cpp/build/reference/arch-x64?view=msvc-170 + ggml_add_cpu_backend_variant(cooperlake SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VNNI AVX512_BF16) + ggml_add_cpu_backend_variant(zen4 SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VBMI AVX512_VNNI AVX512_BF16) + endif() + ggml_add_cpu_backend_variant(alderlake SSE42 AVX F16C FMA AVX2 BMI2 AVX_VNNI) if (NOT MSVC) # MSVC doesn't support AMX - ggml_add_cpu_backend_variant(sapphirerapids SSE42 AVX F16C AVX2 BMI2 FMA AVX512 AVX512_VBMI AVX512_VNNI AVX512_BF16 AMX_TILE AMX_INT8) + ggml_add_cpu_backend_variant(sapphirerapids SSE42 AVX F16C FMA AVX2 BMI2 AVX512 AVX512_VBMI AVX512_VNNI AVX512_BF16 AMX_TILE AMX_INT8) endif() elseif(GGML_SYSTEM_ARCH STREQUAL "ARM") if (CMAKE_SYSTEM_NAME MATCHES "Linux") diff --git a/ggml/src/ggml-cpu/ggml-cpu-impl.h b/ggml/src/ggml-cpu/ggml-cpu-impl.h index 7597377cc..0e8dd0ae0 100644 --- a/ggml/src/ggml-cpu/ggml-cpu-impl.h +++ b/ggml/src/ggml-cpu/ggml-cpu-impl.h @@ -328,7 +328,7 @@ inline static int32x4_t ggml_vdotq_s32(int32x4_t acc, int8x16_t a, int8x16_t b) #if defined(_MSC_VER) || defined(__MINGW32__) #include -#elif defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__SSSE3__) || defined(__SSE3__) || defined(__SSE__) +#elif defined(__SSE__) || defined(__SSE3__) || defined(__SSSE3__) || defined(__AVX__) || defined(__F16C__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__AVX512BF16__) #include #endif diff --git a/ggml/src/ggml-cpu/simd-mappings.h b/ggml/src/ggml-cpu/simd-mappings.h index 101a9c086..a7a827220 100644 --- a/ggml/src/ggml-cpu/simd-mappings.h +++ b/ggml/src/ggml-cpu/simd-mappings.h @@ -14,10 +14,6 @@ #include #endif -#if defined(__F16C__) -#include -#endif - #if defined(__riscv_v_intrinsic) #include #endif diff --git a/ggml/src/ggml-impl.h b/ggml/src/ggml-impl.h index fe57d4c58..80e0fd2ff 100644 --- a/ggml/src/ggml-impl.h +++ b/ggml/src/ggml-impl.h @@ -24,10 +24,6 @@ #include #endif -#if defined(__F16C__) -#include -#endif - #ifdef __cplusplus extern "C" { #endif From cffa5c46eaf6eddb0db27eeef1c76204fae98984 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sun, 28 Dec 2025 09:57:04 +0100 Subject: [PATCH 08/12] mtmd: clarify that we no longer accept AI-generated PRs (#18406) --- tools/mtmd/models/models.h | 5 +++++ tools/mtmd/mtmd.h | 3 +++ 2 files changed, 8 insertions(+) diff --git a/tools/mtmd/models/models.h b/tools/mtmd/models/models.h index 8d6d4ef67..e08c33f35 100644 --- a/tools/mtmd/models/models.h +++ b/tools/mtmd/models/models.h @@ -2,6 +2,11 @@ #include "../clip-graph.h" +/* + * IMPORTANT: The mtmd module does NOT accept pull requests that are fully or predominantly AI-generated. + * We encourage human contributors to ensure the quality and reliability of the codebase. + */ + struct clip_graph_siglip : clip_graph { clip_graph_siglip(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} ggml_cgraph * build() override; diff --git a/tools/mtmd/mtmd.h b/tools/mtmd/mtmd.h index 9f7e861e9..44d05ceae 100644 --- a/tools/mtmd/mtmd.h +++ b/tools/mtmd/mtmd.h @@ -27,6 +27,9 @@ * - Make sure the C API is aligned with the libllama C API (as in llama.h) * - Do not include model name (e.g., qwen, gemma) in the API, use generic terms instead * - Keep the API minimal, do not expose internal details unless necessary + * + * IMPORTANT: The mtmd module does NOT accept pull requests that are fully or predominantly AI-generated. + * We encourage human contributors to ensure the quality and reliability of the codebase. */ #ifdef LLAMA_SHARED From e59efe6a78748e8447bc69db2d0fe51332f92c02 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sun, 28 Dec 2025 10:50:56 +0100 Subject: [PATCH 09/12] github: update issue templates [no ci] (#18410) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * github: update issue templates [no ci] * Apply suggestions from code review Co-authored-by: Sigbjørn Skjæret --------- Co-authored-by: Sigbjørn Skjæret --- .github/ISSUE_TEMPLATE/010-bug-compilation.yml | 3 ++- .github/ISSUE_TEMPLATE/011-bug-results.yml | 15 +++++++++++++-- .github/ISSUE_TEMPLATE/019-bug-misc.yml | 15 +++++++++++++-- 3 files changed, 28 insertions(+), 5 deletions(-) diff --git a/.github/ISSUE_TEMPLATE/010-bug-compilation.yml b/.github/ISSUE_TEMPLATE/010-bug-compilation.yml index feb0d5120..c106f47a2 100644 --- a/.github/ISSUE_TEMPLATE/010-bug-compilation.yml +++ b/.github/ISSUE_TEMPLATE/010-bug-compilation.yml @@ -8,7 +8,8 @@ body: value: > Thanks for taking the time to fill out this bug report! This issue template is intended for bug reports where the compilation of llama.cpp fails. - Before opening an issue, please confirm that the compilation still fails with `-DGGML_CCACHE=OFF`. + Before opening an issue, please confirm that the compilation still fails + after recreating the CMake build directory and with `-DGGML_CCACHE=OFF`. If the compilation succeeds with ccache disabled you should be able to permanently fix the issue by clearing `~/.cache/ccache` (on Linux). - type: textarea diff --git a/.github/ISSUE_TEMPLATE/011-bug-results.yml b/.github/ISSUE_TEMPLATE/011-bug-results.yml index b815e70a8..31202dfa8 100644 --- a/.github/ISSUE_TEMPLATE/011-bug-results.yml +++ b/.github/ISSUE_TEMPLATE/011-bug-results.yml @@ -98,7 +98,18 @@ body: label: Relevant log output description: > Please copy and paste any relevant log output, including the command that you entered and any generated text. - This will be automatically formatted into code, so no need for backticks. - render: shell + For very long logs (thousands of lines), preferably upload them as files instead. + On Linux you can redirect console output into a file by appending ` > llama.log 2>&1` to your command. + value: | +
+ Logs + + + ```console + + ``` +
+ + validations: required: true diff --git a/.github/ISSUE_TEMPLATE/019-bug-misc.yml b/.github/ISSUE_TEMPLATE/019-bug-misc.yml index e1bd08ddd..8e867e7f6 100644 --- a/.github/ISSUE_TEMPLATE/019-bug-misc.yml +++ b/.github/ISSUE_TEMPLATE/019-bug-misc.yml @@ -85,8 +85,19 @@ body: label: Relevant log output description: > If applicable, please copy and paste any relevant log output, including any generated text. - This will be automatically formatted into code, so no need for backticks. If you are encountering problems specifically with the `llama_params_fit` module, always upload `--verbose` logs as well. - render: shell + For very long logs (thousands of lines), please upload them as files instead. + On Linux you can redirect console output into a file by appending ` > llama.log 2>&1` to your command. + value: | +
+ Logs + + + ```console + + ``` +
+ + validations: required: false From f8d561eb87cf689afb32ae5ee72118ddffaef12b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sun, 28 Dec 2025 10:52:09 +0100 Subject: [PATCH 10/12] llama-fit-params: fix step size for last device (#18415) --- src/llama.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/llama.cpp b/src/llama.cpp index 93a9c408b..76b3acbad 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -512,6 +512,9 @@ static void llama_params_fit_impl( if (mem_high[id] > targets[id]) { assert(ngl_per_device_high[id].n_layer > ngl_per_device[id].n_layer); uint32_t delta = ngl_per_device_high[id].n_layer - ngl_per_device[id].n_layer; + if (hp_nex > 0 && size_t(id) == nd - 1) { + delta--; + } LLAMA_LOG_DEBUG("%s: start filling device %" PRIu32 ", delta=%" PRIu32 "\n", __func__, id, delta); while (delta > 1) { uint32_t step_size = int64_t(delta) * (targets[id] - mem[id]) / (mem_high[id] - mem[id]); From 60f17f56da78d8fabc658279fd25e109700122da Mon Sep 17 00:00:00 2001 From: o7si <32285332+o7si@users.noreply.github.com> Date: Sun, 28 Dec 2025 18:34:41 +0800 Subject: [PATCH 11/12] rpc: fix segfault on invalid endpoint format (#18387) * rpc: fix segfault on invalid endpoint format * rpc: add error log for failed endpoint connection --- common/arg.cpp | 2 +- ggml/src/ggml-rpc/ggml-rpc.cpp | 5 +++++ 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/common/arg.cpp b/common/arg.cpp index 189470182..62d31393c 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2017,7 +2017,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex if (llama_supports_rpc()) { add_opt(common_arg( {"--rpc"}, "SERVERS", - "comma separated list of RPC servers", + "comma separated list of RPC servers (host:port)", [](common_params & params, const std::string & value) { add_rpc_devices(value); GGML_UNUSED(params); diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp index e7890a5ee..164b39d01 100644 --- a/ggml/src/ggml-rpc/ggml-rpc.cpp +++ b/ggml/src/ggml-rpc/ggml-rpc.cpp @@ -524,6 +524,7 @@ static std::shared_ptr get_socket(const std::string & endpoint) { std::string host; int port; if (!parse_endpoint(endpoint, host, port)) { + GGML_LOG_ERROR("Failed to parse endpoint: %s\n", endpoint.c_str()); return nullptr; } #ifdef _WIN32 @@ -2053,6 +2054,10 @@ ggml_backend_reg_t ggml_backend_rpc_reg(void) { static uint32_t ggml_backend_rpc_get_device_count(const char * endpoint) { auto sock = get_socket(endpoint); + if (sock == nullptr) { + GGML_LOG_ERROR("Failed to connect to %s\n", endpoint); + return 0; + } rpc_msg_device_count_rsp response; bool status = send_rpc_cmd(sock, RPC_CMD_DEVICE_COUNT, nullptr, 0, &response, sizeof(response)); RPC_STATUS_ASSERT(status); From 07a0c4ba923a8caea3eeba060d399a28c461deac Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Sun, 28 Dec 2025 20:53:36 +0800 Subject: [PATCH 12/12] Revert "ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON (#18413)" (#18426) --- ggml/src/ggml-cuda/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt index f49121754..3b438c30c 100644 --- a/ggml/src/ggml-cuda/CMakeLists.txt +++ b/ggml/src/ggml-cuda/CMakeLists.txt @@ -37,13 +37,14 @@ if (CUDAToolkit_FOUND) endif() endif() endif() + message(STATUS "Using CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}") enable_language(CUDA) # Replace any 12x-real architectures with 12x{a}-real. FP4 ptx instructions are not available in just 12x if (GGML_NATIVE) set(PROCESSED_ARCHITECTURES "") - if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND CMAKE_CUDA_ARCHITECTURES_NATIVE) + if (CMAKE_CUDA_ARCHITECTURES_NATIVE) set(ARCH_LIST ${CMAKE_CUDA_ARCHITECTURES_NATIVE}) else() set(ARCH_LIST ${CMAKE_CUDA_ARCHITECTURES}) @@ -65,7 +66,6 @@ if (CUDAToolkit_FOUND) endif() endforeach() endif() - message(STATUS "Using CUDA architectures: ${CMAKE_CUDA_ARCHITECTURES}") file(GLOB GGML_HEADERS_CUDA "*.cuh") list(APPEND GGML_HEADERS_CUDA "../../include/ggml-cuda.h")