diff --git a/.gitignore b/.gitignore index cdf3cfcf89..a86e3511e4 100644 --- a/.gitignore +++ b/.gitignore @@ -28,6 +28,8 @@ /tests/test_metal_tp_spec /tests/test_metal_tp_cancel /tests/test_metal_moe_prefill +/tests/test_metal_slab_residency +/speed-bench/metal_slab_residency_bench /tests/test_metal_ssd_experts /tests/test_metal_command_memory /tests/test_metal_dense_mpp diff --git a/Makefile b/Makefile index 2d970935c3..994dbb0787 100644 --- a/Makefile +++ b/Makefile @@ -166,6 +166,19 @@ tests/test_metal_moe_prefill: tests/test_metal_moe_prefill.o $(CORE_OBJS) test-metal-moe-prefill: tests/test_metal_moe_prefill ./tests/test_metal_moe_prefill +tests/test_metal_slab_residency: tests/test_metal_slab_residency.m ds4_metal.m ds4_gpu.h ds4_image.o $(METAL_SRCS) + $(CC) $(OBJCFLAGS) -I. -o $@ $< ds4_image.o $(METAL_LDLIBS) + +.PHONY: test-metal-slab-residency +test-metal-slab-residency: tests/test_metal_slab_residency + MTL_DEBUG_LAYER=1 ./tests/test_metal_slab_residency + +speed-bench/metal_slab_residency_bench: speed-bench/metal_slab_residency_bench.m + $(CC) $(OBJCFLAGS) -o $@ $< $(METAL_LDLIBS) + +.PHONY: metal-slab-residency-bench +metal-slab-residency-bench: speed-bench/metal_slab_residency_bench + tests/test_qwen4_moe_mm_specialize.o: tests/test_qwen4_moe_mm_specialize.c ds4_gpu.h $(CC) $(CFLAGS) -fno-fast-math -I. -c -o $@ $< @@ -1023,6 +1036,7 @@ clean: rm -f tests/test_web_recovery rm -f tests/test_metal_ssd_experts rm -f tests/test_metal_command_memory + rm -f tests/test_metal_slab_residency speed-bench/metal_slab_residency_bench rm -f tests/test_deepseek41_metal rm -f tests/test_deepseek41_cuda rm -f tests/test_cuda_q8_rows diff --git a/ds4_metal.m b/ds4_metal.m index 0d7536993f..8771714036 100644 --- a/ds4_metal.m +++ b/ds4_metal.m @@ -1014,6 +1014,7 @@ static void ds4_gpu_print_device_summary(void) { static id g_stream_expert_cache_up_addr_buffers[DS4_METAL_STREAM_EXPERT_CACHE_MAX_LAYER]; static id g_stream_expert_cache_down_addr_buffers[DS4_METAL_STREAM_EXPERT_CACHE_MAX_LAYER]; static id g_stream_expert_cache_slabs[DS4_METAL_STREAM_EXPERT_CACHE_MAX_SLABS]; +static id g_stream_slab_residency_set; static uint32_t g_stream_expert_cache_slab_start_slot[DS4_METAL_STREAM_EXPERT_CACHE_MAX_SLABS]; static uint32_t g_stream_expert_cache_slab_slot_count[DS4_METAL_STREAM_EXPERT_CACHE_MAX_SLABS]; static uint32_t g_stream_expert_cache_slab_slots_used[DS4_METAL_STREAM_EXPERT_CACHE_MAX_SLABS]; @@ -4421,6 +4422,13 @@ void ds4_gpu_print_memory_report(const char *label) { (unsigned long long)g_stream_expert_cache_buffer_allocs, (unsigned long long)g_stream_expert_cache_buffer_reuses); } + if (g_stream_slab_residency_set) { + uint64_t bytes = 0; + for (uint32_t i = 0; i < g_stream_expert_cache_slab_count; i++) + bytes += [g_stream_expert_cache_slabs[i] length]; + fprintf(stderr, "ds4: streaming slab residency: %u slabs, %.2f GiB allocations\n", + g_stream_expert_cache_slab_count, ds4_gpu_gib(bytes)); + } if (g_stream_expert_cache_mlock_bytes != 0 || g_stream_expert_cache_mlock_failures != 0) { fprintf(stderr, @@ -13877,6 +13885,17 @@ static uint64_t ds4_gpu_stream_expert_slab_target_bytes(void) { return target; } +static void ds4_gpu_stream_slab_residency_clear(void) { +#if TARGET_OS_OSX + if (@available(macOS 15.0, *)) { + if (g_stream_slab_residency_set) { + [g_queue removeResidencySet:g_stream_slab_residency_set]; + g_stream_slab_residency_set = nil; + } + } +#endif +} + static id ds4_gpu_stream_expert_alloc_slab_buffer( uint64_t len, NSString *label) { @@ -13896,6 +13915,36 @@ static uint64_t ds4_gpu_stream_expert_slab_target_bytes(void) { } buffer.label = label; g_stream_expert_cache_buffer_allocs++; + if (getenv("DS4_METAL_STREAMING_SLAB_RESIDENCY") && + !g_stream_expert_cache_mlock_relief_applied) { +#if TARGET_OS_OSX + if (@available(macOS 15.0, *)) { + const BOOL fresh = g_stream_slab_residency_set == nil; + if (fresh) { + MTLResidencySetDescriptor *desc = [[MTLResidencySetDescriptor alloc] init]; + desc.label = @"ds4_streaming_expert_slabs"; + desc.initialCapacity = 64; + NSError *error = nil; + g_stream_slab_residency_set = [g_device newResidencySetWithDescriptor:desc error:&error]; + if (!g_stream_slab_residency_set) { + fprintf(stderr, "ds4: streaming slab residency set failed: %s\n", + [[error localizedDescription] UTF8String]); + return nil; + } + } + /* Only owned cache slabs belong here, never disk-backed model + * views. Keep registration across slot reuse; clear it with the + * physical slab pool after outstanding GPU work has drained. + * Queue attachment covers each submission without a separate + * explicit requestResidency lifetime. */ + [g_stream_slab_residency_set addAllocation:buffer]; + [g_stream_slab_residency_set commit]; + if (fresh) { + [g_queue addResidencySet:g_stream_slab_residency_set]; + } + } +#endif + } return buffer; } @@ -15184,6 +15233,7 @@ static void ds4_gpu_stream_expert_cache_clear_all(int reset_stats) { } g_stream_expert_cache_bytes = 0; g_stream_expert_cache_entry_count = 0; + ds4_gpu_stream_slab_residency_clear(); for (uint32_t i = 0; i < g_stream_expert_cache_slab_count; i++) { g_stream_expert_cache_slabs[i] = nil; g_stream_expert_cache_slab_start_slot[i] = 0; @@ -15725,6 +15775,9 @@ static uint32_t ds4_gpu_stream_expert_cache_release_mlock_margin( if (released == 0) return 0; g_stream_expert_cache_mlock_relief_applied = 1; + /* Let released slots become reclaimable instead of requesting the whole + * pool again on every submission. A cache rebuild can enable residency. */ + ds4_gpu_stream_slab_residency_clear(); uint32_t cap = g_stream_expert_cache_entry_count; const uint32_t locked_after = diff --git a/speed-bench/README.md b/speed-bench/README.md index 28666c0f3c..2410e124c3 100644 --- a/speed-bench/README.md +++ b/speed-bench/README.md @@ -27,6 +27,27 @@ python3 speed-bench/plot_speed.py speed-bench/m3_max.csv --title "M3 Max t/s" The script uses only the Python standard library. By default it writes a file next to the CSV using the `_ts.svg` suffix, such as `speed-bench/m3_max_ts.svg`. +### DeepSeek V4.1 Flash streaming slab residency + +On macOS 15+, `DS4_METAL_STREAMING_SLAB_RESIDENCY=1` attaches owned expert-cache +slabs to the existing Metal queue. It is opt-in; disk-backed model views and +Engram tables are excluded. Memory-pressure relief detaches the set until the +cache is rebuilt. See [the M2 Ultra investigation](v41_slab_residency_m2_ultra.md). + +A model-free reproducer checks every GPU result and toggles queue attachment +OFF/ON/OFF/ON. The following allocates and locks approximately 104 GiB; run it +alone on a host with enough available memory: + +```sh +make metal-slab-residency-bench test-metal-slab-residency +DS4_SLAB_BENCH_ALTERNATE_SMALL=1 ./speed-bench/metal_slab_residency_bench \ + toggle 26 4096 4080 24 6 > /tmp/slab-residency.csv +``` + +Arguments are mode, slab count, allocation MiB per slab, filled/locked MiB per +slab, iterations per phase and selected slabs per command buffer. Use eight +slabs instead of 26 for the 32 GiB control. No model file is opened. + ### Metal decode schedule A/B Build the balanced, same-engine Metal decode comparison with: diff --git a/speed-bench/metal_slab_residency_bench.m b/speed-bench/metal_slab_residency_bench.m new file mode 100644 index 0000000000..3fe34055b7 --- /dev/null +++ b/speed-bench/metal_slab_residency_bench.m @@ -0,0 +1,164 @@ +/* Model-free large-slab residency reproducer; see speed-bench/README.md. */ +#import +#import +#include +#include +#include +#include +#include +#include + +static double now_ms(void) { + struct timespec t; + clock_gettime(CLOCK_MONOTONIC, &t); + return t.tv_sec * 1000.0 + t.tv_nsec / 1e6; +} + +static uint64_t number(const char *text) { + char *end = NULL; + unsigned long long value = strtoull(text, &end, 10); + return text[0] && text[0] != '-' && end && !*end ? value : 0; +} + +int main(int argc, char **argv) { + if (argc != 7) { + fprintf(stderr, "usage: %s none|queue|toggle SLABS SLAB_MIB FILLED_MIB ITERATIONS SELECTED\n", + argv[0]); + return 2; + } + const char *mode = argv[1]; + const uint64_t count64 = number(argv[2]), mib = number(argv[3]); + const uint64_t filled_mib = number(argv[4]); + const uint64_t iterations64 = number(argv[5]), selected64 = number(argv[6]); + if ((strcmp(mode, "none") && strcmp(mode, "queue") && strcmp(mode, "toggle")) || + !count64 || count64 > UINT32_MAX || !mib || mib > UINT64_MAX / 1048576 || + filled_mib < 4 || filled_mib > mib || !iterations64 || iterations64 > UINT32_MAX || + !selected64 || selected64 > count64 || selected64 > UINT32_MAX / 1024) return 2; + const uint32_t count = (uint32_t)count64, iterations = (uint32_t)iterations64; + const uint32_t selected = (uint32_t)selected64; + const uint64_t bytes = mib * 1048576, filled = filled_mib * 1048576; + const BOOL alternate = getenv("DS4_SLAB_BENCH_ALTERNATE_SMALL") != NULL; + + if (@available(macOS 15.0, *)) { + @autoreleasepool { + id dev = MTLCreateSystemDefaultDevice(); + id queue = [dev newCommandQueue]; + if (!dev || !queue || bytes > dev.maxBufferLength) return 1; + fprintf(stderr, "device=%s mode=%s slabs=%u slab_MiB=%llu filled_MiB=%llu " + "selected=%u alternate_small=%d\n", dev.name.UTF8String, mode, + count, mib, filled_mib, selected, alternate); + NSString *source = @"#include \n" + "using namespace metal;\n" + "kernel void probe(device const ulong *a [[buffer(0)]], " + "device uint *out [[buffer(1)]], uint i [[thread_position_in_grid]]) { " + "device const uint *b = reinterpret_cast(a[i/1024]); " + "out[i] = b[(i%1024)*1024]; }"; + NSError *error = nil; + id lib = [dev newLibraryWithSource:source options:nil error:&error]; + if (!lib) { + fprintf(stderr, "Metal library: %s\n", error.description.UTF8String); + return 1; + } + id pipeline = [dev + newComputePipelineStateWithFunction:[lib newFunctionWithName:@"probe"] error:&error]; + if (!pipeline) return 1; + + NSMutableArray> *pool = [NSMutableArray array]; + for (uint32_t i = 0; i < count; i++) { + id b = [dev newBufferWithLength:bytes options:MTLResourceStorageModeShared]; + if (!b) return 1; + memset(b.contents, i % 127 + 1, filled); + /* Match the Q2 expert-slot lock granularity. No model file IO. */ + const uint64_t slot_bytes = 9961472; + for (uint64_t off = 0; off < filled; off += slot_bytes) { + const uint64_t n = filled - off < slot_bytes ? filled - off : slot_bytes; + if (mlock((char *)b.contents + off, n)) { + perror("mlock"); + return 1; + } + } + [pool addObject:b]; + } + id small = [dev newBufferWithLength:4194304 options:MTLResourceStorageModeShared]; + id addresses = [dev newBufferWithLength:selected * sizeof(uint64_t) + options:MTLResourceStorageModeShared]; + id out = [dev newBufferWithLength:selected * 1024 * sizeof(uint32_t) + options:MTLResourceStorageModeShared]; + if (!small || !addresses || !out) return 1; + memset(small.contents, 1, 4194304); + id set = nil; + BOOL queued = NO; + if (strcmp(mode, "none")) { + MTLResidencySetDescriptor *desc = [MTLResidencySetDescriptor new]; + desc.initialCapacity = count; + set = [dev newResidencySetWithDescriptor:desc error:&error]; + if (!set) return 1; + for (id b in pool) [set addAllocation:b]; + [set commit]; + if (!strcmp(mode, "queue")) { + [queue addResidencySet:set]; + queued = YES; + } + } + printf("phase,iteration,kind,wall_ms,gpu_ms,driver_ms,checksum\n"); + const uint32_t phases = !strcmp(mode, "toggle") ? 4 : 1; + for (uint32_t phase = 0; phase < phases; phase++) { + if (phases > 1) { + if (phase % 2) { + [queue addResidencySet:set]; + queued = YES; + } else if (queued) { + [queue removeResidencySet:set]; + queued = NO; + } + } + for (uint32_t it = 0; it < iterations; it++) { + @autoreleasepool { + const BOOL use_small = alternate && (it % 2); + const double begin = now_ms(); + id cb = [queue commandBuffer]; + id enc = [cb computeCommandEncoder]; + uint64_t expected = 0; + for (uint32_t j = 0; j < selected; j++) { + const uint32_t k = ((uint64_t)it * selected + j) % count; + id b = use_small ? small : pool[k]; + ((uint64_t *)addresses.contents)[j] = b.gpuAddress; + [enc useResource:b usage:MTLResourceUsageRead]; + expected += (uint64_t)(0x01010101u * (use_small ? 1 : k % 127 + 1)) * 1024; + } + [enc setComputePipelineState:pipeline]; + [enc setBuffer:addresses offset:0 atIndex:0]; + [enc setBuffer:out offset:0 atIndex:1]; + [enc dispatchThreadgroups:MTLSizeMake(selected * 4, 1, 1) + threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; + [enc endEncoding]; + [cb commit]; + [cb waitUntilCompleted]; + const double end = now_ms(); + if (cb.status != MTLCommandBufferStatusCompleted) { + fprintf(stderr, "command buffer: %s\n", cb.error.description.UTF8String); + return 1; + } + uint64_t checksum = 0; + for (uint32_t j = 0; j < selected * 1024; j++) checksum += ((uint32_t *)out.contents)[j]; + if (checksum != expected) { + fprintf(stderr, "checksum mismatch\n"); + return 1; + } + printf("%u,%u,%s,%.6f,%.6f,%.6f,%llu\n", phase, it, + use_small ? "small" : "slabs", end - begin, + (cb.GPUEndTime - cb.GPUStartTime) * 1000, + (cb.kernelEndTime - cb.kernelStartTime) * 1000, checksum); + } + } + fflush(stdout); + } + if (queued) [queue removeResidencySet:set]; + for (id b in pool) munlock(b.contents, filled); + } + } else { + fprintf(stderr, "requires macOS 15 or later\n"); + return 1; + } + return 0; +} diff --git a/speed-bench/v41_slab_residency_m2_ultra.csv b/speed-bench/v41_slab_residency_m2_ultra.csv new file mode 100644 index 0000000000..aed28dc051 --- /dev/null +++ b/speed-bench/v41_slab_residency_m2_ultra.csv @@ -0,0 +1,5 @@ +run,variant,ctx_tokens,prefill_tokens,prefill_tps,gen_tokens,gen_tps,gen_first_ms,gen_steady_tokens,gen_steady_tps,kvcache_bytes +bench-0-control,control,2048,2048,76.15,8,0.25,4103.550,7,0.25,0 +bench-1-candidate,candidate,2048,2048,70.95,8,9.57,329.048,7,13.83,0 +bench-2-candidate,candidate,2048,2048,69.16,8,9.67,316.315,7,13.71,0 +bench-3-control,control,2048,2048,60.11,8,0.25,3935.309,7,0.25,0 diff --git a/speed-bench/v41_slab_residency_m2_ultra.md b/speed-bench/v41_slab_residency_m2_ultra.md new file mode 100644 index 0000000000..7e79141ab2 --- /dev/null +++ b/speed-bench/v41_slab_residency_m2_ultra.md @@ -0,0 +1,154 @@ +# DeepSeek V4.1 Flash: owned Metal slab residency + +For current-upstream results and the known upstream router test failure, see +[2026-09-14 revalidation](#upstream-revalidation-2026-09-14). +The diagnosis and validation below describe the original 2026-09-13 base. + +## Problem and diagnosis + +DeepSeek V4.1 Flash SSD decoding on this host fell to approximately 0.26 tokens/s +with the default owned expert-cache slabs. GPU execution itself did not explain +the multi-second token latency. Disabling slabs removed the severe cliff. +Registering allocations in an unattached residency set did not help; attaching +the set to the existing command queue did. Additional CPU `mlock` and warmup +alone did not remove the delay. + +The model-free benchmark in this change allocates shared Metal buffers, writes +and locks them on the CPU, then reads small, dispersed samples through GPU +addresses with `useResource`. It performs no model-file IO. With a 104 GiB pool, +the initial queue-detached phase spent about 250 ms per large-buffer submission, +of which approximately 249.7 ms was in Metal's reported kernel/driver interval; +GPU execution was approximately 0.013 ms. Queue attachment reduced submission +wall time to approximately 0.180 ms. The 32 GiB control did not exhibit the cliff. +These per-phase means exclude the first four iterations of each phase. +These observations locate the delay in large-allocation submission/residency +handling, without claiming a particular private driver implementation or a +hardware-wide defect. + +The queue-toggle run retains the residency set while detached. After its first +attachment, phase 2 remains fast: detachment alone is not an eviction command. +Do not interpret it as a reversible OFF/ON/OFF/ON performance experiment. +Fresh-process model comparisons below supply independent controls. + +## Change + +`DS4_METAL_STREAMING_SLAB_RESIDENCY=1` registers only owned expert-cache slab +buffers and attaches the set once to the existing queue on macOS 15+. It adds +no explicit `requestResidency`/`endResidency` pair. Cache teardown removes and +releases the set with the physical pool. Successful mlock-margin relief also +removes/releases it, preventing later slab allocation from reattaching until +the cache is rebuilt. Model mappings and disk-only Engram tables are excluded. +The feature remains opt-in pending evidence on other hosts and memory budgets. + +## Environment and results + +Measured on 2026-09-13, Apple M2 Ultra, 192 GiB unified memory, macOS 15.7.4, +Metal, against upstream `bd66c402070042bf0a79ad6ece8242de4c93680c`. +Model: DeepSeek V4.1 Flash calibrated IQ2_XXS/Q2_K, 365,713,686,528 bytes; +SHA-256 `1ce6a8f8806205c13330d7ca287bd198331dc5ca35ccc5d8a9a92a188a6f6f42`. +The model was reused without conversion. The machine's existing +`iogpu.wired_limit_mb=188000` was unchanged. One inference/Metal benchmark ran +at a time; ordinary desktop background services remained running. + +All model benchmarks use `speed-bench/promessi_sposi.txt`, SSD streaming, +a 2,048-token prefill, 8,257 allocated context, default power and automatic +expert-cache sizing. The planner reports 135.26 GiB dynamic expert cache plus +7.12 GiB prefill headroom. Each process starts a fresh engine/cache; the OS/file +cache is not flushed. Order is control, candidate, candidate, control. Controls +use the same branch binary with the opt-in flag absent. The production path +with the flag absent is unchanged from the upstream base. + +[Raw model CSV](v41_slab_residency_m2_ultra.csv); model-free [104 GiB](v41_slab_residency_m2_ultra_repro_26.csv) and [32 GiB](v41_slab_residency_m2_ultra_repro_8.csv) CSVs. + +| Run | Variant | Prefill t/s | Decode t/s (8 tokens) | First token ms | Steady t/s (7 tokens) | +| --- | --- | ---: | ---: | ---: | ---: | +| 1 | control | 76.15 | 0.25 | 4103.550 | 0.25 | +| 2 | candidate | 70.95 | 9.57 | 329.048 | 13.83 | +| 3 | candidate | 69.16 | 9.67 | 316.315 | 13.71 | +| 4 | control | 60.11 | 0.25 | 3935.309 | 0.25 | + +All four decoded outputs are identical. Eight-token decode intentionally keeps +the multi-second control bounded; its first-token cost materially affects the +aggregate. This is not a 512-token sustained-throughput claim. Prefill varies +from 60.11 to 76.15 t/s across controls and 69.16 to 70.95 t/s across candidates; +these two pairs do not establish a precise prefill speedup. + +Reproduce each arm with the flag absent or set to `1`: + +```sh +DS4_METAL_STREAMING_SLAB_RESIDENCY=1 ./ds4-bench -m MODEL --ssd-streaming \ + --prompt-file speed-bench/promessi_sposi.txt --ctx-start 2048 --ctx-max 2048 \ + --ctx-alloc 8257 --gen-tokens 8 --show-output --csv /tmp/slab-on.csv +``` + +## Validation + +- Clean default Metal build, CPU compile, and restored Metal executable links. +- `make test-metal-slab-residency` with Metal API validation: protected expert + survives relief; an unprotected slot is unlocked/evicted; the queue releases + the set; allocation cannot reattach after relief; rebuilding reenables it; + weak references confirm the set and old physical buffers are destroyed. +- `make test-metal-ssd-experts test-metal-moe-prefill`, plus SSD + `--table-admission` and MoE `--ssd-address`, under Metal API validation. +- With the residency flag enabled and Metal API validation: + `tests/test_deepseek41_graph MODEL --session-fixture` (real weights, snapshots, + restored logits, prefix reuse, cancellation, boundary and malformed-state checks). +- `make test-frontends test-engram test-deepseek41-gguf test-quality-api`. +- Both 104 GiB and 32 GiB model-free runs check every GPU checksum. + +The macOS SDK 15 build reports the same two pre-existing unused Metal 4 symbol +warnings as the upstream base. Full legacy `make test` was not run: its default +model vectors target the older Flash checkpoint; the V4.1 checks above were +used. This is single-host Metal evidence, not a release sign-off or validation +of CUDA, ROCm, RDMA, other Macs, or all memory-pressure conditions. + +## Upstream revalidation (2026-09-14) + +Revalidated after merging upstream `a04f46fa423e45712c8c7e430eff422479f314a3` +(DeepSeek V4.1 CUDA support). The original measurements above remain historical. +The host, exact model, prompt, context allocation, cache policy, and independent +ABBA procedure are unchanged. Each variant was measured twice on this host. + +[New raw model CSV](v41_slab_residency_m2_ultra_20260914.csv). + +| Run | Variant | Prefill t/s | Decode t/s | First token ms | Steady t/s | +| --- | --- | ---: | ---: | ---: | ---: | +| 1 | control | 72.13 | 0.24 | 4289.607 | 0.24 | +| 2 | candidate | 67.99 | 9.35 | 333.096 | 13.41 | +| 3 | candidate | 67.81 | 9.27 | 336.455 | 13.31 | +| 4 | control | 61.97 | 0.24 | 4037.232 | 0.24 | + +Mean 8-token decode: 0.240 → 9.310 t/s (38.8×). +All four decoded outputs match exactly. + +Validation rerun: clean Metal build, CPU compilation and restored Metal links; +frontend, Engram, V4.1 GGUF and quality-tool unit tests. Under Metal API +validation, compact carry, index scores/top-k, general top-k, index projection, +embedding and TP attention subtests pass. Slab lifecycle, SSD expert kernels, +MoE prefill, admission/address checks and the real-model session fixture pass. + +The full `test-deepseek41-metal` suite fails in the new upstream router test +at `tests/test_deepseek41_metal.c:105`. An unmodified checkout of the same +upstream commit reproduces exactly the same failure under Metal API validation: + +```text +router n=256 mode=0 rows=1 row=0 expert=4 +logit=-11.8886719 actual=0.00263455603 ref=0.00262947031 +``` + +No kernel or tolerance was changed to bypass it. The full kernel suite is +therefore not green. CUDA/ROCm hardware and full legacy `make test` were not +executed. The two existing SDK 15 unused Metal 4 symbol warnings remain. + +## Upstream integration (2026-09-15) + +Merged upstream `9139e2ae58a41503968a500f36f75895c1ba63fc` and retained both +sets of test targets/documentation at the conflict. Clean Metal and CPU builds, +frontend/Engram/V4.1 GGUF/quality-tool tests and the real-model session fixture +with this optimization enabled pass. The session fixture ran with Metal API +validation. +Slab lifecycle and SSD expert/MoE kernel tests also pass. + +Performance numbers above are the 2026-09-14 measurements, not a new benchmark +on this base. The unchanged upstream router failure is tracked by #1039 and +the separate accuracy fix #1044; neither is bundled into this optimization. diff --git a/speed-bench/v41_slab_residency_m2_ultra_20260914.csv b/speed-bench/v41_slab_residency_m2_ultra_20260914.csv new file mode 100644 index 0000000000..70e47db6a6 --- /dev/null +++ b/speed-bench/v41_slab_residency_m2_ultra_20260914.csv @@ -0,0 +1,5 @@ +run,variant,ctx_tokens,prefill_tokens,prefill_tps,gen_tokens,gen_tps,gen_first_ms,gen_steady_tokens,gen_steady_tps,kvcache_bytes +bench-0-control,control,2048,2048,72.13,8,0.24,4289.607,7,0.24,0 +bench-1-candidate,candidate,2048,2048,67.99,8,9.35,333.096,7,13.41,0 +bench-2-candidate,candidate,2048,2048,67.81,8,9.27,336.455,7,13.31,0 +bench-3-control,control,2048,2048,61.97,8,0.24,4037.232,7,0.24,0 diff --git a/speed-bench/v41_slab_residency_m2_ultra_repro_26.csv b/speed-bench/v41_slab_residency_m2_ultra_repro_26.csv new file mode 100644 index 0000000000..6551d6f624 --- /dev/null +++ b/speed-bench/v41_slab_residency_m2_ultra_repro_26.csv @@ -0,0 +1,97 @@ +phase,iteration,kind,wall_ms,gpu_ms,driver_ms,checksum +0,0,slabs,684.585000,0.013125,683.521750,362192065536 +0,1,small,0.389000,0.011791,0.181875,103483447296 +0,2,slabs,684.771000,0.010584,683.770208,1603993433088 +0,3,small,0.281000,0.011250,0.078750,103483447296 +0,4,slabs,229.034000,0.013625,228.267083,1052081714176 +0,5,small,0.269000,0.010792,0.077875,103483447296 +0,6,slabs,226.538000,0.013000,225.672958,1397026538496 +0,7,small,0.267000,0.010834,0.069875,103483447296 +0,8,slabs,256.071000,0.012875,255.331125,1741971362816 +0,9,small,0.262000,0.011375,0.067250,103483447296 +0,10,slabs,226.239000,0.013000,225.374542,1190059643904 +0,11,small,0.286000,0.011083,0.081625,103483447296 +0,12,slabs,284.960000,0.013417,284.096084,2431861011456 +0,13,small,0.263000,0.011250,0.069875,103483447296 +0,14,slabs,262.161000,0.012875,261.410959,983092749312 +0,15,small,0.226000,0.011208,0.039000,103483447296 +0,16,slabs,258.342000,0.013375,257.519125,2224894116864 +0,17,small,0.235000,0.011667,0.047500,103483447296 +0,18,slabs,245.338000,0.012833,244.584625,776125854720 +0,19,small,0.245000,0.011625,0.046125,103483447296 +0,20,slabs,233.212000,0.013125,232.383750,2017927222272 +0,21,small,0.228000,0.010875,0.041167,103483447296 +0,22,slabs,283.379000,0.013042,282.521875,569158960128 +0,23,small,0.224000,0.011042,0.038250,103483447296 +1,0,slabs,1775.149000,0.013375,1774.253125,362192065536 +1,1,small,0.375000,0.011708,0.174375,103483447296 +1,2,slabs,0.199000,0.008375,0.036250,1603993433088 +1,3,small,0.195000,0.007000,0.036083,103483447296 +1,4,slabs,0.206000,0.006625,0.031875,1052081714176 +1,5,small,0.176000,0.006375,0.030375,103483447296 +1,6,slabs,0.184000,0.006833,0.029708,1397026538496 +1,7,small,0.176000,0.006333,0.029125,103483447296 +1,8,slabs,0.188000,0.007000,0.041666,1741971362816 +1,9,small,0.174000,0.006292,0.029583,103483447296 +1,10,slabs,0.174000,0.006833,0.030667,1190059643904 +1,11,small,0.172000,0.006250,0.029625,103483447296 +1,12,slabs,0.183000,0.006917,0.029375,2431861011456 +1,13,small,0.181000,0.006250,0.035125,103483447296 +1,14,slabs,0.175000,0.006750,0.030500,983092749312 +1,15,small,0.175000,0.006208,0.028833,103483447296 +1,16,slabs,0.161000,0.007000,0.023041,2224894116864 +1,17,small,0.156000,0.006250,0.021583,103483447296 +1,18,slabs,0.148000,0.006000,0.023041,776125854720 +1,19,small,0.190000,0.006333,0.031834,103483447296 +1,20,slabs,0.190000,0.006000,0.025959,2017927222272 +1,21,small,0.189000,0.006333,0.028125,103483447296 +1,22,slabs,0.192000,0.005958,0.033000,569158960128 +1,23,small,0.176000,0.006375,0.027500,103483447296 +2,0,slabs,0.177000,0.005750,0.025625,362192065536 +2,1,small,0.180000,0.006250,0.027500,103483447296 +2,2,slabs,0.184000,0.006000,0.028000,1603993433088 +2,3,small,0.185000,0.006375,0.028000,103483447296 +2,4,slabs,0.180000,0.006000,0.026250,1052081714176 +2,5,small,0.176000,0.006500,0.026084,103483447296 +2,6,slabs,0.179000,0.005833,0.026250,1397026538496 +2,7,small,0.177000,0.006208,0.025750,103483447296 +2,8,slabs,0.182000,0.005875,0.027125,1741971362816 +2,9,small,0.180000,0.006750,0.027041,103483447296 +2,10,slabs,0.180000,0.006167,0.028042,1190059643904 +2,11,small,0.178000,0.006250,0.027000,103483447296 +2,12,slabs,0.181000,0.006375,0.028042,2431861011456 +2,13,small,0.176000,0.006667,0.026917,103483447296 +2,14,slabs,0.181000,0.006375,0.026625,983092749312 +2,15,small,0.174000,0.006250,0.026875,103483447296 +2,16,slabs,0.177000,0.005917,0.027000,2224894116864 +2,17,small,0.169000,0.006209,0.024000,103483447296 +2,18,slabs,0.174000,0.005875,0.024875,776125854720 +2,19,small,0.175000,0.006250,0.026000,103483447296 +2,20,slabs,0.175000,0.005958,0.026458,2017927222272 +2,21,small,0.174000,0.006333,0.026584,103483447296 +2,22,slabs,0.174000,0.005875,0.027083,569158960128 +2,23,small,0.175000,0.006041,0.026375,103483447296 +3,0,slabs,0.177000,0.005667,0.027125,362192065536 +3,1,small,0.277000,0.006041,0.026583,103483447296 +3,2,slabs,0.184000,0.005833,0.027750,1603993433088 +3,3,small,0.179000,0.006375,0.026833,103483447296 +3,4,slabs,0.182000,0.005875,0.027250,1052081714176 +3,5,small,0.176000,0.006458,0.026875,103483447296 +3,6,slabs,0.180000,0.005792,0.027667,1397026538496 +3,7,small,0.180000,0.006500,0.026583,103483447296 +3,8,slabs,0.181000,0.006000,0.027250,1741971362816 +3,9,small,0.176000,0.006375,0.026792,103483447296 +3,10,slabs,0.179000,0.006000,0.027750,1190059643904 +3,11,small,0.178000,0.006375,0.024708,103483447296 +3,12,slabs,0.176000,0.005750,0.024750,2431861011456 +3,13,small,0.176000,0.006375,0.025250,103483447296 +3,14,slabs,0.181000,0.005875,0.027250,983092749312 +3,15,small,0.179000,0.006250,0.026625,103483447296 +3,16,slabs,0.182000,0.005792,0.027375,2224894116864 +3,17,small,0.180000,0.006250,0.026375,103483447296 +3,18,slabs,0.182000,0.005917,0.028250,776125854720 +3,19,small,0.173000,0.006125,0.026000,103483447296 +3,20,slabs,0.171000,0.005875,0.026375,2017927222272 +3,21,small,0.171000,0.006250,0.026125,103483447296 +3,22,slabs,0.166000,0.005958,0.022292,569158960128 +3,23,small,0.166000,0.006375,0.021250,103483447296 diff --git a/speed-bench/v41_slab_residency_m2_ultra_repro_8.csv b/speed-bench/v41_slab_residency_m2_ultra_repro_8.csv new file mode 100644 index 0000000000..e78daa9c11 --- /dev/null +++ b/speed-bench/v41_slab_residency_m2_ultra_repro_8.csv @@ -0,0 +1,97 @@ +phase,iteration,kind,wall_ms,gpu_ms,driver_ms,checksum +0,0,slabs,696.287000,0.013167,695.189625,362192065536 +0,1,small,0.377000,0.010750,0.168708,103483447296 +0,2,slabs,224.870000,0.012792,224.003500,500169995264 +0,3,small,0.266000,0.010750,0.071833,103483447296 +0,4,slabs,0.241000,0.007541,0.070750,362192065536 +0,5,small,0.237000,0.006583,0.055375,103483447296 +0,6,slabs,0.221000,0.005959,0.060792,500169995264 +0,7,small,0.197000,0.006333,0.052209,103483447296 +0,8,slabs,0.207000,0.005708,0.054333,362192065536 +0,9,small,0.198000,0.006333,0.053625,103483447296 +0,10,slabs,0.194000,0.005959,0.055292,500169995264 +0,11,small,0.195000,0.006083,0.045375,103483447296 +0,12,slabs,0.191000,0.006000,0.048625,362192065536 +0,13,small,0.191000,0.006208,0.046750,103483447296 +0,14,slabs,0.193000,0.006250,0.050208,500169995264 +0,15,small,0.171000,0.006125,0.028750,103483447296 +0,16,slabs,0.167000,0.005750,0.028125,362192065536 +0,17,small,0.165000,0.006167,0.027083,103483447296 +0,18,slabs,0.276000,0.005917,0.023875,500169995264 +0,19,small,0.172000,0.006167,0.025500,103483447296 +0,20,slabs,0.182000,0.005667,0.032375,362192065536 +0,21,small,0.167000,0.006125,0.027042,103483447296 +0,22,slabs,0.170000,0.005625,0.027000,500169995264 +0,23,small,0.161000,0.006250,0.025792,103483447296 +1,0,slabs,0.167000,0.005791,0.026667,362192065536 +1,1,small,0.162000,0.006333,0.026167,103483447296 +1,2,slabs,0.167000,0.005917,0.026917,500169995264 +1,3,small,0.164000,0.006500,0.025542,103483447296 +1,4,slabs,0.167000,0.005667,0.027209,362192065536 +1,5,small,0.166000,0.006583,0.026292,103483447296 +1,6,slabs,0.166000,0.005750,0.027583,500169995264 +1,7,small,0.161000,0.006083,0.025917,103483447296 +1,8,slabs,0.167000,0.005875,0.027750,362192065536 +1,9,small,0.163000,0.006167,0.025833,103483447296 +1,10,slabs,0.165000,0.005833,0.027208,500169995264 +1,11,small,0.163000,0.006416,0.025833,103483447296 +1,12,slabs,0.167000,0.005875,0.027625,362192065536 +1,13,small,0.164000,0.006500,0.025875,103483447296 +1,14,slabs,0.165000,0.005958,0.026584,500169995264 +1,15,small,0.163000,0.006167,0.025375,103483447296 +1,16,slabs,0.171000,0.005833,0.026542,362192065536 +1,17,small,0.183000,0.006250,0.027833,103483447296 +1,18,slabs,0.176000,0.005875,0.025125,500169995264 +1,19,small,0.174000,0.006125,0.024625,103483447296 +1,20,slabs,0.181000,0.005833,0.028167,362192065536 +1,21,small,0.177000,0.006250,0.027625,103483447296 +1,22,slabs,0.180000,0.005875,0.028084,500169995264 +1,23,small,0.176000,0.006583,0.027083,103483447296 +2,0,slabs,0.179000,0.005875,0.027875,362192065536 +2,1,small,0.186000,0.006375,0.025875,103483447296 +2,2,slabs,0.181000,0.005875,0.028042,500169995264 +2,3,small,0.173000,0.006500,0.026250,103483447296 +2,4,slabs,0.179000,0.006125,0.028292,362192065536 +2,5,small,0.172000,0.006125,0.025125,103483447296 +2,6,slabs,0.187000,0.005875,0.027500,500169995264 +2,7,small,0.176000,0.006500,0.026500,103483447296 +2,8,slabs,0.177000,0.005750,0.027375,362192065536 +2,9,small,0.180000,0.006125,0.026667,103483447296 +2,10,slabs,0.176000,0.005875,0.026583,500169995264 +2,11,small,0.172000,0.006125,0.025125,103483447296 +2,12,slabs,0.181000,0.005792,0.027625,362192065536 +2,13,small,0.177000,0.006209,0.026166,103483447296 +2,14,slabs,0.178000,0.005833,0.027459,500169995264 +2,15,small,0.175000,0.006250,0.026875,103483447296 +2,16,slabs,0.183000,0.005875,0.027125,362192065536 +2,17,small,0.175000,0.006000,0.027208,103483447296 +2,18,slabs,0.179000,0.005792,0.027958,500169995264 +2,19,small,0.172000,0.006167,0.025625,103483447296 +2,20,slabs,0.182000,0.005875,0.028000,362192065536 +2,21,small,0.177000,0.006250,0.027000,103483447296 +2,22,slabs,0.178000,0.005875,0.027916,500169995264 +2,23,small,0.178000,0.006416,0.026750,103483447296 +3,0,slabs,0.174000,0.005625,0.028250,362192065536 +3,1,small,0.168000,0.006000,0.026625,103483447296 +3,2,slabs,0.173000,0.005709,0.027500,500169995264 +3,3,small,0.168000,0.006250,0.026375,103483447296 +3,4,slabs,0.171000,0.005833,0.027417,362192065536 +3,5,small,0.167000,0.006208,0.027000,103483447296 +3,6,slabs,0.173000,0.006250,0.027250,500169995264 +3,7,small,0.206000,0.006292,0.025625,103483447296 +3,8,slabs,0.175000,0.005625,0.027333,362192065536 +3,9,small,0.171000,0.005875,0.026500,103483447296 +3,10,slabs,0.172000,0.005875,0.027833,500169995264 +3,11,small,0.172000,0.006167,0.026084,103483447296 +3,12,slabs,0.173000,0.006167,0.028042,362192065536 +3,13,small,0.169000,0.006125,0.026583,103483447296 +3,14,slabs,0.175000,0.005959,0.027417,500169995264 +3,15,small,0.169000,0.006250,0.026166,103483447296 +3,16,slabs,0.176000,0.005875,0.027375,362192065536 +3,17,small,0.170000,0.006250,0.026625,103483447296 +3,18,slabs,0.175000,0.005875,0.027916,500169995264 +3,19,small,0.172000,0.006167,0.026000,103483447296 +3,20,slabs,0.176000,0.005750,0.027458,362192065536 +3,21,small,0.164000,0.006000,0.024375,103483447296 +3,22,slabs,0.169000,0.005875,0.026250,500169995264 +3,23,small,0.194000,0.006208,0.025750,103483447296 diff --git a/tests/test_metal_slab_residency.m b/tests/test_metal_slab_residency.m new file mode 100644 index 0000000000..cc12fca4c2 --- /dev/null +++ b/tests/test_metal_slab_residency.m @@ -0,0 +1,104 @@ +/* Real Metal objects with a small synthetic cache: exercise the existing + * relief/clear paths without exhausting the host's memory. */ +#include "../ds4_metal.m" +#include + +bool ds4_log_is_tty(FILE *fp) { + (void)fp; + return false; +} + +static void add_slot(uint32_t expert) { + const uint64_t part_bytes = 256u * 1024u; + id gate = nil, up = nil, down = nil; + NSUInteger gate_inner = 0, up_inner = 0, down_inner = 0; + assert(ds4_gpu_stream_expert_alloc_slab_slot(part_bytes, part_bytes, + &gate, &up, &down, &gate_inner, &up_inner, &down_inner)); + uint32_t slot; + assert(ds4_gpu_stream_expert_slab_slot_for_buffer(gate, gate_inner, &slot)); + memset((char *)gate.contents + gate_inner, expert + 1, 3 * part_bytes); + assert(ds4_gpu_stream_expert_slab_lock_slot(slot)); + + ds4_gpu_stream_expert_cache_entry *entry = &g_stream_expert_cache[0][expert]; + entry->valid = 1; + entry->slab_backed = 1; + entry->slab_slot = slot; + entry->gate_buffer = gate; + entry->up_buffer = up; + entry->down_buffer = down; + entry->gate_inner = gate_inner; + entry->up_inner = up_inner; + entry->down_inner = down_inner; + entry->gate_expert_bytes = part_bytes; + entry->down_expert_bytes = part_bytes; + entry->logical_bytes = 3 * part_bytes; + entry->last_used = expert; + g_stream_expert_cache_entry_count++; + g_stream_expert_cache_layer_count[0]++; + g_stream_expert_cache_bytes += entry->logical_bytes; +} + +static void submit_and_drain(void) { + id cb = [g_queue commandBuffer]; + id enc = [cb blitCommandEncoder]; + [enc fillBuffer:g_stream_expert_cache_slabs[0] range:NSMakeRange(0, 16) value:7]; + [enc endEncoding]; + [cb commit]; + [cb waitUntilCompleted]; + assert(cb.status == MTLCommandBufferStatusCompleted); +} + +int main(void) { + if (@available(macOS 15.0, *)) { + @autoreleasepool { + assert(ds4_gpu_init()); + unsetenv("DS4_METAL_DISABLE_STREAMING_EXPERT_SLABS"); + setenv("DS4_METAL_STREAMING_SLAB_RESIDENCY", "1", 1); + setenv("DS4_METAL_STREAMING_EXPERT_SLAB_MB", "16", 1); + ds4_gpu_set_ssd_streaming(true); + ds4_gpu_set_streaming_expert_cache_budget(32); + __weak id old_set = nil; + __weak id old_buffer = nil; + @autoreleasepool { + for (uint32_t i = 0; i < 10; i++) add_slot(i); + assert(g_stream_slab_residency_set); + old_set = g_stream_slab_residency_set; + old_buffer = g_stream_expert_cache_slabs[0]; + submit_and_drain(); + const uint64_t locked = g_stream_expert_cache_mlock_bytes; + const int32_t protected = 0; + assert(ds4_gpu_stream_expert_cache_release_mlock_margin(0, &protected, 1) == 1); + assert(g_stream_expert_cache[0][0].valid); + assert(g_stream_expert_cache_mlock_bytes == locked - 3 * 256 * 1024); + assert(g_stream_expert_cache_entry_count == 9); + assert(!g_stream_slab_residency_set); + submit_and_drain(); + /* Even another allocation cannot reattach after relief. */ + id extra = ds4_gpu_stream_expert_alloc_slab_buffer( + 1024 * 1024, @"after relief"); + assert(extra && !g_stream_slab_residency_set); + assert(ds4_gpu_stream_expert_cache_release_mlock_margin(0, &protected, 1) == 0); + } + assert(!old_set); /* The queue no longer retains the set. */ + @autoreleasepool { + ds4_gpu_set_streaming_expert_cache_budget(32); + } + assert(!old_buffer); + assert(g_stream_expert_cache_mlock_bytes == 0); + assert(!g_stream_expert_cache_mlock_relief_applied); + @autoreleasepool { + add_slot(0); + assert(g_stream_slab_residency_set); + submit_and_drain(); + old_set = g_stream_slab_residency_set; + old_buffer = g_stream_expert_cache_slabs[0]; + ds4_gpu_set_streaming_expert_cache_budget(16); + } + assert(!old_set && !old_buffer); + puts("Metal slab residency: protected-slot relief, no reattach, rebuild and deallocation: PASS"); + } + } else { + puts("Metal slab residency: SKIP (requires macOS 15)"); + } + return 0; +}