mirror of
https://github.com/zenlm/zen5-engine.git
synced 2026-07-26 22:09:00 +00:00
CUDA support
This commit is contained in:
Codex
@@ -12,8 +12,12 @@ METAL_LDLIBS := $(LDLIBS) -framework Foundation -framework Metal
|
||||
CORE_OBJS = ds4.o ds4_metal.o
|
||||
NATIVE_CORE_OBJS = ds4_native.o
|
||||
else
|
||||
CFLAGS += -DDS4_NO_METAL
|
||||
CORE_OBJS = ds4.o
|
||||
CFLAGS += -D_GNU_SOURCE -fno-finite-math-only
|
||||
CUDA_HOME ?= /usr/local/cuda
|
||||
NVCC ?= $(CUDA_HOME)/bin/nvcc
|
||||
NVCCFLAGS ?= -O3 --use_fast_math -Xcompiler -mcpu=native -Xcompiler -pthread
|
||||
CUDA_LDLIBS ?= -lm -Xcompiler -pthread -L$(CUDA_HOME)/targets/sbsa-linux/lib -L$(CUDA_HOME)/lib64 -lcudart -lcublas
|
||||
CORE_OBJS = ds4.o ds4_cuda.o
|
||||
NATIVE_CORE_OBJS = ds4_native.o
|
||||
METAL_LDLIBS := $(LDLIBS)
|
||||
endif
|
||||
@@ -33,10 +37,10 @@ ds4_native: ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS)
|
||||
$(CC) $(CFLAGS) -o $@ ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS) $(NATIVE_LDLIBS)
|
||||
else
|
||||
ds4: ds4_cli.o linenoise.o $(CORE_OBJS)
|
||||
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
|
||||
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
|
||||
|
||||
ds4-server: ds4_server.o rax.o $(CORE_OBJS)
|
||||
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
|
||||
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
|
||||
|
||||
ds4_native: ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS)
|
||||
$(CC) $(CFLAGS) -o $@ ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS) $(LDLIBS)
|
||||
@@ -69,8 +73,15 @@ ds4_cli_native.o: ds4_cli.c ds4.h linenoise.h
|
||||
ds4_metal.o: ds4_metal.m ds4_metal.h $(METAL_SRCS)
|
||||
$(CC) $(OBJCFLAGS) -c -o $@ ds4_metal.m
|
||||
|
||||
ds4_cuda.o: ds4_cuda.cu ds4_metal.h ds4_iq2_tables_cuda.inc
|
||||
$(NVCC) $(NVCCFLAGS) -c -o $@ ds4_cuda.cu
|
||||
|
||||
ds4_test: ds4_test.o rax.o $(CORE_OBJS)
|
||||
ifeq ($(UNAME_S),Darwin)
|
||||
$(CC) $(CFLAGS) -o $@ ds4_test.o rax.o $(CORE_OBJS) $(METAL_LDLIBS)
|
||||
else
|
||||
$(NVCC) $(NVCCFLAGS) -o $@ ds4_test.o rax.o $(CORE_OBJS) $(CUDA_LDLIBS)
|
||||
endif
|
||||
|
||||
test: ds4_test
|
||||
./ds4_test
|
||||
|
||||
@@ -70,6 +70,10 @@ static const char DS4_REASONING_EFFORT_MAX_PREFIX[] =
|
||||
* asks for a reasoning budget the allocated context is not meant to hold. */
|
||||
#define DS4_THINK_MAX_MIN_CONTEXT 393216u
|
||||
|
||||
static bool ds4_backend_uses_graph(ds4_backend backend) {
|
||||
return backend == DS4_BACKEND_METAL || backend == DS4_BACKEND_CUDA;
|
||||
}
|
||||
|
||||
/* =========================================================================
|
||||
* Fixed DeepSeek V4 Flash Shape.
|
||||
* =========================================================================
|
||||
@@ -1329,6 +1333,132 @@ static ds4_tensor *model_find_tensor(const ds4_model *m, const char *name) {
|
||||
return NULL;
|
||||
}
|
||||
|
||||
#ifndef DS4_NO_METAL
|
||||
#ifndef __APPLE__
|
||||
typedef struct {
|
||||
uint64_t off;
|
||||
uint64_t end;
|
||||
} accelerator_tensor_span;
|
||||
|
||||
static int accelerator_tensor_span_cmp(const void *a, const void *b) {
|
||||
const accelerator_tensor_span *sa = a;
|
||||
const accelerator_tensor_span *sb = b;
|
||||
if (sa->off < sb->off) return -1;
|
||||
if (sa->off > sb->off) return 1;
|
||||
if (sa->end < sb->end) return -1;
|
||||
if (sa->end > sb->end) return 1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
static uint64_t accelerator_cuda_preload_span_bytes(void) {
|
||||
uint64_t mb = 1024;
|
||||
const char *env = getenv("DS4_CUDA_WEIGHT_PRELOAD_SPAN_MB");
|
||||
if (env && env[0]) {
|
||||
char *end = NULL;
|
||||
unsigned long long v = strtoull(env, &end, 10);
|
||||
if (end != env && v > 0) mb = (uint64_t)v;
|
||||
}
|
||||
if (mb < 64) mb = 64;
|
||||
if (mb > 4096) mb = 4096;
|
||||
return mb * 1048576ull;
|
||||
}
|
||||
|
||||
static bool accelerator_cache_model_tensor_spans(const ds4_model *m, uint64_t *cached_out) {
|
||||
accelerator_tensor_span *spans = xmalloc((size_t)m->n_tensors * sizeof(spans[0]));
|
||||
uint64_t nspan = 0;
|
||||
for (uint64_t i = 0; i < m->n_tensors; i++) {
|
||||
const ds4_tensor *t = &m->tensors[i];
|
||||
if (t->bytes == 0) continue;
|
||||
if (t->abs_offset > m->size || t->bytes > m->size - t->abs_offset) {
|
||||
free(spans);
|
||||
return false;
|
||||
}
|
||||
spans[nspan++] = (accelerator_tensor_span){
|
||||
.off = t->abs_offset,
|
||||
.end = t->abs_offset + t->bytes,
|
||||
};
|
||||
}
|
||||
qsort(spans, (size_t)nspan, sizeof(spans[0]), accelerator_tensor_span_cmp);
|
||||
|
||||
const uint64_t max_span = accelerator_cuda_preload_span_bytes();
|
||||
uint64_t cached = 0;
|
||||
uint64_t merged = 0;
|
||||
for (uint64_t i = 0; i < nspan;) {
|
||||
uint64_t off = spans[i].off;
|
||||
uint64_t end = spans[i].end;
|
||||
i++;
|
||||
while (i < nspan && spans[i].off <= end + 65536u && spans[i].end - off <= max_span) {
|
||||
if (spans[i].end > end) end = spans[i].end;
|
||||
i++;
|
||||
}
|
||||
while (off < end) {
|
||||
uint64_t chunk_end = end;
|
||||
if (chunk_end - off > max_span) chunk_end = off + max_span;
|
||||
char label[96];
|
||||
snprintf(label, sizeof(label), "tensor-span:%" PRIu64, merged);
|
||||
if (ds4_metal_cache_model_range(m->map, m->size, off, chunk_end - off, label) == 0) {
|
||||
fprintf(stderr,
|
||||
"ds4: accelerator failed to cache model tensor span %" PRIu64
|
||||
" at offset %" PRIu64 "\n",
|
||||
merged, off);
|
||||
free(spans);
|
||||
return false;
|
||||
}
|
||||
cached += chunk_end - off;
|
||||
merged++;
|
||||
off = chunk_end;
|
||||
}
|
||||
}
|
||||
free(spans);
|
||||
if (cached_out) *cached_out = cached;
|
||||
return true;
|
||||
}
|
||||
|
||||
static bool accelerator_cache_model_tensors(ds4_backend backend, const ds4_model *m) {
|
||||
if (backend != DS4_BACKEND_CUDA) return true;
|
||||
if (!m || !m->map || m->size == 0) return false;
|
||||
if (getenv("DS4_CUDA_DIRECT_MODEL") != NULL) {
|
||||
return true;
|
||||
}
|
||||
|
||||
const double t0 = now_sec();
|
||||
uint64_t cached = 0;
|
||||
if (!accelerator_cache_model_tensor_spans(m, &cached)) return false;
|
||||
if (getenv("DS4_CUDA_Q8_F16_PRELOAD") != NULL ||
|
||||
getenv("DS4_CUDA_Q8_F32_PRELOAD") != NULL) {
|
||||
for (uint64_t i = 0; i < m->n_tensors; i++) {
|
||||
const ds4_tensor *t = &m->tensors[i];
|
||||
if (t->bytes == 0) continue;
|
||||
if (t->abs_offset > m->size || t->bytes > m->size - t->abs_offset) return false;
|
||||
char label[128];
|
||||
snprintf(label, sizeof(label), "tensor:%.*s", (int)t->name.len, t->name.ptr);
|
||||
if (t->type == DS4_TENSOR_Q8_0 && t->ndim == 2 &&
|
||||
ds4_metal_cache_q8_f16_range(m->map, m->size, t->abs_offset, t->bytes, t->dim[0], t->dim[1], label) == 0) {
|
||||
fprintf(stderr, "ds4: accelerator failed to cache dequantized Q8 tensor %.*s\n",
|
||||
(int)t->name.len, t->name.ptr);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (cached != 0) {
|
||||
const double t1 = now_sec();
|
||||
if (ds4_log_is_tty(stderr)) fputc('\n', stderr);
|
||||
fprintf(stderr,
|
||||
"ds4: CUDA startup model cache prepared %.2f GiB of tensor spans in %.3fs\n",
|
||||
(double)cached / 1073741824.0,
|
||||
t1 - t0);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
#else
|
||||
static bool accelerator_cache_model_tensors(ds4_backend backend, const ds4_model *m) {
|
||||
(void)backend;
|
||||
(void)m;
|
||||
return true;
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
|
||||
/* Return the in-place tensor payload inside the mapped GGUF. */
|
||||
static const void *tensor_data(const ds4_model *m, const ds4_tensor *t) {
|
||||
return m->map + t->abs_offset;
|
||||
@@ -13298,7 +13428,7 @@ ds4_context_memory ds4_context_memory_estimate(ds4_backend backend, int ctx_size
|
||||
ds4_context_memory m = {0};
|
||||
uint32_t ctx = ctx_size > 0 ? (uint32_t)ctx_size : 1u;
|
||||
|
||||
if (backend == DS4_BACKEND_METAL) {
|
||||
if (ds4_backend_uses_graph(backend)) {
|
||||
m.prefill_cap = metal_graph_prefill_cap_for_prompt((int)ctx);
|
||||
m.raw_cap = metal_graph_raw_cap_for_context((int)ctx, m.prefill_cap);
|
||||
|
||||
@@ -14838,7 +14968,12 @@ ds4_context_memory ds4_context_memory_estimate(ds4_backend backend, int ctx_size
|
||||
*/
|
||||
|
||||
const char *ds4_backend_name(ds4_backend backend) {
|
||||
return backend == DS4_BACKEND_METAL ? "metal" : "cpu";
|
||||
switch (backend) {
|
||||
case DS4_BACKEND_METAL: return "metal";
|
||||
case DS4_BACKEND_CUDA: return "cuda";
|
||||
case DS4_BACKEND_CPU: return "cpu";
|
||||
}
|
||||
return "unknown";
|
||||
}
|
||||
|
||||
bool ds4_think_mode_enabled(ds4_think_mode mode) {
|
||||
@@ -15644,10 +15779,11 @@ int ds4_engine_generate_argmax(
|
||||
const ds4_vocab *vocab = &e->vocab;
|
||||
const ds4_weights *weights = &e->weights;
|
||||
|
||||
if (e->backend == DS4_BACKEND_METAL) {
|
||||
if (ds4_backend_uses_graph(e->backend)) {
|
||||
#ifndef DS4_NO_METAL
|
||||
if (!e->metal_ready) {
|
||||
fprintf(stderr, "ds4: Metal generation requested but Metal is unavailable\n");
|
||||
fprintf(stderr, "ds4: %s generation requested but the graph backend is unavailable\n",
|
||||
ds4_backend_name(e->backend));
|
||||
return 1;
|
||||
}
|
||||
return generate_metal_graph_raw_swa(model, vocab, weights, prompt,
|
||||
@@ -15658,7 +15794,8 @@ int ds4_engine_generate_argmax(
|
||||
emit, done, emit_ud,
|
||||
progress, progress_ud);
|
||||
#else
|
||||
fprintf(stderr, "ds4: Metal generation requested but this build has no Metal support\n");
|
||||
fprintf(stderr, "ds4: %s generation requested but this build has no graph backend support\n",
|
||||
ds4_backend_name(e->backend));
|
||||
return 1;
|
||||
#endif
|
||||
}
|
||||
@@ -15877,15 +16014,14 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) {
|
||||
if (opt->n_threads > 0) g_requested_threads = (uint32_t)opt->n_threads;
|
||||
ds4_acquire_instance_lock();
|
||||
|
||||
model_open(&e->model, opt->model_path,
|
||||
opt->backend == DS4_BACKEND_METAL, true);
|
||||
const bool graph_backend = ds4_backend_uses_graph(opt->backend);
|
||||
model_open(&e->model, opt->model_path, graph_backend, true);
|
||||
if (opt->warm_weights) model_warm_weights(&e->model);
|
||||
vocab_load(&e->vocab, &e->model);
|
||||
config_validate_model(&e->model);
|
||||
weights_bind(&e->weights, &e->model);
|
||||
if (opt->mtp_path && opt->mtp_path[0]) {
|
||||
model_open(&e->mtp_model, opt->mtp_path,
|
||||
opt->backend == DS4_BACKEND_METAL, true);
|
||||
model_open(&e->mtp_model, opt->mtp_path, graph_backend, true);
|
||||
mtp_weights_bind(&e->mtp_weights, &e->mtp_model);
|
||||
e->mtp_ready = true;
|
||||
fprintf(stderr, "ds4: MTP support model loaded: %s (draft=%d)\n",
|
||||
@@ -15894,23 +16030,42 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) {
|
||||
}
|
||||
|
||||
#ifndef DS4_NO_METAL
|
||||
if (e->backend == DS4_BACKEND_CUDA) {
|
||||
#ifdef __APPLE__
|
||||
fprintf(stderr, "ds4: CUDA backend requested but this build is linked with Metal, not CUDA\n");
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
#endif
|
||||
}
|
||||
if (e->backend == DS4_BACKEND_METAL) {
|
||||
#ifndef __APPLE__
|
||||
fprintf(stderr, "ds4: Metal backend requested but this build is linked with CUDA, not Metal\n");
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
#endif
|
||||
}
|
||||
if (graph_backend) {
|
||||
e->metal_ready = ds4_metal_init() != 0;
|
||||
if (!e->metal_ready) {
|
||||
fprintf(stderr, "ds4: Metal backend unavailable; aborting startup\n");
|
||||
fprintf(stderr, "ds4: %s backend unavailable; aborting startup\n",
|
||||
ds4_backend_name(e->backend));
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
}
|
||||
ds4_metal_set_quality(e->quality);
|
||||
(void)ds4_metal_set_model_fd(e->model.fd);
|
||||
if (!ds4_metal_set_model_map_range(e->model.map,
|
||||
e->model.size,
|
||||
e->model.tensor_data_pos,
|
||||
e->model.size - e->model.tensor_data_pos))
|
||||
{
|
||||
fprintf(stderr,
|
||||
"ds4: Metal failed to map model views; aborting startup. "
|
||||
"This is commonly caused by insufficient memory or Metal VM budget.\n");
|
||||
"ds4: %s failed to map model views; aborting startup. "
|
||||
"This is commonly caused by insufficient memory or accelerator VM budget.\n",
|
||||
ds4_backend_name(e->backend));
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
@@ -15922,17 +16077,27 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) {
|
||||
e->mtp_model.size - e->mtp_model.tensor_data_pos))
|
||||
{
|
||||
fprintf(stderr,
|
||||
"ds4: Metal failed to map MTP model views; aborting startup. "
|
||||
"This is commonly caused by insufficient memory or Metal VM budget.\n");
|
||||
"ds4: %s failed to map MTP model views; aborting startup. "
|
||||
"This is commonly caused by insufficient memory or accelerator VM budget.\n",
|
||||
ds4_backend_name(e->backend));
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
}
|
||||
fprintf(stderr, "ds4: Metal backend initialized for graph diagnostics\n");
|
||||
if (!e->mtp_ready && !accelerator_cache_model_tensors(e->backend, &e->model)) {
|
||||
fprintf(stderr, "ds4: %s failed to prepare startup model cache\n",
|
||||
ds4_backend_name(e->backend));
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
}
|
||||
fprintf(stderr, "ds4: %s backend initialized for graph diagnostics\n",
|
||||
ds4_backend_name(e->backend));
|
||||
}
|
||||
#else
|
||||
if (e->backend == DS4_BACKEND_METAL) {
|
||||
fprintf(stderr, "ds4: Metal backend requested but this build has no Metal support; aborting startup\n");
|
||||
if (graph_backend) {
|
||||
fprintf(stderr, "ds4: %s backend requested but this build has no graph backend support; aborting startup\n",
|
||||
ds4_backend_name(e->backend));
|
||||
ds4_engine_close(e);
|
||||
*out = NULL;
|
||||
return 1;
|
||||
@@ -15969,7 +16134,7 @@ int ds4_session_create(ds4_session **out, ds4_engine *e, int ctx_size) {
|
||||
(void)ctx_size;
|
||||
return 1;
|
||||
#else
|
||||
if (e->backend != DS4_BACKEND_METAL || !e->metal_ready) return 1;
|
||||
if (!ds4_backend_uses_graph(e->backend) || !e->metal_ready) return 1;
|
||||
|
||||
ds4_session *s = xcalloc(1, sizeof(*s));
|
||||
s->engine = e;
|
||||
|
||||
@@ -16,6 +16,7 @@
|
||||
|
||||
typedef enum {
|
||||
DS4_BACKEND_METAL,
|
||||
DS4_BACKEND_CUDA,
|
||||
DS4_BACKEND_CPU,
|
||||
} ds4_backend;
|
||||
|
||||
|
||||
@@ -88,11 +88,13 @@ static void usage(FILE *fp) {
|
||||
" -c, --ctx N\n"
|
||||
" Context size allocated for the session. Default: 32768\n"
|
||||
" --metal\n"
|
||||
" Use the Metal graph backend. This is the normal fast path and the default.\n"
|
||||
" Use the Metal graph backend. This is the normal fast path on macOS.\n"
|
||||
" --cuda\n"
|
||||
" Use the CUDA graph backend. This is the normal fast path on CUDA builds.\n"
|
||||
" --cpu\n"
|
||||
" Use the CPU reference/debug backend. Not recommended for normal inference.\n"
|
||||
" --backend NAME\n"
|
||||
" Select backend explicitly: metal or cpu. Default: metal\n"
|
||||
" Select backend explicitly: metal, cuda, or cpu.\n"
|
||||
" -t, --threads N\n"
|
||||
" CPU helper threads for host-side or reference work.\n"
|
||||
" --quality\n"
|
||||
@@ -168,7 +170,7 @@ static void usage(FILE *fp) {
|
||||
" ./ds4 --think-max --prompt-file prompt.txt --ctx 393216\n"
|
||||
"\n"
|
||||
"Notes:\n"
|
||||
" The CLI keeps KV cache state across interactive turns on the Metal backend.\n"
|
||||
" The CLI keeps KV cache state across interactive turns on graph backends.\n"
|
||||
" Long added input is processed with batched prefill; short continuations use decode.\n"
|
||||
" Startup prints the extra context-buffer memory for the selected context size.\n"
|
||||
"\n"
|
||||
@@ -208,12 +210,23 @@ static float parse_float_range(const char *s, const char *opt, float min, float
|
||||
|
||||
static ds4_backend parse_backend(const char *s) {
|
||||
if (!strcmp(s, "metal")) return DS4_BACKEND_METAL;
|
||||
if (!strcmp(s, "cuda")) return DS4_BACKEND_CUDA;
|
||||
if (!strcmp(s, "cpu")) return DS4_BACKEND_CPU;
|
||||
fprintf(stderr, "ds4: invalid backend: %s\n", s);
|
||||
fprintf(stderr, "ds4: valid backends are: metal, cpu\n");
|
||||
fprintf(stderr, "ds4: valid backends are: metal, cuda, cpu\n");
|
||||
exit(2);
|
||||
}
|
||||
|
||||
static ds4_backend default_backend(void) {
|
||||
#ifdef DS4_NO_METAL
|
||||
return DS4_BACKEND_CPU;
|
||||
#elif defined(__APPLE__)
|
||||
return DS4_BACKEND_METAL;
|
||||
#else
|
||||
return DS4_BACKEND_CUDA;
|
||||
#endif
|
||||
}
|
||||
|
||||
static void log_context_memory(ds4_backend backend, int ctx_size) {
|
||||
ds4_context_memory m = ds4_context_memory_estimate(backend, ctx_size);
|
||||
fprintf(stderr,
|
||||
@@ -433,7 +446,7 @@ static void build_prompt(ds4_engine *engine, const cli_generation_options *gen,
|
||||
static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, const ds4_tokens *prompt) {
|
||||
ds4_session *session = NULL;
|
||||
if (ds4_session_create(&session, engine, cfg->gen.ctx_size) != 0) {
|
||||
fprintf(stderr, "ds4: sampled CLI generation requires the Metal session backend\n");
|
||||
fprintf(stderr, "ds4: sampled CLI generation requires a graph session backend\n");
|
||||
return 1;
|
||||
}
|
||||
|
||||
@@ -606,7 +619,7 @@ static void json_write_token(FILE *fp, ds4_engine *engine, int token) {
|
||||
static int run_logprob_dump(ds4_engine *engine, const cli_config *cfg, const ds4_tokens *prompt) {
|
||||
ds4_session *session = NULL;
|
||||
if (ds4_session_create(&session, engine, cfg->gen.ctx_size) != 0) {
|
||||
fprintf(stderr, "ds4: --dump-logprobs requires the Metal session backend\n");
|
||||
fprintf(stderr, "ds4: --dump-logprobs requires a graph session backend\n");
|
||||
return 1;
|
||||
}
|
||||
|
||||
@@ -839,7 +852,7 @@ static void repl_chat_apply_max_prefix(ds4_engine *engine, repl_chat *chat, bool
|
||||
static int repl_chat_create_session(ds4_engine *engine, repl_chat *chat, int ctx_size) {
|
||||
ds4_session *session = NULL;
|
||||
if (ds4_session_create(&session, engine, ctx_size) != 0) {
|
||||
fprintf(stderr, "ds4: interactive chat KV cache requires the Metal backend\n");
|
||||
fprintf(stderr, "ds4: interactive chat KV cache requires a graph backend\n");
|
||||
return 1;
|
||||
}
|
||||
if (chat->session) ds4_session_free(chat->session);
|
||||
@@ -1161,7 +1174,7 @@ static cli_config parse_options(int argc, char **argv) {
|
||||
cli_config c = {
|
||||
.engine = {
|
||||
.model_path = "ds4flash.gguf",
|
||||
.backend = DS4_BACKEND_METAL,
|
||||
.backend = default_backend(),
|
||||
.mtp_draft_tokens = 1,
|
||||
.mtp_margin = 3.0f,
|
||||
},
|
||||
@@ -1234,6 +1247,8 @@ static cli_config parse_options(int argc, char **argv) {
|
||||
c.engine.backend = DS4_BACKEND_CPU;
|
||||
} else if (!strcmp(arg, "--metal")) {
|
||||
c.engine.backend = DS4_BACKEND_METAL;
|
||||
} else if (!strcmp(arg, "--cuda")) {
|
||||
c.engine.backend = DS4_BACKEND_CUDA;
|
||||
} else if (!strcmp(arg, "--dump-tokens")) {
|
||||
c.gen.dump_tokens = true;
|
||||
} else if (!strcmp(arg, "--dump-logprobs")) {
|
||||
|
||||
+9666
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,77 @@
|
||||
__device__ __constant__ uint8_t cuda_ksigns_iq2xs[128] = {
|
||||
0, 129, 130, 3, 132, 5, 6, 135, 136, 9, 10, 139, 12, 141, 142, 15,
|
||||
144, 17, 18, 147, 20, 149, 150, 23, 24, 153, 154, 27, 156, 29, 30, 159,
|
||||
160, 33, 34, 163, 36, 165, 166, 39, 40, 169, 170, 43, 172, 45, 46, 175,
|
||||
48, 177, 178, 51, 180, 53, 54, 183, 184, 57, 58, 187, 60, 189, 190, 63,
|
||||
192, 65, 66, 195, 68, 197, 198, 71, 72, 201, 202, 75, 204, 77, 78, 207,
|
||||
80, 209, 210, 83, 212, 85, 86, 215, 216, 89, 90, 219, 92, 221, 222, 95,
|
||||
96, 225, 226, 99, 228, 101, 102, 231, 232, 105, 106, 235, 108, 237, 238, 111,
|
||||
240, 113, 114, 243, 116, 245, 246, 119, 120, 249, 250, 123, 252, 125, 126, 255,
|
||||
};
|
||||
|
||||
__device__ __constant__ uint64_t cuda_iq2xxs_grid[256] = {
|
||||
0x0808080808080808, 0x080808080808082b, 0x0808080808081919, 0x0808080808082b08,
|
||||
0x0808080808082b2b, 0x0808080808190819, 0x0808080808191908, 0x08080808082b0808,
|
||||
0x08080808082b082b, 0x08080808082b2b08, 0x08080808082b2b2b, 0x0808080819080819,
|
||||
0x0808080819081908, 0x0808080819190808, 0x0808080819192b08, 0x08080808192b0819,
|
||||
0x08080808192b1908, 0x080808082b080808, 0x080808082b08082b, 0x080808082b082b2b,
|
||||
0x080808082b2b082b, 0x0808081908080819, 0x0808081908081908, 0x0808081908190808,
|
||||
0x0808081908191919, 0x0808081919080808, 0x080808192b081908, 0x080808192b192b08,
|
||||
0x0808082b08080808, 0x0808082b0808082b, 0x0808082b082b082b, 0x0808082b2b08082b,
|
||||
0x0808190808080819, 0x0808190808081908, 0x0808190808190808, 0x08081908082b0819,
|
||||
0x08081908082b1908, 0x0808190819080808, 0x080819081908082b, 0x0808190819082b08,
|
||||
0x08081908192b0808, 0x080819082b080819, 0x080819082b081908, 0x080819082b190808,
|
||||
0x080819082b2b1908, 0x0808191908080808, 0x080819190808082b, 0x0808191908082b08,
|
||||
0x08081919082b0808, 0x080819191908192b, 0x08081919192b2b19, 0x080819192b080808,
|
||||
0x080819192b190819, 0x0808192b08082b19, 0x0808192b08190808, 0x0808192b19080808,
|
||||
0x0808192b2b081908, 0x0808192b2b2b1908, 0x08082b0808080808, 0x08082b0808081919,
|
||||
0x08082b0808082b08, 0x08082b0808191908, 0x08082b08082b2b08, 0x08082b0819080819,
|
||||
0x08082b0819081908, 0x08082b0819190808, 0x08082b081919082b, 0x08082b082b082b08,
|
||||
0x08082b1908081908, 0x08082b1919080808, 0x08082b2b0808082b, 0x08082b2b08191908,
|
||||
0x0819080808080819, 0x0819080808081908, 0x0819080808190808, 0x08190808082b0819,
|
||||
0x0819080819080808, 0x08190808192b0808, 0x081908082b081908, 0x081908082b190808,
|
||||
0x081908082b191919, 0x0819081908080808, 0x0819081908082b08, 0x08190819082b0808,
|
||||
0x0819081919190808, 0x0819081919192b2b, 0x081908192b080808, 0x0819082b082b1908,
|
||||
0x0819082b19081919, 0x0819190808080808, 0x0819190808082b08, 0x08191908082b0808,
|
||||
0x08191908082b1919, 0x0819190819082b19, 0x081919082b080808, 0x0819191908192b08,
|
||||
0x08191919192b082b, 0x0819192b08080808, 0x0819192b0819192b, 0x08192b0808080819,
|
||||
0x08192b0808081908, 0x08192b0808190808, 0x08192b0819080808, 0x08192b082b080819,
|
||||
0x08192b1908080808, 0x08192b1908081919, 0x08192b192b2b0808, 0x08192b2b19190819,
|
||||
0x082b080808080808, 0x082b08080808082b, 0x082b080808082b2b, 0x082b080819081908,
|
||||
0x082b0808192b0819, 0x082b08082b080808, 0x082b08082b08082b, 0x082b0819082b2b19,
|
||||
0x082b081919082b08, 0x082b082b08080808, 0x082b082b0808082b, 0x082b190808080819,
|
||||
0x082b190808081908, 0x082b190808190808, 0x082b190819080808, 0x082b19081919192b,
|
||||
0x082b191908080808, 0x082b191919080819, 0x082b1919192b1908, 0x082b192b2b190808,
|
||||
0x082b2b0808082b08, 0x082b2b08082b0808, 0x082b2b082b191908, 0x082b2b2b19081908,
|
||||
0x1908080808080819, 0x1908080808081908, 0x1908080808190808, 0x1908080808192b08,
|
||||
0x19080808082b0819, 0x19080808082b1908, 0x1908080819080808, 0x1908080819082b08,
|
||||
0x190808081919192b, 0x19080808192b0808, 0x190808082b080819, 0x190808082b081908,
|
||||
0x190808082b190808, 0x1908081908080808, 0x19080819082b0808, 0x19080819192b0819,
|
||||
0x190808192b080808, 0x190808192b081919, 0x1908082b08080819, 0x1908082b08190808,
|
||||
0x1908082b19082b08, 0x1908082b1919192b, 0x1908082b192b2b08, 0x1908190808080808,
|
||||
0x1908190808082b08, 0x19081908082b0808, 0x190819082b080808, 0x190819082b192b19,
|
||||
0x190819190819082b, 0x19081919082b1908, 0x1908192b08080808, 0x19082b0808080819,
|
||||
0x19082b0808081908, 0x19082b0808190808, 0x19082b0819080808, 0x19082b0819081919,
|
||||
0x19082b1908080808, 0x19082b1919192b08, 0x19082b19192b0819, 0x19082b192b08082b,
|
||||
0x19082b2b19081919, 0x19082b2b2b190808, 0x1919080808080808, 0x1919080808082b08,
|
||||
0x1919080808190819, 0x1919080808192b19, 0x19190808082b0808, 0x191908082b080808,
|
||||
0x191908082b082b08, 0x1919081908081908, 0x191908191908082b, 0x191908192b2b1908,
|
||||
0x1919082b2b190819, 0x191919082b190808, 0x191919082b19082b, 0x1919191908082b2b,
|
||||
0x1919192b08080819, 0x1919192b19191908, 0x19192b0808080808, 0x19192b0808190819,
|
||||
0x19192b0808192b19, 0x19192b08192b1908, 0x19192b1919080808, 0x19192b2b08082b08,
|
||||
0x192b080808081908, 0x192b080808190808, 0x192b080819080808, 0x192b0808192b2b08,
|
||||
0x192b081908080808, 0x192b081919191919, 0x192b082b08192b08, 0x192b082b192b0808,
|
||||
0x192b190808080808, 0x192b190808081919, 0x192b191908190808, 0x192b19190819082b,
|
||||
0x192b19192b081908, 0x192b2b081908082b, 0x2b08080808080808, 0x2b0808080808082b,
|
||||
0x2b08080808082b2b, 0x2b08080819080819, 0x2b0808082b08082b, 0x2b08081908081908,
|
||||
0x2b08081908192b08, 0x2b08081919080808, 0x2b08082b08190819, 0x2b08190808080819,
|
||||
0x2b08190808081908, 0x2b08190808190808, 0x2b08190808191919, 0x2b08190819080808,
|
||||
0x2b081908192b0808, 0x2b08191908080808, 0x2b0819191908192b, 0x2b0819192b191908,
|
||||
0x2b08192b08082b19, 0x2b08192b19080808, 0x2b08192b192b0808, 0x2b082b080808082b,
|
||||
0x2b082b1908081908, 0x2b082b2b08190819, 0x2b19080808081908, 0x2b19080808190808,
|
||||
0x2b190808082b1908, 0x2b19080819080808, 0x2b1908082b2b0819, 0x2b1908190819192b,
|
||||
0x2b1908192b080808, 0x2b19082b19081919, 0x2b19190808080808, 0x2b191908082b082b,
|
||||
0x2b19190819081908, 0x2b19191919190819, 0x2b192b082b080819, 0x2b192b19082b0808,
|
||||
0x2b2b08080808082b, 0x2b2b080819190808, 0x2b2b08082b081919, 0x2b2b081908082b19,
|
||||
0x2b2b082b08080808, 0x2b2b190808192b08, 0x2b2b2b0819190808, 0x2b2b2b1908081908,
|
||||
};
|
||||
@@ -35,7 +35,10 @@ int ds4_metal_end_commands(void);
|
||||
int ds4_metal_synchronize(void);
|
||||
|
||||
int ds4_metal_set_model_map(const void *model_map, uint64_t model_size);
|
||||
int ds4_metal_set_model_fd(int fd);
|
||||
int ds4_metal_set_model_map_range(const void *model_map, uint64_t model_size, uint64_t map_offset, uint64_t map_size);
|
||||
int ds4_metal_cache_model_range(const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes, const char *label);
|
||||
int ds4_metal_cache_q8_f16_range(const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes, uint64_t in_dim, uint64_t out_dim, const char *label);
|
||||
void ds4_metal_set_quality(bool quality);
|
||||
void ds4_metal_print_memory_report(const char *label);
|
||||
|
||||
|
||||
@@ -4392,6 +4392,11 @@ int ds4_metal_set_model_map(const void *model_map, uint64_t model_size) {
|
||||
return ds4_metal_set_model_map_range(model_map, model_size, 0, model_size);
|
||||
}
|
||||
|
||||
int ds4_metal_set_model_fd(int fd) {
|
||||
(void)fd;
|
||||
return 1;
|
||||
}
|
||||
|
||||
static id<MTLBuffer> ds4_metal_wrap_model_range(
|
||||
const void *model_map,
|
||||
uint64_t model_size,
|
||||
|
||||
+30
-4
@@ -7893,6 +7893,8 @@ static void usage(FILE *fp) {
|
||||
" Apply steering after attention outputs. Default: 0\n"
|
||||
" --warm-weights\n"
|
||||
" Touch mapped tensor pages before serving. Slower startup, fewer first-use stalls.\n"
|
||||
" --metal | --cuda | --backend NAME\n"
|
||||
" Select graph backend explicitly. Defaults to Metal on macOS and CUDA on CUDA builds.\n"
|
||||
"\n"
|
||||
"HTTP API:\n"
|
||||
" --host HOST\n"
|
||||
@@ -7942,7 +7944,7 @@ static void usage(FILE *fp) {
|
||||
" ./ds4-server --ctx 100000 --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 8192\n"
|
||||
"\n"
|
||||
"Notes:\n"
|
||||
" The server is Metal-only. Use /v1/chat/completions, /v1/completions, or /v1/messages.\n"
|
||||
" Use /v1/chat/completions, /v1/completions, or /v1/messages.\n"
|
||||
" Larger --ctx values allocate more KV memory at startup; the startup log prints the estimate.\n"
|
||||
" Disk KV caching is best for agents that resend long prompts with stable prefixes.\n"
|
||||
"\n"
|
||||
@@ -7950,11 +7952,29 @@ static void usage(FILE *fp) {
|
||||
" Show this help.\n");
|
||||
}
|
||||
|
||||
static ds4_backend parse_backend_arg(const char *s, const char *arg) {
|
||||
if (!strcmp(s, "metal")) return DS4_BACKEND_METAL;
|
||||
if (!strcmp(s, "cuda")) return DS4_BACKEND_CUDA;
|
||||
server_log(DS4_LOG_DEFAULT, "ds4-server: invalid %s value: %s", arg, s);
|
||||
server_log(DS4_LOG_DEFAULT, "ds4-server: valid server backends are: metal, cuda");
|
||||
exit(2);
|
||||
}
|
||||
|
||||
static ds4_backend default_server_backend(void) {
|
||||
#ifdef DS4_NO_METAL
|
||||
return DS4_BACKEND_CPU;
|
||||
#elif defined(__APPLE__)
|
||||
return DS4_BACKEND_METAL;
|
||||
#else
|
||||
return DS4_BACKEND_CUDA;
|
||||
#endif
|
||||
}
|
||||
|
||||
static server_config parse_options(int argc, char **argv) {
|
||||
server_config c = {
|
||||
.engine = {
|
||||
.model_path = "ds4flash.gguf",
|
||||
.backend = DS4_BACKEND_METAL,
|
||||
.backend = default_server_backend(),
|
||||
.mtp_draft_tokens = 1,
|
||||
.mtp_margin = 3.0f,
|
||||
},
|
||||
@@ -8024,8 +8044,14 @@ static server_config parse_options(int argc, char **argv) {
|
||||
directional_steering_scale_set = true;
|
||||
} else if (!strcmp(arg, "--warm-weights")) {
|
||||
c.engine.warm_weights = true;
|
||||
} else if (!strcmp(arg, "--cpu") || !strcmp(arg, "--backend")) {
|
||||
server_log(DS4_LOG_DEFAULT, "ds4-server: server mode is Metal-only");
|
||||
} else if (!strcmp(arg, "--metal")) {
|
||||
c.engine.backend = DS4_BACKEND_METAL;
|
||||
} else if (!strcmp(arg, "--cuda")) {
|
||||
c.engine.backend = DS4_BACKEND_CUDA;
|
||||
} else if (!strcmp(arg, "--backend")) {
|
||||
c.engine.backend = parse_backend_arg(need_arg(&i, argc, argv, arg), arg);
|
||||
} else if (!strcmp(arg, "--cpu")) {
|
||||
server_log(DS4_LOG_DEFAULT, "ds4-server: server mode requires a graph backend");
|
||||
exit(2);
|
||||
} else {
|
||||
server_log(DS4_LOG_DEFAULT, "ds4-server: unknown option: %s", arg);
|
||||
|
||||
Reference in New Issue
Block a user