mirror of
https://github.com/LostRuins/koboldcpp.git
synced 2026-10-04 03:56:54 +00:00
Merge branch 'upstream' into concedo_experimental
# Conflicts: # .github/workflows/server-self-hosted.yml # .github/workflows/ui-build-self-hosted.yml # .github/workflows/ui-build.yml # .github/workflows/ui.yml # ggml/src/ggml-sycl/mmvq.cpp # ggml/src/ggml-sycl/vecdotq.hpp # tests/test-backend-ops.cpp # tests/test-sampling.cpp # tools/cli/README.md # tools/server/README.md
This commit is contained in:
Concedo
commit
ad1eee0c5f
18 changed files with 337 additions and 25 deletions
|
|
@ -895,6 +895,12 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
|
|||
postprocess_cpu_params(params.speculative.draft.cpuparams, ¶ms.cpuparams);
|
||||
postprocess_cpu_params(params.speculative.draft.cpuparams_batch, ¶ms.cpuparams_batch);
|
||||
|
||||
// default the mmproj device to the global device selection if not set explicitly with -mmdev
|
||||
if (params.mmproj_use_gpu && params.mmproj_device == nullptr && !params.devices.empty()) {
|
||||
params.mmproj_device = params.devices.front();
|
||||
params.mmproj_use_gpu = params.mmproj_device != nullptr;
|
||||
}
|
||||
|
||||
if (params.prompt_cache_all && (params.interactive || params.interactive_first)) {
|
||||
throw std::invalid_argument("error: --prompt-cache-all not supported in interactive mode yet\n");
|
||||
}
|
||||
|
|
@ -2611,7 +2617,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
|
|||
add_opt(common_arg(
|
||||
// note: "-mmdev" must sort after "--rpc" in the preset map, else RPC devices are not registered yet
|
||||
{"-mmdev", "--mmproj-device"}, "DEVICE",
|
||||
"device to use for multimodal projector (none = don't offload, default: auto)\n"
|
||||
"device to use for multimodal projector (none = don't offload, default: follows --device)\n"
|
||||
"use --list-devices to see a list of available devices",
|
||||
[](common_params & params, const std::string & value) {
|
||||
if (value == "none") {
|
||||
|
|
@ -4230,7 +4236,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
|
|||
).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_BACKEND_SAMPLING"));
|
||||
add_opt(common_arg(
|
||||
{"--spec-draft-device", "-devd", "--device-draft"}, "<dev1,dev2,..>",
|
||||
"comma-separated list of devices to use for offloading the draft model (none = don't offload)\n"
|
||||
"comma-separated list of devices to use for offloading the draft model (none = don't offload, default: follows --device)\n"
|
||||
"use --list-devices to see a list of available devices",
|
||||
[](common_params & params, const std::string & value) {
|
||||
params.speculative.draft.devices = parse_device_list(value);
|
||||
|
|
|
|||
|
|
@ -2467,11 +2467,22 @@ common_params common_base_params_to_speculative(const common_params & params) {
|
|||
result.pooling_type = LLAMA_POOLING_TYPE_UNSPECIFIED;
|
||||
|
||||
if (has_draft) {
|
||||
result.devices = params_spec.devices;
|
||||
// default to global devices value
|
||||
if (!params_spec.devices.empty()) {
|
||||
result.devices = params_spec.devices;
|
||||
}
|
||||
result.model = params_spec.mparams;
|
||||
result.n_gpu_layers = params_spec.n_gpu_layers;
|
||||
result.tensor_buft_overrides = params_spec.tensor_buft_overrides;
|
||||
|
||||
// a draft pinned to a single device doesn't need the meta wrapper an inherited -sm tensor would give it
|
||||
// (the device list is null-terminated, so a single device means size 2)
|
||||
const size_t n_devs = std::count_if(params_spec.devices.begin(), params_spec.devices.end(),
|
||||
[](ggml_backend_dev_t d) { return d != nullptr; });
|
||||
if (n_devs == 1) {
|
||||
result.split_mode = LLAMA_SPLIT_MODE_LAYER;
|
||||
}
|
||||
|
||||
if (params_spec.cpuparams.n_threads > 0) {
|
||||
result.cpuparams.n_threads = params_spec.cpuparams.n_threads;
|
||||
result.cpuparams_batch.n_threads = params_spec.cpuparams_batch.n_threads;
|
||||
|
|
|
|||
|
|
@ -78,7 +78,7 @@ struct ggml_compute_params {
|
|||
#if defined(__ARM_NEON)
|
||||
|
||||
// ref: https://github.com/ggml-org/llama.cpp/pull/5404
|
||||
#ifdef _MSC_VER
|
||||
#if defined(_MSC_VER) && !defined(__clang__)
|
||||
#define ggml_vld1q_u32(w,x,y,z) { ((w) + ((uint64_t)(x) << 32)), ((y) + ((uint64_t)(z) << 32)) }
|
||||
#else
|
||||
#define ggml_vld1q_u32(w,x,y,z) { (w), (x), (y), (z) }
|
||||
|
|
|
|||
|
|
@ -306,11 +306,7 @@ static ggml_cuda_device_info ggml_cuda_init() {
|
|||
|
||||
info.default_tensor_split[id] = total_vram;
|
||||
total_vram += device_vram;
|
||||
#if defined(GGML_USE_HIP)
|
||||
info.devices[id].integrated = prop.integrated;
|
||||
#else
|
||||
info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034)
|
||||
#endif
|
||||
info.devices[id].nsm = prop.multiProcessorCount;
|
||||
info.devices[id].smpb = prop.sharedMemPerBlock;
|
||||
info.devices[id].warp_size = prop.warpSize;
|
||||
|
|
|
|||
|
|
@ -839,6 +839,8 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
|||
|
||||
const char * suffix = "";
|
||||
|
||||
bool split = false;
|
||||
|
||||
// use custom matrix x vector kernel
|
||||
switch (tsrc0) {
|
||||
case GGML_TYPE_F32:
|
||||
|
|
@ -942,6 +944,13 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
|||
nsg = N_SG_IQ3_XXS;
|
||||
nr0 = N_R0_IQ3_XXS;
|
||||
smem = 256*4+128;
|
||||
|
||||
// split the rows across threads when there are fewer than 32 chunks per row
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ3_XXS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ3_S:
|
||||
{
|
||||
|
|
@ -993,7 +1002,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
|||
const int16_t r3 = (int16_t) (ne13 / ne03);
|
||||
|
||||
snprintf(base, 256, "kernel_mul_mv_%s_%s%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1), suffix);
|
||||
snprintf(name, 256, "%s_nsg=%d_ne12=%d_r2=%d_r3=%d", base, nsg, ne12, r2, r3);
|
||||
snprintf(name, 256, "%s_nsg=%d_ne12=%d_r2=%d_r3=%d_split=%d", base, nsg, ne12, r2, r3, split);
|
||||
|
||||
ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
|
||||
if (!res.pipeline) {
|
||||
|
|
@ -1003,6 +1012,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
|||
ggml_metal_cv_set_int16(cv, (int16_t) ne12, FC_MUL_MV + 2);
|
||||
ggml_metal_cv_set_int16(cv, r2, FC_MUL_MV + 3);
|
||||
ggml_metal_cv_set_int16(cv, r3, FC_MUL_MV + 4);
|
||||
ggml_metal_cv_set_bool (cv, split, FC_MUL_MV + 5);
|
||||
|
||||
res = ggml_metal_library_compile_pipeline(lib, base, name, cv);
|
||||
|
||||
|
|
@ -1081,6 +1091,8 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
|||
|
||||
const char * suffix = "";
|
||||
|
||||
bool split = false;
|
||||
|
||||
// use custom matrix x vector kernel
|
||||
switch (tsrc0) {
|
||||
case GGML_TYPE_F32:
|
||||
|
|
@ -1177,6 +1189,13 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
|||
nsg = N_SG_IQ3_XXS;
|
||||
nr0 = N_R0_IQ3_XXS;
|
||||
smem = 256*4+128;
|
||||
|
||||
// split the rows across threads when there are fewer than 32 chunks per row
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ3_XXS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ3_S:
|
||||
{
|
||||
|
|
@ -1224,7 +1243,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
|||
};
|
||||
|
||||
snprintf(base, 256, "kernel_mul_mv_id_%s_%s%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1), suffix);
|
||||
snprintf(name, 256, "%s_nsg=%d", base, nsg);
|
||||
snprintf(name, 256, "%s_nsg=%d_split=%d", base, nsg, split);
|
||||
|
||||
ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
|
||||
if (!res.pipeline) {
|
||||
|
|
@ -1234,6 +1253,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
|||
ggml_metal_cv_set_int16(cv, 1, FC_MUL_MV + 2);
|
||||
ggml_metal_cv_set_int16(cv, 1, FC_MUL_MV + 3);
|
||||
ggml_metal_cv_set_int16(cv, 1, FC_MUL_MV + 4);
|
||||
ggml_metal_cv_set_bool (cv, split, FC_MUL_MV + 5);
|
||||
|
||||
res = ggml_metal_library_compile_pipeline(lib, base, name, cv);
|
||||
|
||||
|
|
|
|||
|
|
@ -77,6 +77,7 @@
|
|||
|
||||
#define N_R0_IQ3_XXS 4
|
||||
#define N_SG_IQ3_XXS 2
|
||||
#define N_R0_IQ3_XXS_SPLIT 8
|
||||
|
||||
#define N_R0_IQ3_S 4
|
||||
#define N_SG_IQ3_S 2
|
||||
|
|
|
|||
|
|
@ -213,6 +213,7 @@ constant short FC_mul_mv_nxpsg [[function_constant(FC_MUL_MV + 1)]];
|
|||
constant short FC_mul_mv_ne12 [[function_constant(FC_MUL_MV + 2)]];
|
||||
constant short FC_mul_mv_r2 [[function_constant(FC_MUL_MV + 3)]];
|
||||
constant short FC_mul_mv_r3 [[function_constant(FC_MUL_MV + 4)]];
|
||||
constant bool FC_mul_mv_split [[function_constant(FC_MUL_MV + 5)]];
|
||||
|
||||
template<typename block_q_type, short NR0, typename args_t>
|
||||
void mul_vec_q_n_f32_impl(
|
||||
|
|
@ -2092,6 +2093,7 @@ kernel void kernel_mul_mv_iq2_xs_f32(
|
|||
kernel_mul_mv_iq2_xs_f32_impl<N_R0_IQ2_XS, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
// FC_mul_mv_split: for nb32 < 32 (nb32 divides 32), 32/nb32 threads share each chunk and each takes a slice of the rows
|
||||
template<int nr0, typename args_t>
|
||||
void kernel_mul_mv_iq3_xxs_f32_impl(
|
||||
args_t args,
|
||||
|
|
@ -2138,11 +2140,18 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
|||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
const int ix = tiisg;
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
|
|
@ -2151,11 +2160,11 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
|||
const int ib = ib32 % (QK_K / 32);
|
||||
|
||||
device const block_iq3_xxs * xr = x + ibl;
|
||||
device const uint8_t * q3 = xr->qs + 8 * ib;
|
||||
device const uint16_t * gas = (device const uint16_t *)(xr->qs + QK_K/4) + 2 * ib;
|
||||
device const half * dh = &xr->d;
|
||||
device const uint8_t * q3 = xr->qs + 8 * ib + (uint64_t) row0*args.nb01;
|
||||
device const uint16_t * gas = (device const uint16_t *)(xr->qs + QK_K/4) + 2 * ib + (uint64_t) row0*args.nb01/2;
|
||||
device const half * dh = &xr->d + (uint64_t) row0*args.nb01/2;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
const uint32_t aux32 = gas[0] | (gas[1] << 16);
|
||||
const float d = db * (0.5f + (aux32 >> 28));
|
||||
|
|
@ -2177,7 +2186,7 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
|||
gas += args.nb01/2;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
|
|
@ -2190,6 +2199,23 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
|||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq3_xxs_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq3_xxs_f32_impl<N_R0_IQ3_XXS_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq3_xxs_f32_impl<N_R0_IQ3_XXS, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq3_xxs_f32")]]
|
||||
kernel void kernel_mul_mv_iq3_xxs_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
|
|
@ -2201,7 +2227,7 @@ kernel void kernel_mul_mv_iq3_xxs_f32(
|
|||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq3_xxs_f32_impl<N_R0_IQ3_XXS, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq3_xxs_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int nr0, typename args_t>
|
||||
|
|
@ -3217,7 +3243,7 @@ template [[host_name("kernel_mul_mv_id_iq1_s_f32")]] kernel kernel_mul_mv_id_t
|
|||
template [[host_name("kernel_mul_mv_id_iq1_m_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq1_m_f32_impl <N_R0_IQ1_M>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xxs_f32_impl<N_R0_IQ2_XXS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xs_f32_impl <N_R0_IQ2_XS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_xxs_f32_impl<N_R0_IQ3_XXS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_xxs_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_s_f32_impl <N_R0_IQ3_S>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_s_f32_impl <N_R0_IQ2_S>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq4_nl_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_nl_f32_impl <N_R0_IQ4_NL>>>;
|
||||
|
|
|
|||
|
|
@ -101,6 +101,14 @@ typedef struct VkPhysicalDeviceCooperativeMatrixDecodeVectorFeaturesNV {
|
|||
#include "ggml-vulkan-shaders-noext.hpp"
|
||||
#endif
|
||||
|
||||
// On 32-bit platforms, Vulkan non-dispatchable handles such as VkBuffer are represented as uint64_t,
|
||||
// and Vulkan-Hpp disables implicit conversions for type safety.
|
||||
namespace {
|
||||
inline std::ostream & operator<<(std::ostream & os, vk::Buffer buffer) {
|
||||
return os << static_cast<VkBuffer>(buffer);
|
||||
}
|
||||
}
|
||||
|
||||
// remove this once it's more widely available in the SDK
|
||||
#if !defined(VK_KHR_shader_bfloat16)
|
||||
|
||||
|
|
@ -1069,6 +1077,9 @@ struct vk_device_struct {
|
|||
vk_pipeline pipeline_trunc[2];
|
||||
vk_pipeline pipeline_sgn[2];
|
||||
|
||||
// fused UNARY+MUL pipelines: [op][f16][norepeat][op_on_b]
|
||||
vk_pipeline pipeline_unary_mul[4][2][2][2];
|
||||
|
||||
vk_pipeline pipeline_add1_f16_f16;
|
||||
vk_pipeline pipeline_add1_f16_f32;
|
||||
vk_pipeline pipeline_add1_f32_f32;
|
||||
|
|
@ -5928,6 +5939,26 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
|
|||
CREATE_UNARY(expm1)
|
||||
#undef CREATE_UNARY
|
||||
|
||||
// spec constants: {norepeat, op_on_b}
|
||||
#define CREATE_UNARY_MUL(name, idx) \
|
||||
for (int dt = 0; dt < 2; ++dt) { \
|
||||
const size_t len_ = dt ? name ## _mul_f16_len : name ## _mul_f32_len; \
|
||||
const unsigned char * data_ = dt ? name ## _mul_f16_data : name ## _mul_f32_data; \
|
||||
const std::string dts_ = dt ? "f16" : "f32"; \
|
||||
for (int ob = 0; ob < 2; ++ob) \
|
||||
for (int nr = 0; nr < 2; ++nr) \
|
||||
ggml_vk_create_pipeline(device, device->pipeline_unary_mul[(idx)][dt][nr][ob], \
|
||||
(#name "_mul" + std::string(ob ? "_b" : "") + "_" + dts_ + (nr ? "_norepeat" : "")).c_str(), \
|
||||
len_, data_, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, \
|
||||
{ (uint32_t) nr, (uint32_t) ob }, 1); \
|
||||
}
|
||||
|
||||
CREATE_UNARY_MUL(gelu, 0)
|
||||
CREATE_UNARY_MUL(sigmoid, 1)
|
||||
CREATE_UNARY_MUL(silu, 2)
|
||||
CREATE_UNARY_MUL(softplus, 3)
|
||||
#undef CREATE_UNARY_MUL
|
||||
|
||||
ggml_vk_create_pipeline(device, device->pipeline_add1_f16_f16, "add1_f16_f16", add1_f16_f16_len, add1_f16_f16_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
|
||||
ggml_vk_create_pipeline(device, device->pipeline_add1_f16_f32, "add1_f16_f32", add1_f16_f32_len, add1_f16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
|
||||
ggml_vk_create_pipeline(device, device->pipeline_add1_f32_f32, "add1_f32_f32", add1_f32_f32_len, add1_f32_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {512, 1, 1}, {}, 1);
|
||||
|
|
@ -8775,7 +8806,7 @@ static bool ggml_vk_buffer_write_2d_async(vk_context subctx, vk_buffer& dst, siz
|
|||
}
|
||||
|
||||
ggml_vk_sync_buffers(nullptr, subctx);
|
||||
subctx->s->buffer->buf.copyBuffer((VkBuffer)staging_buffer->buffer, (VkBuffer)dst->buffer, slices);
|
||||
subctx->s->buffer->buf.copyBuffer(staging_buffer->buffer, dst->buffer, slices);
|
||||
|
||||
if (width == spitch) {
|
||||
deferred_memcpy((uint8_t *)staging_buffer->ptr, src, staging_size, &subctx->in_memcpys);
|
||||
|
|
@ -12441,7 +12472,7 @@ template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk
|
|||
}
|
||||
|
||||
template<typename PC>
|
||||
static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst, ggml_op op, PC&& pc) {
|
||||
static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst, ggml_op op, PC&& pc, vk_pipeline pipeline_override = nullptr) {
|
||||
VK_LOG_DEBUG("ggml_vk_op_f32((" << src0 << ", name=" << src0->name << ", type=" << src0->type << ", ne0=" << src0->ne[0] << ", ne1=" << src0->ne[1] << ", ne2=" << src0->ne[2] << ", ne3=" << src0->ne[3] << ", nb0=" << src0->nb[0] << ", nb1=" << src0->nb[1] << ", nb2=" << src0->nb[2] << ", nb3=" << src0->nb[3];
|
||||
if (src1 != nullptr) {
|
||||
std::cerr << "), (" << src1 << ", name=" << src1->name << ", type=" << src1->type << ", ne0=" << src1->ne[0] << ", ne1=" << src1->ne[1] << ", ne2=" << src1->ne[2] << ", ne3=" << src1->ne[3] << ", nb0=" << src1->nb[0] << ", nb1=" << src1->nb[1] << ", nb2=" << src1->nb[2] << ", nb3=" << src1->nb[3];
|
||||
|
|
@ -12472,7 +12503,12 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
|
|||
|
||||
init_pushconst_fastdiv(pc);
|
||||
|
||||
vk_pipeline pipeline = ggml_vk_op_get_pipeline(ctx, src0, src1, src2, dst, op);
|
||||
vk_pipeline pipeline;
|
||||
if (pipeline_override) {
|
||||
pipeline = pipeline_override;
|
||||
} else {
|
||||
pipeline = ggml_vk_op_get_pipeline(ctx, src0, src1, src2, dst, op);
|
||||
}
|
||||
|
||||
if (pipeline == nullptr) {
|
||||
std::cerr << "ggml_vulkan: Error: Missing op: " << ggml_op_name(op) << " for " << ggml_type_name(src0->type);
|
||||
|
|
@ -13057,6 +13093,52 @@ static void ggml_vk_mul(ggml_backend_vk_context * ctx, vk_context& subctx, const
|
|||
});
|
||||
}
|
||||
|
||||
// index into device->pipeline_unary_mul for the supported unary ops, or -1
|
||||
static int ggml_vk_unary_mul_op_index(ggml_unary_op op) {
|
||||
switch (op) {
|
||||
case GGML_UNARY_OP_GELU: return 0;
|
||||
case GGML_UNARY_OP_SIGMOID: return 1;
|
||||
case GGML_UNARY_OP_SILU: return 2;
|
||||
case GGML_UNARY_OP_SOFTPLUS: return 3;
|
||||
default: return -1;
|
||||
}
|
||||
}
|
||||
|
||||
static void ggml_vk_unary_mul(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx) {
|
||||
const ggml_tensor * unary = cgraph->nodes[node_idx];
|
||||
ggml_tensor * mul = cgraph->nodes[node_idx + 1];
|
||||
|
||||
// unary on src1 that tiles into src0
|
||||
const bool op_on_b = mul->src[1] == unary &&
|
||||
!ggml_are_same_shape(unary->src[0], mul->src[0]) &&
|
||||
ggml_can_repeat(unary, mul->src[0]);
|
||||
|
||||
const ggml_tensor * src0 = op_on_b ? mul->src[0] : unary->src[0];
|
||||
const ggml_tensor * src1 = op_on_b ? unary->src[0] :
|
||||
((mul->src[0] == unary) ? mul->src[1] : mul->src[0]);
|
||||
|
||||
const bool f16 = src0->type == GGML_TYPE_F16;
|
||||
const bool norepeat = ggml_are_same_shape(src0, src1);
|
||||
const int oi = ggml_vk_unary_mul_op_index(ggml_get_unary_op(unary));
|
||||
if (oi < 0) {
|
||||
GGML_ABORT("fatal error");
|
||||
}
|
||||
vk_pipeline pipeline = ctx->device->pipeline_unary_mul[oi][f16][norepeat][op_on_b];
|
||||
|
||||
const uint32_t src0_type_size = ggml_type_size(src0->type);
|
||||
const uint32_t src1_type_size = ggml_type_size(src1->type);
|
||||
const uint32_t dst_type_size = ggml_type_size(mul->type);
|
||||
|
||||
ggml_vk_op_f32<vk_op_binary_push_constants>(ctx, subctx, src0, src1, nullptr, nullptr, mul, GGML_OP_UNARY, {
|
||||
(uint32_t)ggml_nelements(op_on_b ? mul : src0),
|
||||
(uint32_t)src0->ne[0], (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], (uint32_t)src0->nb[0] / src0_type_size, (uint32_t)src0->nb[1] / src0_type_size, (uint32_t)src0->nb[2] / src0_type_size, (uint32_t)src0->nb[3] / src0_type_size,
|
||||
(uint32_t)src1->ne[0], (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], (uint32_t)src1->nb[0] / src1_type_size, (uint32_t)src1->nb[1] / src1_type_size, (uint32_t)src1->nb[2] / src1_type_size, (uint32_t)src1->nb[3] / src1_type_size,
|
||||
(uint32_t) mul->ne[0], (uint32_t) mul->ne[1], (uint32_t) mul->ne[2],(uint32_t) mul->ne[3], (uint32_t) mul->nb[0] / dst_type_size, (uint32_t) mul->nb[1] / dst_type_size, (uint32_t) mul->nb[2] / dst_type_size, (uint32_t) mul->nb[3] / dst_type_size,
|
||||
0,
|
||||
0.0f, 0.0f, 0,
|
||||
}, pipeline);
|
||||
}
|
||||
|
||||
static void ggml_vk_div(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
|
||||
const uint32_t src0_type_size = ggml_type_size(src0->type);
|
||||
const uint32_t src1_type_size = ggml_type_size(src1->type);
|
||||
|
|
@ -16339,6 +16421,10 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr
|
|||
ggml_vk_topk_moe(ctx, compute_ctx, cgraph, node_idx);
|
||||
break;
|
||||
}
|
||||
if (ctx->num_additional_fused_ops) {
|
||||
ggml_vk_unary_mul(ctx, compute_ctx, cgraph, node_idx);
|
||||
break;
|
||||
}
|
||||
|
||||
switch (ggml_get_unary_op(node)) {
|
||||
case GGML_UNARY_OP_ELU:
|
||||
|
|
@ -17300,7 +17386,48 @@ static bool ggml_vk_is_empty(ggml_tensor * node) {
|
|||
return ggml_is_empty(node) || node->op == GGML_OP_NONE || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE;
|
||||
}
|
||||
|
||||
static bool ggml_vk_can_fuse_unary_mul(const struct ggml_cgraph * cgraph, int unary_idx, int mul_idx) {
|
||||
const ggml_tensor * unary = cgraph->nodes[unary_idx];
|
||||
const ggml_tensor * mul = cgraph->nodes[mul_idx];
|
||||
|
||||
if (ggml_vk_unary_mul_op_index(ggml_get_unary_op(unary)) < 0) {
|
||||
return false;
|
||||
}
|
||||
if (unary->type != GGML_TYPE_F32 && unary->type != GGML_TYPE_F16) {
|
||||
return false;
|
||||
}
|
||||
if (unary->type != mul->type) {
|
||||
return false;
|
||||
}
|
||||
if (mul->src[0] != unary && mul->src[1] != unary) {
|
||||
return false;
|
||||
}
|
||||
const ggml_tensor * other = (mul->src[0] == unary) ? mul->src[1] : mul->src[0];
|
||||
if (other == nullptr || other->type != unary->type) {
|
||||
return false;
|
||||
}
|
||||
if (!ggml_is_contiguous_1(other) || !ggml_is_contiguous_1(unary->src[0])) {
|
||||
return false;
|
||||
}
|
||||
// fastmod needs src to tile into dst
|
||||
if (mul->src[0] == unary) {
|
||||
return ggml_can_repeat(other, unary);
|
||||
}
|
||||
return ggml_can_repeat(unary, mul->src[0]);
|
||||
}
|
||||
|
||||
static bool ggml_vk_can_fuse_unary_mul_pair(const struct ggml_cgraph * cgraph, int node_idx) {
|
||||
const enum ggml_op ops[] = { GGML_OP_UNARY, GGML_OP_MUL };
|
||||
const int outputs[] = { node_idx + 1 };
|
||||
return ggml_can_fuse_subgraph(cgraph, node_idx, 2, ops, outputs, 1) &&
|
||||
ggml_vk_can_fuse_unary_mul(cgraph, node_idx, node_idx + 1);
|
||||
}
|
||||
|
||||
static bool ggml_vk_can_fuse(const ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops) {
|
||||
if (ops.size() == 2 && ops.begin()[0] == GGML_OP_UNARY && ops.begin()[1] == GGML_OP_MUL) {
|
||||
return ggml_vk_can_fuse_unary_mul_pair(cgraph, node_idx);
|
||||
}
|
||||
|
||||
if (!ggml_can_fuse(cgraph, node_idx, ops)) {
|
||||
return false;
|
||||
}
|
||||
|
|
@ -17366,6 +17493,7 @@ static bool ggml_vk_can_fuse(const ggml_backend_vk_context * ctx, const struct g
|
|||
}
|
||||
}
|
||||
}
|
||||
|
||||
auto const &mm_add_ok = [&](const ggml_tensor *mul, const ggml_tensor *add) {
|
||||
const ggml_tensor *bias = add->src[0] == mul ? add->src[1] : add->src[0];
|
||||
|
||||
|
|
@ -18234,6 +18362,16 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
|
|||
// they are overwritten, and one workgroup per row. So close enough.
|
||||
op_srcs_fused_elementwise[0] = true;
|
||||
op_srcs_fused_elementwise[1] = true;
|
||||
} else if (ggml_vk_can_fuse(ctx, cgraph, i, { GGML_OP_UNARY, GGML_OP_MUL })) {
|
||||
ctx->num_additional_fused_ops = 1;
|
||||
switch (ggml_get_unary_op(cgraph->nodes[i])) {
|
||||
case GGML_UNARY_OP_GELU: fusion_string = "GELU_MUL"; break;
|
||||
case GGML_UNARY_OP_SIGMOID: fusion_string = "SIGMOID_MUL"; break;
|
||||
case GGML_UNARY_OP_SILU: fusion_string = "SILU_MUL"; break;
|
||||
default: fusion_string = "SOFTPLUS_MUL"; break;
|
||||
}
|
||||
op_srcs_fused_elementwise[0] = true;
|
||||
op_srcs_fused_elementwise[1] = true;
|
||||
} else if (ggml_vk_can_fuse_ssm_conv(ctx, cgraph, i, 2)) {
|
||||
ctx->num_additional_fused_ops = 2;
|
||||
fusion_string = "SSM_CONV_BIAS_SILU";
|
||||
|
|
@ -18532,6 +18670,16 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
|
|||
std::set<ggml_tensor *> used_node_set;
|
||||
|
||||
int first_unused = 0;
|
||||
|
||||
// scheduled or zero-compute nodes in [lo, hi)
|
||||
auto const &empty_or_scheduled_between = [&](int lo, int hi) -> bool {
|
||||
for (int v = lo; v < hi; ++v) {
|
||||
if (!used[v] && !is_empty(graph->nodes[v])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
};
|
||||
while (first_unused < graph->n_nodes) {
|
||||
std::vector<int> current_set;
|
||||
|
||||
|
|
@ -18647,7 +18795,8 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
|
|||
for (int c = first_unused; c < j; ++c) {
|
||||
if (!used[c] &&
|
||||
is_src_of(graph->nodes[j], graph->nodes[c]) &&
|
||||
!(j == c+1 && c == current_set.back() && graph->nodes[c]->op == GGML_OP_RMS_NORM && graph->nodes[j]->op == GGML_OP_MUL) &&
|
||||
!(c == current_set.back() && graph->nodes[c]->op == GGML_OP_RMS_NORM && graph->nodes[j]->op == GGML_OP_MUL && empty_or_scheduled_between(c+1, j)) &&
|
||||
!(c == current_set.back() && graph->nodes[c]->op == GGML_OP_UNARY && graph->nodes[j]->op == GGML_OP_MUL && empty_or_scheduled_between(c+1, j)) &&
|
||||
!(j == c+1 && c == current_set.back() && graph->nodes[c]->op == GGML_OP_MUL_MAT && graph->nodes[j]->op == GGML_OP_ADD) &&
|
||||
!(j == c+1 && c == current_set.back() && graph->nodes[c]->op == GGML_OP_MUL_MAT_ID && graph->nodes[j]->op == GGML_OP_ADD_ID) &&
|
||||
!(j == c+1 && c == current_set.back() && graph->nodes[c]->op == GGML_OP_MUL_MAT_ID && graph->nodes[j]->op == GGML_OP_MUL) &&
|
||||
|
|
@ -18760,6 +18909,27 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
|
|||
}
|
||||
}
|
||||
}
|
||||
// UNARY + MUL: pull the consuming MUL forward
|
||||
if (j > 0 &&
|
||||
graph->nodes[j]->op == GGML_OP_UNARY) {
|
||||
for (int k = j + 1; k < std::min(j + 15, graph->n_nodes); ++k) {
|
||||
ggml_tensor * mul = graph->nodes[k];
|
||||
if (mul->op != GGML_OP_MUL || (mul->src[0] != graph->nodes[j] && mul->src[1] != graph->nodes[j])) {
|
||||
continue;
|
||||
}
|
||||
ggml_tensor * other = (mul->src[0] == graph->nodes[j]) ? mul->src[1] : mul->src[0];
|
||||
// the other src must either be weights or already processed
|
||||
if (!(other->op == GGML_OP_NONE || used_node_set.find(other) != used_node_set.end())) {
|
||||
continue;
|
||||
}
|
||||
if (!ggml_vk_can_fuse_unary_mul(graph, j, k)) {
|
||||
continue;
|
||||
}
|
||||
current_set.push_back(k);
|
||||
used[k] = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
// Second pass grabs view nodes.
|
||||
|
|
|
|||
|
|
@ -1,9 +1,23 @@
|
|||
#version 450
|
||||
|
||||
#include "types.glsl"
|
||||
#if defined(UNARY_MUL_FUSION)
|
||||
#include "generic_binary_head.glsl"
|
||||
#else
|
||||
#include "generic_unary_head.glsl"
|
||||
#endif
|
||||
|
||||
#if defined(UNARY_MUL_FUSION)
|
||||
// OP on src1
|
||||
layout(constant_id = 1) const bool op_on_b = false;
|
||||
#endif
|
||||
|
||||
#if defined(UNARY_MUL_FUSION)
|
||||
layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in;
|
||||
const uint num_threads = 256;
|
||||
#else
|
||||
layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;
|
||||
#endif
|
||||
|
||||
float op_abs(float x) {
|
||||
return abs(x);
|
||||
|
|
@ -123,6 +137,7 @@ float op_gelu_erf(float a) {
|
|||
return 0.5f * a * (1.0f + sign_x * y);
|
||||
}
|
||||
|
||||
#if !defined(UNARY_MUL_FUSION)
|
||||
float op_xielu(float x) {
|
||||
const float alpha_n = p.param1;
|
||||
const float alpha_p = p.param2;
|
||||
|
|
@ -136,6 +151,7 @@ float op_xielu(float x) {
|
|||
const float min_x_eps = min(x, eps);
|
||||
return (op_expm1(min_x_eps) - x) * alpha_n + beta * x;
|
||||
}
|
||||
#endif
|
||||
|
||||
float op_floor(float x) {
|
||||
return floor(x);
|
||||
|
|
@ -155,8 +171,28 @@ float op_trunc(float x) {
|
|||
}
|
||||
|
||||
void main() {
|
||||
const uint idx = get_idx();
|
||||
uint idx = get_idx();
|
||||
|
||||
#if defined(UNARY_MUL_FUSION)
|
||||
// keep total threads at 512
|
||||
[[unroll]] for (uint iter = 0; iter < 2; ++iter) {
|
||||
if (idx >= p.ne) {
|
||||
continue;
|
||||
}
|
||||
uint i00, i01, i02, i03;
|
||||
get_indices(idx, i00, i01, i02, i03);
|
||||
|
||||
if (op_on_b) {
|
||||
data_d[get_doffset() + dst_idx(i00, i01, i02, i03)] =
|
||||
D_TYPE(FLOAT_TYPE(OP(float(data_b[get_boffset() + src1_idx(i00, i01, i02, i03)]))) * FLOAT_TYPE(data_a[get_aoffset() + src0_idx(i00, i01, i02, i03)]));
|
||||
} else {
|
||||
data_d[get_doffset() + dst_idx(i00, i01, i02, i03)] =
|
||||
D_TYPE(FLOAT_TYPE(OP(float(data_a[get_aoffset() + src0_idx(i00, i01, i02, i03)]))) * FLOAT_TYPE(data_b[get_boffset() + src1_idx(i00, i01, i02, i03)]));
|
||||
}
|
||||
|
||||
idx += num_threads;
|
||||
}
|
||||
#else
|
||||
if (idx >= p.ne) {
|
||||
return;
|
||||
}
|
||||
|
|
@ -165,4 +201,5 @@ void main() {
|
|||
const uint d_idx = get_doffset() + dst_idx(idx);
|
||||
|
||||
data_d[d_idx] = D_TYPE(OP(float(data_a[a_idx])));
|
||||
#endif
|
||||
}
|
||||
|
|
|
|||
|
|
@ -992,6 +992,15 @@ void process_shaders() {
|
|||
string_to_spv("softplus_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"OP", "op_softplus"}});
|
||||
string_to_spv("softplus_f32", "unary.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}, {"OP", "op_softplus"}});
|
||||
|
||||
string_to_spv("gelu_mul_f32", "unary.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"OP", "op_gelu"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("gelu_mul_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}, {"OP", "op_gelu"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("sigmoid_mul_f32", "unary.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"OP", "op_sigmoid"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("sigmoid_mul_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}, {"OP", "op_sigmoid"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("silu_mul_f32", "unary.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"OP", "op_silu"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("silu_mul_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}, {"OP", "op_silu"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("softplus_mul_f32","unary.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"OP", "op_softplus"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
string_to_spv("softplus_mul_f16","unary.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}, {"OP", "op_softplus"}, {"UNARY_MUL_FUSION", "1"}});
|
||||
|
||||
string_to_spv("add1_f16_f16", "add1.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}});
|
||||
string_to_spv("add1_f16_f32", "add1.comp", {{"A_TYPE", "float16_t"}, {"B_TYPE", "float"}, {"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}});
|
||||
string_to_spv("add1_f32_f32", "add1.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
|
||||
|
|
|
|||
|
|
@ -9,6 +9,7 @@
|
|||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <cerrno>
|
||||
#include <map>
|
||||
#include <new>
|
||||
#include <stdexcept>
|
||||
|
|
|
|||
|
|
@ -3699,6 +3699,9 @@ llama_context * llama_init_from_model(
|
|||
LLAMA_LOG_ERROR("%s: SPLIT_MODE_TENSOR requires flash_attn to be enabled\n", __func__);
|
||||
return nullptr;
|
||||
}
|
||||
if (model->get_split_state_ud.n_devices == 1) {
|
||||
LLAMA_LOG_WARN("%s: SPLIT_MODE_TENSOR being used for a single device is not recommended\n", __func__);
|
||||
}
|
||||
}
|
||||
|
||||
if ((model->hparams.is_mla() || model->arch == LLM_ARCH_DEEPSEEK4) && params.type_k != params.type_v) {
|
||||
|
|
|
|||
|
|
@ -6,6 +6,7 @@
|
|||
#include "llama-adapter.h"
|
||||
|
||||
#include <cstdint>
|
||||
#include <cstdlib>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
#include <set>
|
||||
|
|
|
|||
|
|
@ -6,6 +6,7 @@
|
|||
|
||||
#include <cstring>
|
||||
#include <climits>
|
||||
#include <cstdlib>
|
||||
#include <stdexcept>
|
||||
#include <cerrno>
|
||||
#include <algorithm>
|
||||
|
|
|
|||
|
|
@ -1576,6 +1576,18 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
|
|||
}
|
||||
}
|
||||
|
||||
// resolve AUTO on systems without mmap support (e.g. iGPUs): fall back to OFF; see #28160
|
||||
if (ml.lazy.mode == LLAMA_LAZY_MODE_AUTO) {
|
||||
for (const auto & dev : devices) {
|
||||
ggml_backend_dev_props props;
|
||||
ggml_backend_dev_get_props(dev.dev, &props);
|
||||
if (!props.caps.mmap_support) {
|
||||
ml.lazy.mode = LLAMA_LAZY_MODE_OFF;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
const char * load_mode_name = params.load_mode == LLAMA_LOAD_MODE_AUTO
|
||||
? llama_load_mode_name(ml.use_mmap ? LLAMA_LOAD_MODE_MMAP : LLAMA_LOAD_MODE_NONE)
|
||||
: llama_load_mode_name(params.load_mode);
|
||||
|
|
|
|||
|
|
@ -14,6 +14,7 @@
|
|||
#include <cmath>
|
||||
#include <cstdarg>
|
||||
#include <cstring>
|
||||
#include <cstdlib>
|
||||
#include <forward_list>
|
||||
#include <limits>
|
||||
#include <map>
|
||||
|
|
|
|||
|
|
@ -1,5 +1,6 @@
|
|||
#include <clocale>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <string>
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
|
|
|
|||
|
|
@ -2311,8 +2311,11 @@ private:
|
|||
|
||||
// evict checkpoints within min-step of a previous checkpoint, unless they were
|
||||
// created by the current task
|
||||
// only when the list is full, otherwise short prompts keep just the oldest checkpoint
|
||||
int64_t last = -1;
|
||||
for (auto it = slot.prompt.checkpoints.begin(); it != slot.prompt.checkpoints.end(); ) {
|
||||
for (auto it = slot.prompt.checkpoints.begin();
|
||||
slot.prompt.checkpoints.size() + 1 >= (size_t) params_base.n_ctx_checkpoints &&
|
||||
it != slot.prompt.checkpoints.end(); ) {
|
||||
if (it->id_task != id_task && last >= 0 && it->n_tokens <= last + params_base.checkpoint_min_step) {
|
||||
SLT_TRC(slot, "erasing context checkpoint too close to an earlier one (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", size = %.3f MiB)\n",
|
||||
it->pos_min, it->pos_max, it->n_tokens, (float) it->size() / 1024 / 1024);
|
||||
|
|
@ -2335,6 +2338,19 @@ private:
|
|||
slot.prompt.checkpoints.erase(slot.prompt.checkpoints.begin());
|
||||
}
|
||||
|
||||
// replace an existing checkpoint at the same n_tokens instead of appending a duplicate
|
||||
{
|
||||
const int64_t n_tokens_new = slot.prompt.n_tokens() - n_tokens_cur;
|
||||
for (auto it = slot.prompt.checkpoints.begin(); it != slot.prompt.checkpoints.end(); ) {
|
||||
if (it->n_tokens == n_tokens_new) {
|
||||
SLT_TRC(slot, "superseding context checkpoint at n_tokens = %" PRId64 "\n", it->n_tokens);
|
||||
it = slot.prompt.checkpoints.erase(it);
|
||||
} else {
|
||||
++it;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
auto & cur = slot.prompt.checkpoints.emplace_back();
|
||||
|
||||
cur.id_task = id_task;
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue