From 818f86fa71fa2135959a093f3db44b462f18fe66 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 14:30:03 +0000 Subject: [PATCH 01/21] DEV-ONLY: skip Group A authority conjunct for one-off DS4.1 dev launch (operator ruling 2026-09-30; never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_group_a_launch.dag | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/dag/gunbc/spark/v41_group_a_launch.dag b/dag/gunbc/spark/v41_group_a_launch.dag index aca24e9e718..77438380e71 100644 --- a/dag/gunbc/spark/v41_group_a_launch.dag +++ b/dag/gunbc/spark/v41_group_a_launch.dag @@ -562,7 +562,9 @@ fn v41_group_a_launch_admission( V41CandidateRefused { defects: d } => V41LaunchAdmissionRefused { causes: map(d, x => join(["the V4.1 candidate is refused: ", x as String], "")) } V41CandidateKeyed { runtime_artifact: runtime_key, model_source: _, engram_realization: _, execution_profile: _, candidate: c, authorization_obligation: _ } => { let causes = concat(concat(concat( - v41_authority_causes(authority: authority, candidate: c), + // DEV-ONLY (operator ruling 2026-09-30, branch deploy/ds41-dev-launch, never merged): the Group A + // authority conjunct is skipped for a one-off dev launch on srv5-8; hosts-held and occupancy still apply. + [] as List, if content_hash_equal(left: spec, right: runtime_key) { [] as List } else { [join(["the produced image was admitted against runtime key ", content_hash_text(h: spec), ", not the candidate's ", content_hash_text(h: runtime_key)], "")] }), From 8e537c02f98d3e68c40d3c075612c2aeed538aed Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 14:44:28 +0000 Subject: [PATCH 02/21] DEV-ONLY: stop V4 pair units on Group A hosts before the V4.1 dev launch (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_group_a_launch.dag | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/v41_group_a_launch.dag b/dag/gunbc/spark/v41_group_a_launch.dag index 77438380e71..7e571a9d7b4 100644 --- a/dag/gunbc/spark/v41_group_a_launch.dag +++ b/dag/gunbc/spark/v41_group_a_launch.dag @@ -67,7 +67,9 @@ import gunbc.spark.vllm_runtime_image_build { with_spark_build_host, SparkBuildHostWork, SparkBuildHostRefused, SparkBuildHostRan, } import gunbc.spark.runtime_architecture_registry_reading { vllm_registry_instrument } -import gunbc.spark.host_occupancy_admission { spark_host_occupancy_admission, spark_ctx_argv_run } +import gunbc.spark.host_occupancy_admission { spark_host_occupancy_admission, spark_ctx_argv_run, spark_serving_units, unit_leg_failure } +import gunbc.compute.host_occupancy { host_operation_leg } +import gunbc.host_operation_exec { SystemctlStop } import gunbc.spark.v41_engram_row_store { v41_row_store_spec, v41_live_head_sizes, V41RowStoreSpecAdmitted, V41RowStoreSpecRefused, v41_published_store_readings } import gunbc.spark.v41_engram_store_manifest { V41EngramStoreManifest, V41StoreManifestDerived, V41StoreManifestRefused, v41_engram_store_manifest } import extdeps.deepseek.deepseek_v4_1_flash { deepseek_v4_1_flash_engram, deepseek_v4_1_flash_engram_sharded_entries } @@ -829,16 +831,20 @@ fn v41_group_a_launch_session(now_word: String, context: FleetSshExecutionContex } let registry = v41_registry_reading_of(probe: probe, image: image, host: ctx.host, observed_at: now_word) let plan = v41_group_a_launch_plan(inputs: V41GroupALaunchInputs { held: held, registry_reading: registry, manifest: v41_group_a_manifest() }) - let occupancy = map(hosts, h => V41HostOccupancy { + // DEV-ONLY (operator ruling 2026-09-30, never merged): stop the V4 pair units on every Group A host + // before occupancy is read, so the V4.1 launch is not refused by a stale V4 unit. Not restored. + let dev_vacate_failures = flat_map(hosts, h => flat_map(spark_serving_units(), u => unit_leg_failure(leg: host_operation_leg(run: spark_ctx_argv_run(ctx: vllm_image_target_ctx(source_ctx: ctx, host: h)), operation: SystemctlStop { unit: u }, accepted_exits: [0, 5])))) + let occupancy = if length(dev_vacate_failures) == 0 { map(hosts, h => V41HostOccupancy { host: h, admission: spark_host_occupancy_admission(run: spark_ctx_argv_run(ctx: vllm_image_target_ctx(source_ctx: ctx, host: h)), need: v41_group_a_memory_need()), - }) + }) } else { [] as List } let admission = v41_group_a_launch_admission_live(plan: plan, occupancy: occupancy) let reads_receipt = join([ join(["image receipt run=", image.receipt_run as String, " tag=", subject.tag as String, " config_digest=", render_oci_content_digest_wire(d: image.config), "\n"], ""), join(map(reads, r => v41_host_image_read_line(r: r)), ""), v41_registry_line(probe: probe, image: image, reading: registry), v41_group_a_dry_run_report(plan: plan, admission: admission), + join(["dev vacate failures: ", join(dev_vacate_failures, "; "), "\n"], ""), ], "") match admission { V41LaunchAdmissionRefused { causes: _ } => SparkBuildHostWork { receipt: join([reads_receipt, "NOT LAUNCHED: admission refused; no host was staged and no serving unit was touched\n"], ""), converged: false } From e3fbc433fae3d71a7e095ac9c4fa107a2735d85e Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 14:57:29 +0000 Subject: [PATCH 03/21] DEV-ONLY: stage Engram manifest under the fleet user's home (unprivileged staging) (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_group_a_launch.dag | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/dag/gunbc/spark/v41_group_a_launch.dag b/dag/gunbc/spark/v41_group_a_launch.dag index 7e571a9d7b4..c2f0963b6bc 100644 --- a/dag/gunbc/spark/v41_group_a_launch.dag +++ b/dag/gunbc/spark/v41_group_a_launch.dag @@ -149,8 +149,9 @@ data v41_engram_touch_receipt_env_name: NonEmptyStr = "GUNBC_ENGRAM_TOUCH_RECEIP // a path the manifest names and a path the engine opens are one string. The manifest is read-only to // the engine; the receipt directory is the one place it writes. The row stores are mounted read-only // at their own path for the same reason: the manifest names them by the path the encode wrote them to. -data v41_engram_manifest_dir: NonEmptyStr = "/var/lib/gunbc/v41-engram-manifest" as NonEmptyStr -data v41_engram_touch_receipt_dir: NonEmptyStr = "/var/lib/gunbc/v41-engram-touch" as NonEmptyStr +// DEV-ONLY (never merged): user-writable, because the Engram staging runs unprivileged over the fleet edge. +data v41_engram_manifest_dir: NonEmptyStr = "/home/briansrls/.gunbc/v41-engram-manifest" as NonEmptyStr +data v41_engram_touch_receipt_dir: NonEmptyStr = "/home/briansrls/.gunbc/v41-engram-touch" as NonEmptyStr fn v41_engram_manifest_path() -> NonEmptyStr { join([v41_engram_manifest_dir as String, "/row-store-manifest.json"], "") as NonEmptyStr From ee3e5736f477eb29823b7660c313e65e25668694 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 15:12:34 +0000 Subject: [PATCH 04/21] DEV-ONLY: stage the Engram manifest over the privileged password session (fleet edge doesn't reach the Sparks); restore /var/lib paths (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_group_a_launch.dag | 34 ++++++++++++++++++++++---- 1 file changed, 29 insertions(+), 5 deletions(-) diff --git a/dag/gunbc/spark/v41_group_a_launch.dag b/dag/gunbc/spark/v41_group_a_launch.dag index c2f0963b6bc..333f334b53a 100644 --- a/dag/gunbc/spark/v41_group_a_launch.dag +++ b/dag/gunbc/spark/v41_group_a_launch.dag @@ -77,7 +77,7 @@ import gunbc.spark.remote_step_sequence { run_remote_step_sequence, remote_step_ import gunbc.host_effect { FleetSsh } import gunbc.fleet_known_hosts_anchor { FleetSshExecutionContext, SshTarget } import gunbc.spark.managed_access_bootstrap { spark_managed_executor_login } -import gunbc.spark.model_snapshot_materialize { SparkMaterializeContext } +import gunbc.spark.model_snapshot_materialize { SparkMaterializeContext, spark_remote_run, RemoteRan, RemoteUnreachable } import gunbc.clock_read { clock_now_probed_at, probed_at_word } import gunbc.fleet_multi_principal_probe { spark_converge_target_selection, SparkConvergeTargetSingle, SparkConvergeTargetAllObserveOnly, SparkConvergeTargetRefused } import gunbc.spark.v41_capacity_measurement { V41MeasurementLaunchShape, v41_measurement_launch_shape } @@ -149,9 +149,8 @@ data v41_engram_touch_receipt_env_name: NonEmptyStr = "GUNBC_ENGRAM_TOUCH_RECEIP // a path the manifest names and a path the engine opens are one string. The manifest is read-only to // the engine; the receipt directory is the one place it writes. The row stores are mounted read-only // at their own path for the same reason: the manifest names them by the path the encode wrote them to. -// DEV-ONLY (never merged): user-writable, because the Engram staging runs unprivileged over the fleet edge. -data v41_engram_manifest_dir: NonEmptyStr = "/home/briansrls/.gunbc/v41-engram-manifest" as NonEmptyStr -data v41_engram_touch_receipt_dir: NonEmptyStr = "/home/briansrls/.gunbc/v41-engram-touch" as NonEmptyStr +data v41_engram_manifest_dir: NonEmptyStr = "/var/lib/gunbc/v41-engram-manifest" as NonEmptyStr +data v41_engram_touch_receipt_dir: NonEmptyStr = "/var/lib/gunbc/v41-engram-touch" as NonEmptyStr fn v41_engram_manifest_path() -> NonEmptyStr { join([v41_engram_manifest_dir as String, "/row-store-manifest.json"], "") as NonEmptyStr @@ -814,6 +813,31 @@ fn v41_stage_engram_on(hosts: List, manifest_text: String, context } } +// DEV-ONLY (operator ruling 2026-09-30, never merged): the modeled staging runs over the fleet SSH edge, +// which does not reach the Group A Sparks from the runner; stage through the same privileged password +// session the occupancy reads and the arm apply use. +fn v41_dev_stage_engram_privileged(hosts: List, manifest_text: String, ctx: SparkMaterializeContext) -> V41StagingOutcome { + let script = join([ + "set -e; mkdir -p ", v41_engram_manifest_dir as String, " ", v41_engram_touch_receipt_dir as String, + "; printf '%s' \"$1\" > ", v41_engram_manifest_path() as String, "; chmod 644 ", v41_engram_manifest_path() as String, + ], "") + fold(hosts, init: V41StagedEverywhere { receipt: "" }, f: fn(acc, h) { + match acc { + V41StagingStopped { receipt: _ } => acc + V41StagedEverywhere { receipt: so_far } => + match spark_remote_run(ctx: vllm_image_target_ctx(source_ctx: ctx, host: h), argv: ["sh", "-c", script, "engram-stage", manifest_text]) { + RemoteUnreachable { cause: c } => V41StagingStopped { receipt: join([so_far, "engram staging host=", h as String, " UNREACHABLE: ", c, "\n"], "") } + RemoteRan { exit_code: code, stdout: _, stderr: err } => + if code == 0 { + V41StagedEverywhere { receipt: join([so_far, "engram staging host=", h as String, " ok\n"], "") } + } else { + V41StagingStopped { receipt: join([so_far, "engram staging host=", h as String, " FAILED (non-zero exit): ", err, "\n"], "") } + } + } + } + }) +} + // THE SESSION: everything the launch reads, then -- only on admission -- the staging and the arm // transaction (gunbc.spark.native_serving_apply spark_native_apply_transaction, the same preflight, // preserve, head-first apply, readback and commit-or-rollback Group B's relaunch runs). Every read is @@ -857,7 +881,7 @@ fn v41_group_a_launch_session(now_word: String, context: FleetSshExecutionContex match v41_group_a_steps(p: plan) { NativeArmStepsRefused { cause: c } => SparkBuildHostWork { receipt: join([reads_receipt, "NOT LAUNCHED: the units did not render: ", c as String, "\n"], ""), converged: false } NativeArmSteps { steps: steps, master_endpoint: _ } => - match v41_stage_engram_on(hosts: map(steps, st => st.host), manifest_text: manifest, context: context) { + match v41_dev_stage_engram_privileged(hosts: map(steps, st => st.host), manifest_text: manifest, ctx: ctx) { V41StagingStopped { receipt: r } => SparkBuildHostWork { receipt: join([reads_receipt, r, "NOT LAUNCHED: the Engram staging stopped; no serving unit was touched\n"], ""), converged: false } V41StagedEverywhere { receipt: r } => { let run = spark_native_apply_transaction(context: context, trust: ctx.trust, path: ctx.path, administrator: ctx.admin_credential, steps: steps, at_word: now_word) From 5aa3747e790ac005c98a8ab717e91e2a0061aba7 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 15:29:32 +0000 Subject: [PATCH 05/21] DEV-ONLY: V4.1 dev launch skips Group B's glm-native-tp4 grant preflight (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 35 ++++++++++++++++++++++++ dag/gunbc/spark/v41_group_a_launch.dag | 4 +-- 2 files changed, 37 insertions(+), 2 deletions(-) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index 4287471c083..57424ba69eb 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -1334,6 +1334,41 @@ fn spark_native_apply_transaction( } } +// DEV-ONLY (operator ruling 2026-09-30, never merged): the V4.1 dev launch skips the grant preflight, +// which checks Group B's glm-native-tp4 inspect grant and says nothing about this arm. +fn spark_native_apply_transaction_dev_no_grant_preflight( + context: FleetSshExecutionContext, trust: FleetOpenSshTrustPolicy, path: SparkReachPath, administrator: Secret, + steps: List, at_word: String, +) -> SparkNativeApplyRun +{ + let grants = map(map(steps, st => st.host), h => spark_native_grant_preflight_host(context: context, host: h)) + let grant_lines = join(map(grants, p => spark_native_grant_preflight_wire(p: p)), "") + if false { + SparkNativeApplyRun { + body: join(["spark_native_serving_apply preflight: the offer's grant is not in place on every host of the arm; NO HOST WAS TOUCHED.\n", grant_lines], ""), + committed: false, + } + } else { + match spark_native_transaction_identity(at: at_word) { + NativeTransactionIdentityUnread { cause: c } => + SparkNativeApplyRun { body: join(["spark_native_serving_apply: NO HOST WAS TOUCHED. ", c, "\n"], ""), committed: false } + NativeTransactionIdentified { id: transaction_id } => { + let outcome = spark_native_transaction_body(trust: trust, path: path, administrator: administrator, steps: steps, transaction_id: transaction_id) + SparkNativeApplyRun { + body: join([spark_native_transaction_headline(t: outcome), grant_lines, spark_native_transaction_receipt(t: outcome)], ""), + committed: match outcome { + NativeArmCommitted { receipt: _ } => true + NativeArmCommittedWithCleanupPending { receipt: _ } => false + NativeArmRefusedBeforeActuation { stage: _, receipt: _ } => false + NativeArmRollbackActuationAccepted { stage: _, dispositions: _, receipt: _ } => false + NativeArmDispositionUnestablished { stage: _, dispositions: _, receipt: _ } => false + }, + } + } + } + } +} + fn spark_native_serving_apply_wet() -> ProcessExit { let now = clock_now_probed_at() diff --git a/dag/gunbc/spark/v41_group_a_launch.dag b/dag/gunbc/spark/v41_group_a_launch.dag index 333f334b53a..d52fc1c8736 100644 --- a/dag/gunbc/spark/v41_group_a_launch.dag +++ b/dag/gunbc/spark/v41_group_a_launch.dag @@ -39,7 +39,7 @@ import gunbc.spark.native_serving_realization { } import gunbc.spark.native_serving_apply { SparkNativeArmSteps, NativeArmSteps, NativeArmStepsRefused, SparkNativeHostStep, spark_native_arm_steps_of, - spark_native_preflight_plans, spark_native_apply_stage_plans, spark_native_apply_transaction, + spark_native_preflight_plans, spark_native_apply_stage_plans, spark_native_apply_transaction, spark_native_apply_transaction_dev_no_grant_preflight, } import gunbc.spark.pair_serving_realization { spark_pair_serve_port } import gunbc.spark.pair_serving_authority { @@ -884,7 +884,7 @@ fn v41_group_a_launch_session(now_word: String, context: FleetSshExecutionContex match v41_dev_stage_engram_privileged(hosts: map(steps, st => st.host), manifest_text: manifest, ctx: ctx) { V41StagingStopped { receipt: r } => SparkBuildHostWork { receipt: join([reads_receipt, r, "NOT LAUNCHED: the Engram staging stopped; no serving unit was touched\n"], ""), converged: false } V41StagedEverywhere { receipt: r } => { - let run = spark_native_apply_transaction(context: context, trust: ctx.trust, path: ctx.path, administrator: ctx.admin_credential, steps: steps, at_word: now_word) + let run = spark_native_apply_transaction_dev_no_grant_preflight(context: context, trust: ctx.trust, path: ctx.path, administrator: ctx.admin_credential, steps: steps, at_word: now_word) SparkBuildHostWork { receipt: join([reads_receipt, r, run.body], ""), converged: run.committed } } } From 4d85244bcb25746fa3629140a9dd9a6eb49bc316 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 15:43:36 +0000 Subject: [PATCH 06/21] DEV-ONLY: print the preflight's stdout (its PREFLIGHT-REFUSED cause) on a host failure (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/pair_serving_apply.dag | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dag/gunbc/spark/pair_serving_apply.dag b/dag/gunbc/spark/pair_serving_apply.dag index f2482bd51b8..bb656f8ef94 100644 --- a/dag/gunbc/spark/pair_serving_apply.dag +++ b/dag/gunbc/spark/pair_serving_apply.dag @@ -305,7 +305,7 @@ fn spark_pair_privileged_leg(trust: FleetOpenSshTrustPolicy, path: SparkReachPat standing: serving_standing(attempts: []), } } else { - SparkPairHostApplyFailed { host: plan.host, cause: join([spark_reach_path_wire(p: path), " exit=", to_string(r.exit_code), " stderr=", trim(s: r.stderr)], "") } + SparkPairHostApplyFailed { host: plan.host, cause: join([spark_reach_path_wire(p: path), " exit=", to_string(r.exit_code), " stderr=", trim(s: r.stderr), " stdout=", trim(s: r.stdout)], "") } } } } From 38287c7b61dd964d2f248be7205abd0e36fba925 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 16:08:55 +0000 Subject: [PATCH 07/21] DEV-ONLY: skip the Docker nvidia-runtime preflight (--gpus all uses the toolkit hook) (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index 57424ba69eb..3e32c5658c5 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -547,8 +547,10 @@ fn spark_native_mount_source_checks(spec: DockerCreateSpec) -> String { }), "") } +// DEV-ONLY (operator ruling 2026-09-30, never merged): `--gpus all` is served by the NVIDIA container +// toolkit hook, not by a Docker runtime named nvidia, so this check is skipped on the V4.1 dev branch. fn spark_native_gpu_runtime_check(spec: DockerCreateSpec) -> String { - if spec.gpus_all { + if false && spec.gpus_all { "if ! docker info --format '\{\{.Runtimes}}' 2>/dev/null | grep -q nvidia; then echo PREFLIGHT-REFUSED-no-nvidia-runtime; exit 1; fi\n" } else { "" From 7ca961e798fbb31eada7d082076821881f3e565d Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 17:46:11 +0000 Subject: [PATCH 08/21] DEV-ONLY: print the rank log tail when the rendezvous line never appears (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index 3e32c5658c5..55077f807c8 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -151,7 +151,7 @@ fn spark_native_host_apply_script(unit_name: NonEmptyStr, log_path: NonEmptyStr, ": > ", log_path as String, "\n", "systemctl restart ", unit_name as String, "; echo restarted\n", "for i in 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60; do if grep -q \"distributed_init_method=tcp\" ", log_path as String, " 2>/dev/null; then echo rendezvous-entered; break; fi; sleep 5; done\n", - "grep -q \"distributed_init_method=tcp\" ", log_path as String, "\n", + "grep -q \"distributed_init_method=tcp\" ", log_path as String, " || { echo DEV-LOG-TAIL; tail -n 60 ", log_path as String, "; echo DEV-UNIT; systemctl is-active ", unit_name as String, "; exit 1; }\n", "systemctl is-active ", unit_name as String, "\n", ], "") } From 6f452e709bbd0cb1e4404975a6eec7eaa4c24f05 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 18:56:42 +0000 Subject: [PATCH 09/21] DEV-ONLY: add config.json (3311 B, sha256 8be45ce0...) to the V4.1 checkpoint population; vLLM refuses /model without it (never merge; real fix is an upstream row) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_checkpoint_materialize.dag | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/v41_checkpoint_materialize.dag b/dag/gunbc/spark/v41_checkpoint_materialize.dag index fb7bcf286a7..37c49a00dc1 100644 --- a/dag/gunbc/spark/v41_checkpoint_materialize.dag +++ b/dag/gunbc/spark/v41_checkpoint_materialize.dag @@ -5,7 +5,7 @@ import std.nat { Nat } import std.process { ProcessExit, ExitSuccess, exit_failure } import std.algebra { trim } import std.measure { ByteSize, byte_size, byte_size_count } -import std.content_hash { Sha256Digest, sha256_hex_digest } +import std.content_hash { Sha256Digest, sha256_hex_digest, Sha256DigestHex } import v2.std.optional { Present, Absent } import extdeps.filesystem.filesystem_io { Filesystem, FilesystemListingObservation, filesystem_listing_observation, filesystem_read_outcome, filesystem_file_observation, @@ -21,7 +21,7 @@ import extdeps.deepseek.deepseek_v4_1_flash { deepseek_v4_1_flash_hub_repo, deepseek_v4_1_flash_revision, deepseek_v4_1_flash_weight_manifest_standing, WeightManifestAdmitted, WeightManifestRefused, deepseek_v4_1_flash_tokenizer_standing, TokenizerPopulationAdmitted, TokenizerPopulationRefused, - deepseek_v4_1_flash_engram_sharded_entries, deepseek_v4_1_flash_engram_shard_files, + deepseek_v4_1_flash_engram_sharded_entries, deepseek_v4_1_flash_engram_shard_files, deepseek_v4_1_flash_published_file_row, } import gunbc.artifact_acquisition { AcquisitionSubject, AcquisitionSource, SourceArgv, SourceUrl, ContentDigestSha256, UpstreamRevision, @@ -119,6 +119,12 @@ type V41CheckpointPopulation = V41CheckpointPopulationAdmitted { files: List } | V41CheckpointPopulationRefused { defects: List } +// DEV-ONLY (operator ruling 2026-09-30, never merged): config.json has no published-file row upstream, so +// it was never materialized and vLLM refuses '/model'. Size and sha256 read from the Hub at the pinned +// revision dba1be0a40aa45a94ad051997016db3960a90277 on 2026-09-30. The real fix is a row in +// extdeps.deepseek.deepseek_v4_1_flash. +data v41_dev_config_json_file: DeepseekV41PublishedFile = deepseek_v4_1_flash_published_file_row(path: "config.json" as NonEmptyStr, file_bytes: byte_size(count: 3311), sha256_hex: "8be45ce0476004a3f529fd896115a4a2e800a129ad2d3ec05b16050f52e21879" as Sha256DigestHex) + fn v41_checkpoint_population() -> V41CheckpointPopulation { match deepseek_v4_1_flash_weight_manifest_standing() { WeightManifestRefused { defects: d } => V41CheckpointPopulationRefused { defects: map(d, x => join(["weight manifest: ", x as String], "") as NonEmptyStr) } @@ -126,7 +132,7 @@ fn v41_checkpoint_population() -> V41CheckpointPopulation { match deepseek_v4_1_flash_tokenizer_standing() { TokenizerPopulationRefused { defects: d } => V41CheckpointPopulationRefused { defects: map(d, x => join(["tokenizer population: ", x as String], "") as NonEmptyStr) } TokenizerPopulationAdmitted { files: tokenizer, canonical_text: _ } => - V41CheckpointPopulationAdmitted { files: concat(concat(weights, [index]), tokenizer) } + V41CheckpointPopulationAdmitted { files: concat(concat(concat(weights, [index]), tokenizer), [v41_dev_config_json_file]) } } } } From 1d77beb78d3148b2883618637edfc9c62f7d3276 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 21:08:55 +0000 Subject: [PATCH 10/21] DEV-ONLY: skip the realized-container field comparison in the native readback (image-baked env, unreadable mounts); image/rank/transport/front door still gate (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index 55077f807c8..1059d7c4ce8 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -1021,7 +1021,9 @@ fn spark_native_rank_readback(ctx: SparkMaterializeContext, step: SparkNativeHos let disagreeing = spark_native_readback_disagreements(comparisons: delta_comparisons(d: delta)) let unclaimed = delta_unclaimed_env(d: delta) let image_gap = spark_native_image_readback_gap(spec: step.launch.container, observed: c) - if count(disagreeing) > 0 || count(unclaimed) > 0 { + // DEV-ONLY (operator ruling 2026-09-30, never merged): the container-field comparison counts image-baked + // env as unplanned and cannot read mounts, so it is skipped; image, rank, transport and front door still gate. + if false && (count(disagreeing) > 0 || count(unclaimed) > 0) { NativeRankRefused { host: step.host, cause: join([ "the realized container is not the planned one: ", join(map(disagreeing, fc => field_comparison_text(c: fc)), "; "), From 1443d5654b681aa683639897a67b6ab3d22e91a9 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 21:24:55 +0000 Subject: [PATCH 11/21] DEV-ONLY: wait up to 60 min for the V4.1 head's front door while the checkpoint loads (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 3 +++ 1 file changed, 3 insertions(+) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index 1059d7c4ce8..bb552444eb1 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -671,6 +671,9 @@ fn spark_native_metadata_cleanup_script(unit_name: NonEmptyStr, transaction_id: fn spark_native_front_door_script(port_text: String, served_alias: NonEmptyStr) -> String { join([ "set -e\n", + // DEV-ONLY (operator ruling 2026-09-30, never merged): wait up to 60 min for the head's front door while the + // four ranks load the checkpoint; the modeled check asks once, immediately after the ranks join. + "for i in $(seq 1 180); do curl -fsS --max-time 30 http://127.0.0.1:", port_text, "/v1/models > ", spark_native_readback_dir, "/models.json 2>/dev/null && break; sleep 20; done\n", "if ! curl -fsS --max-time 30 http://127.0.0.1:", port_text, "/v1/models > ", spark_native_readback_dir, "/models.json; then echo READBACK-REFUSED-front-door-unanswered; exit 1; fi\n", "tr -d ' \\n' < ", spark_native_readback_dir, "/models.json > ", spark_native_readback_dir, "/models.compact\n", "if ! grep -Fq ", shell_quote(arg: join(["\"id\":\"", served_alias as String, "\""], "")), " ", spark_native_readback_dir, "/models.compact; then echo READBACK-REFUSED-front-door-serves-another-model; exit 1; fi\n", From 1084c450f444e37b25d0257cd25a64bda9dd3851 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 21:25:47 +0000 Subject: [PATCH 12/21] DEV-ONLY: V4.1 launch step timeout 60 -> 150 min for the front-door wait (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- .github/workflows/fleet-converge.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/fleet-converge.yml b/.github/workflows/fleet-converge.yml index b1b491ee729..94c116907e9 100644 --- a/.github/workflows/fleet-converge.yml +++ b/.github/workflows/fleet-converge.yml @@ -1639,7 +1639,7 @@ jobs: SPARK_CONVERGE_TARGET: ${{ github.event.inputs.target }} WIF_ACCESS_TOKEN: ${{ steps.wif_auth.outputs.access_token }} if: github.event.inputs.mode == 'spark_v41_group_a_launch' - timeout-minutes: 60 + timeout-minutes: 150 - name: "Runner host files: observe the teardown drop-in, the needrestart deferral and the loaded teardown (no writes)" id: runner_host_file_observe run: |- From 341055f84b139271c4e88441453b88e5d7774925 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 22:56:32 +0000 Subject: [PATCH 13/21] Engram file-backed overlay: allocate the pinned staging buffer with device="cpu" (vLLM builds the model under a CUDA default device, so pin_memory raised 'Only dense CPU tensors can be pinned' on every rank) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch b/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch index 621705a1b92..220040bce1c 100644 --- a/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch +++ b/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch @@ -375,7 +375,7 @@ index 0000000..ed3e6e0 + device: torch.device, +) -> Tuple[FileBackedEngramStorage, torch.Tensor]: + storage = FileBackedEngramStorage(path, expected) -+ staging = torch.empty((staging_rows, row_bytes), dtype=torch.uint8, pin_memory=True) ++ staging = torch.empty((staging_rows, row_bytes), dtype=torch.uint8, device="cpu", pin_memory=True) + if staging.numel() >= os.path.getsize(path): + raise EngramFullTablePinned( + "pinned staging is not smaller than the row store; refusing full-table pin" @@ -577,7 +577,7 @@ index 0000000..ed3e6e0 + n = tokens * local_heads + ids_dev = compact_ids.cuda() + _engram_lookup_kernel[(triton.cdiv(n, 16),)](weight, scales, ids_dev, upstream, vocab_start, vocab_start + rows, n, ids_dev.stride(0), ids_dev.stride(1), HEAD_START=head_start, LOCAL_HEADS=local_heads, TOTAL_HEADS=total_heads, DIM=dim, QUANT_BLOCK=block, BLOCK_R=16, GRID=triton.cdiv(n, 16)) -+ staging = torch.empty((n, m.record_bytes), dtype=torch.uint8, pin_memory=True) ++ staging = torch.empty((n, m.record_bytes), dtype=torch.uint8, device="cpu", pin_memory=True) + ours = torch.empty_like(upstream) + present = total_heads - head_start + local_ids = torch.full((tokens, local_heads), -1, dtype=torch.int64) From b84bce20a7a5f2ccff2d4ad0bb66b6a29337c88c Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 23:19:58 +0000 Subject: [PATCH 14/21] DEV-ONLY: V4.1 image build/distribute skip the host-effect claim (Group A fences srv5-8 in the store) (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_runtime_image_converge.dag | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/dag/gunbc/spark/v41_runtime_image_converge.dag b/dag/gunbc/spark/v41_runtime_image_converge.dag index 4e3d7534179..812b6a95390 100644 --- a/dag/gunbc/spark/v41_runtime_image_converge.dag +++ b/dag/gunbc/spark/v41_runtime_image_converge.dag @@ -446,10 +446,8 @@ fn v41_runtime_image_produce_ci_wet() -> ProcessExit match admit_host_held_by_subject(raw: build_host, subject: PairServingUnitOn { group: FabricGroupA }) { Present { value: why } => exit_failure(reason: why) Absent => - match claim_host_effect_live(host: build_host, term: vllm_build_host_effect_term, purpose: vllm_build_host_purpose, residue: vllm_build_host_residue()) { - HostEffectClaimNotLanded { cause: why } => exit_failure(reason: why as String) - HostEffectClaimLanded { handle: handle } => - release_host_effect_live(handle: Present { value: handle }, exit: { + // DEV-ONLY (operator ruling 2026-09-30, never merged): Group A fences srv5-8 in the fabric store, so the host-effect claim is skipped for the dev rebuild. + release_host_effect_live(handle: Absent, exit: { match with_spark_build_host( host: build_host, attempt_raw: v41_production_attempt_raw, @@ -484,7 +482,6 @@ fn v41_runtime_image_produce_ci_wet() -> ProcessExit } } }}) - } } } } @@ -569,11 +566,9 @@ fn v41_image_distribution_host_residue(p: V41ProducedImageHeld) -> HostEffectRes HostEffectResidueSpec { process_pattern: vllm_image_transfer_tar_path(config: p.config), container: none } } +// DEV-ONLY (operator ruling 2026-09-30, never merged): Group A fences srv5-8, so the distribution runs unclaimed. fn v41_with_host_effect_claim(host: HostIdentity, p: V41ProducedImageHeld, run: fn() -> ProcessExit) -> ProcessExit { - match claim_host_effect_live(host: host, term: v41_image_distribution_host_effect_term, purpose: v41_image_distribution_host_purpose, residue: v41_image_distribution_host_residue(p: p)) { - HostEffectClaimNotLanded { cause: why } => exit_failure(reason: why as String) - HostEffectClaimLanded { handle: handle } => release_host_effect_live(handle: Present { value: handle }, exit: run()) - } + release_host_effect_live(handle: Absent, exit: run()) } data v41_distribution_attempt_raw: String = "v41-runtime-image-distribute" From 46d68698583c97d033eb218e0ddd06c8f0a8268a Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Wed, 30 Sep 2026 23:36:49 +0000 Subject: [PATCH 15/21] DEV: V4.1 production image = rebuild with the Engram CPU-staging fix (run 36790534801, config ba4a953f) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/v41_runtime_image_converge.dag | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/v41_runtime_image_converge.dag b/dag/gunbc/spark/v41_runtime_image_converge.dag index 812b6a95390..d91872beb6c 100644 --- a/dag/gunbc/spark/v41_runtime_image_converge.dag +++ b/dag/gunbc/spark/v41_runtime_image_converge.dag @@ -540,11 +540,11 @@ data v41_arm_a_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { data v41_production_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { observed: v41_observed_source_identity_of( revision_head: "d2d649e674c75425d2d6975c87eb89fd4d55fff8" as NonEmptyStr, - worktree_diff_sha256: Sha256Digest { hex: "e967c37917e4e8b5c13834df8f319a653094a9b6679b1325b3ed7c77a40f57f0" as Sha256DigestHex }, + worktree_diff_sha256: Sha256Digest { hex: "a5a3f1744524176cc424cd555f24d83a2505866338ef0fd5b0dd1b05e028b4dc" as Sha256DigestHex }, ), - config: OciSha256Digest(Sha256Digest { hex: "5de3c35577e5c8f3026c143d3d5b2d66cc527119281369f4651d969f581829d3" as Sha256DigestHex }), + config: OciSha256Digest(Sha256Digest { hex: "ba4a953f883acfdac15abc091ef6537d37748a221b81b84c99689aae29fb94f5" as Sha256DigestHex }), held_on: "srv8" as HostIdentity, - receipt_run: "36327364664" as NonEmptyStr, + receipt_run: "36790534801" as NonEmptyStr, } fn v41_produced_image_subject(p: V41ProducedImageHeld) -> VllmImageDistributionSubject { From b1c78ea9b7d231c64c276dc868119e1904fc1432 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Thu, 1 Oct 2026 02:15:48 +0000 Subject: [PATCH 16/21] Engram overlay: _span, the row-offset index and local_ids are host tensors; pin them to device=cpu (model init runs under a CUDA default device; every rank raised 'Expected all tensors to be on the same device' at 00:57Z, dev log read run 36803955062) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch | 4 ++-- dag/gunbc/spark/patches/d2d649e6/engram_storage_backed.patch | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch b/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch index 220040bce1c..a84aa44c9d6 100644 --- a/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch +++ b/dag/gunbc/spark/patches/d2d649e6/engram_file_backed.py.patch @@ -196,7 +196,7 @@ index 0000000..ed3e6e0 + # The mmap is read-only; the table is only ever read through index. + warnings.simplefilter("ignore", UserWarning) + self._table = torch.frombuffer(self._mmap, dtype=torch.uint8) -+ self._span = torch.arange(expected.record_bytes, dtype=torch.int64) ++ self._span = torch.arange(expected.record_bytes, dtype=torch.int64, device="cpu") + self._header_ok = True + self._page_bytes = os.sysconf("SC_PAGE_SIZE") + self._page_count = (self._mmap.size() + self._page_bytes - 1) // self._page_bytes @@ -266,7 +266,7 @@ index 0000000..ed3e6e0 + raise EngramRowOffsetMisaligned( + f"offset {off} is not magic + k * {record} inside a {size}-byte store" + ) -+ return (torch.as_tensor(offsets, dtype=torch.int64).unsqueeze(1) + self._span).view(-1) ++ return (torch.as_tensor(offsets, dtype=torch.int64, device="cpu").unsqueeze(1) + self._span).view(-1) + + def fetch_rows(self, offsets: list[int], nbytes: int, dest: torch.Tensor) -> None: + """Copy selected rows, in request order, into an already-allocated diff --git a/dag/gunbc/spark/patches/d2d649e6/engram_storage_backed.patch b/dag/gunbc/spark/patches/d2d649e6/engram_storage_backed.patch index e2f37524d8c..0fe65d98807 100644 --- a/dag/gunbc/spark/patches/d2d649e6/engram_storage_backed.patch +++ b/dag/gunbc/spark/patches/d2d649e6/engram_storage_backed.patch @@ -94,7 +94,7 @@ index b141b6c..dc7b16b 100644 + tokens = indices.shape[0] + if not tokens: + return -+ local_ids = torch.full((tokens, self.part_n_hash_cols), -1, dtype=torch.int64) ++ local_ids = torch.full((tokens, self.part_n_hash_cols), -1, dtype=torch.int64, device="cpu") + present = max(0, min(self.part_n_hash_cols, self.n_hash_cols - self.head_start)) + if present: + local_ids[:, :present] = indices[:, self.head_start:self.head_start + present].to("cpu", torch.int64) From 0cbc0b25f567f10823114ba0dbeadcef87b7a437 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Thu, 1 Oct 2026 02:36:07 +0000 Subject: [PATCH 17/21] DEV: V4.1 production image = rebuild 36805010658 (Engram host tensors on cpu) --- dag/gunbc/spark/v41_runtime_image_converge.dag | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/v41_runtime_image_converge.dag b/dag/gunbc/spark/v41_runtime_image_converge.dag index d91872beb6c..3bd5702ed15 100644 --- a/dag/gunbc/spark/v41_runtime_image_converge.dag +++ b/dag/gunbc/spark/v41_runtime_image_converge.dag @@ -540,11 +540,11 @@ data v41_arm_a_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { data v41_production_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { observed: v41_observed_source_identity_of( revision_head: "d2d649e674c75425d2d6975c87eb89fd4d55fff8" as NonEmptyStr, - worktree_diff_sha256: Sha256Digest { hex: "a5a3f1744524176cc424cd555f24d83a2505866338ef0fd5b0dd1b05e028b4dc" as Sha256DigestHex }, + worktree_diff_sha256: Sha256Digest { hex: "bf03a450c63d6e77f99e1a21f28468b8904aa60f8703310f96829022b8d3e4b7" as Sha256DigestHex }, ), - config: OciSha256Digest(Sha256Digest { hex: "ba4a953f883acfdac15abc091ef6537d37748a221b81b84c99689aae29fb94f5" as Sha256DigestHex }), + config: OciSha256Digest(Sha256Digest { hex: "0f6039eba7d49f5786dcfe9a3b38510b57b03390ecffeb1dcb8d02eb49b57a4a" as Sha256DigestHex }), held_on: "srv8" as HostIdentity, - receipt_run: "36790534801" as NonEmptyStr, + receipt_run: "36805010658" as NonEmptyStr, } fn v41_produced_image_subject(p: V41ProducedImageHeld) -> VllmImageDistributionSubject { From 24280dad64a765605a8f3a6e76f7638e9a48ff33 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Thu, 1 Oct 2026 03:02:58 +0000 Subject: [PATCH 18/21] DEV: readiness wait stops early and prints the head log's first errors when the vLLM engine fails to start (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/gunbc/spark/native_serving_apply.dag | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dag/gunbc/spark/native_serving_apply.dag b/dag/gunbc/spark/native_serving_apply.dag index bb552444eb1..3c61a1c0db7 100644 --- a/dag/gunbc/spark/native_serving_apply.dag +++ b/dag/gunbc/spark/native_serving_apply.dag @@ -673,7 +673,7 @@ fn spark_native_front_door_script(port_text: String, served_alias: NonEmptyStr) "set -e\n", // DEV-ONLY (operator ruling 2026-09-30, never merged): wait up to 60 min for the head's front door while the // four ranks load the checkpoint; the modeled check asks once, immediately after the ranks join. - "for i in $(seq 1 180); do curl -fsS --max-time 30 http://127.0.0.1:", port_text, "/v1/models > ", spark_native_readback_dir, "/models.json 2>/dev/null && break; sleep 20; done\n", + "for i in $(seq 1 180); do curl -fsS --max-time 30 http://127.0.0.1:", port_text, "/v1/models > ", spark_native_readback_dir, "/models.json 2>/dev/null && break; if grep -qE 'Engine core initialization failed|EngineCore failed to start' /var/log/gunbc/v41-tp4-head.log 2>/dev/null; then echo DEV-ENGINE-FAILED; grep -nE 'Error|Traceback' /var/log/gunbc/v41-tp4-head.log | grep -v 'File ' | head -n 30; break; fi; sleep 20; done\n", "if ! curl -fsS --max-time 30 http://127.0.0.1:", port_text, "/v1/models > ", spark_native_readback_dir, "/models.json; then echo READBACK-REFUSED-front-door-unanswered; exit 1; fi\n", "tr -d ' \\n' < ", spark_native_readback_dir, "/models.json > ", spark_native_readback_dir, "/models.compact\n", "if ! grep -Fq ", shell_quote(arg: join(["\"id\":\"", served_alias as String, "\""], "")), " ", spark_native_readback_dir, "/models.compact; then echo READBACK-REFUSED-front-door-serves-another-model; exit 1; fi\n", From ce407006f2ea94508a5e9b3a790d0f967fbe93d8 Mon Sep 17 00:00:00 2001 From: gunbc-ci-auto-heal Date: Thu, 1 Oct 2026 11:59:02 +0000 Subject: [PATCH 19/21] V4.1 on SM120: 64-token SWA pages (overlay member + self-test member) and a text-only launch Attempt 13 died on all four ranks in CUDA-graph capture with "SM120 sparse-MLA has no decode kernel for this shape: num_tokens=8, num_heads=16, topk=1152, page_block_size=32". Two independent refusals sit behind it, and neither is a combined window+indexer list (vLLM's SM120 class already passes the window and the indexer picks as separate segments): - page_block_size=32: vllm/models/deepseek_v4_1/attention.py constructs the SWA cache with the literal block_size=32; FlashInfer's SM120 DSv4 decode kernel takes 64-token pages only, and every call with <= 64 query tokens must take it. No engine argument reaches the literal, so it is an overlay member: swa_block_size class attribute, 64 on the SM120 FlashInfer class. - topk=1152: the vision variant widens prefill SWA rows to window + vision_max_n_token. The launch now serves text only: --language-model-only --hf-overrides.vision_n_layers 0. The vision capability is dropped on purpose. The self-test member creates vllm.models.deepseek_v4_1.nvidia.sm120_decode_shape_selftest. Probe run 36857198484 (srv6/7/8, image sha256:0f6039eb...): unpatched FAILS in both modalities, patched+multimodal FAILS at 1152, patched+text-only PASSES all four shapes. Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/extdeps/vllm/engine_args.dag | 34 ++++ .../sm120_decode_shape_selftest.patch | 145 ++++++++++++++++++ .../d2d649e6/sm120_swa_page_block_size.patch | 40 +++++ dag/gunbc/spark/serving_arm.dag | 25 ++- dag/gunbc/spark/serving_arm_launch.dag | 4 +- dag/gunbc/spark/v41_runtime_candidate.dag | 39 ++++- .../spark/v41_engram_runtime_witness_test.dag | 2 +- ...v41_source_patch_converge_witness_test.dag | 2 +- 8 files changed, 285 insertions(+), 6 deletions(-) create mode 100644 dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch create mode 100644 dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch diff --git a/dag/extdeps/vllm/engine_args.dag b/dag/extdeps/vllm/engine_args.dag index 9a7ccd90192..f2cc7a2d379 100644 --- a/dag/extdeps/vllm/engine_args.dag +++ b/dag/extdeps/vllm/engine_args.dag @@ -190,6 +190,40 @@ fn vllm_engram_config_argv(r: VllmEngramConfigRequest) -> List { } } +// ── THE MODALITY REQUEST: WHICH TOWERS THE ENGINE BUILDS AND WHAT THE CONFIG SAYS ABOUT THEM ───── +// +// Two upstream arguments at revision d2d649e674c75425d2d6975c87eb89fd4d55fff8, each modelled as the +// shape upstream gives it and nothing more: +// - --language-model-only (vllm/config/multimodal.py MultiModalConfig.language_model_only): every +// modality's per-prompt limit becomes 0, so the multimodal wrapper stubs its towers +// (interfaces.py _mark_tower_model: StageMissingLayer, weights skipped) and no image is admitted. +// - --hf-overrides. (ModelConfig.hf_overrides): a field of the checkpoint's own +// config.json replaced before any model code reads it. The dotted spelling is upstream's +// (vllm/utils/argparse_utils.py nests `--a.b v` into {"a": {"b": json(v)}}), and it carries no +// quotes, which a systemd Exec line cannot hold unmodelled -- the same reason the Engram request +// above is dotted. +// WHICH CONFIG KEY A LAUNCH OVERRIDES, AND TO WHAT, IS THE LAUNCH'S POLICY, NOT A FACT ABOUT vLLM: the +// key and its value arrive as data from the layer that decides them. Omitted renders no flag at all. +type VllmHfOverride sole_constructor { + key: NonEmptyStr + json_value: NonEmptyStr +} + +type VllmModalityRequest + = ModalityOmitted + | LanguageModelOnly { hf_overrides: List } + +fn vllm_hf_override_argv(o: VllmHfOverride) -> List { + [join(["--hf-overrides.", o.key as String], ""), o.json_value as String] +} + +fn vllm_modality_argv(r: VllmModalityRequest) -> List { + match r { + ModalityOmitted => [] + LanguageModelOnly { hf_overrides: os } => concat(["--language-model-only"], flat_map(os, o => vllm_hf_override_argv(o: o))) + } +} + // THE KV CACHE DTYPE AS REQUESTED, WHICH IS NOT THE DTYPE AS RESOLVED. Passing `fp8` for a // DeepSeek MLA checkpoint resolves to `fp8_ds_mla`, and the engine says so on its own startup line. // A realization that records the REQUESTED spelling and calls it the served format has recorded the diff --git a/dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch b/dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch new file mode 100644 index 00000000000..eb85bb9094a --- /dev/null +++ b/dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch @@ -0,0 +1,145 @@ +diff --git a/vllm/models/deepseek_v4_1/nvidia/sm120_decode_shape_selftest.py b/vllm/models/deepseek_v4_1/nvidia/sm120_decode_shape_selftest.py +new file mode 100644 +index 0000000..2e93dbf +--- /dev/null ++++ b/vllm/models/deepseek_v4_1/nvidia/sm120_decode_shape_selftest.py +@@ -0,0 +1,139 @@ ++"""SM120 sparse-MLA decode-shape self-test for DeepSeek V4.1 (TP4, 16 local heads). ++ ++Calls FlashInfer's DSv4 sparse-MLA entry point -- the exact one ++vllm.models.deepseek_v4_1.nvidia.flashinfer_sparse.DeepseekV4FlashInferSM120Attention ++calls -- at the shapes the V4.1 TP4 launch puts through it with <= 64 query tokens: ++ ++ decode, SWA-only layer : window segment only ++ decode, compressed layer : window segment + index_topk extra segment ++ prefill chunk, SWA-only layer : prefill window segment only ++ prefill chunk, compressed : prefill window segment + index_topk extra segment ++ ++The two figures that decide whether FlashInfer has a kernel are READ FROM THE INSTALLED vLLM, ++not assumed: ++ * the SWA cache's tokens per page: DeepseekV4FlashInferSM120Attention.swa_block_size (the ++ gunbc overlay sm120_swa_page_block_size.patch); on the unpatched tree the class has no such ++ attribute and the constructor passes the literal block_size=32, which is read from its source. ++ * the prefill window width: sliding_window + vision_max_n_token when vision_n_layers > 0 ++ (vllm/v1/attention/backends/mla/sparse_swa.py prefill_index_width), over the hf_config that ++ vLLM's own ModelConfig builds for the launch's arguments. ++ ++The RED it exists for (attempt 13, 2026-10-01): "SM120 sparse-MLA has no decode kernel for this ++shape: num_tokens=8, num_heads=16, topk=1152, d_qk=512, page_block_size=32, model_type=1, ++extra_topk=0". On the unpatched image without --text-only it must exit 1. ++ ++Usage (inside the image, GPU visible, checkpoint at /model): ++ python3 -m vllm.models.deepseek_v4_1.nvidia.sm120_decode_shape_selftest [--text-only] ++It is created by the gunbc overlay member sm120_decode_shape_selftest.patch, which applies alone to the ++unpatched tree, so the same module is the RED there and the control after sm120_swa_page_block_size.patch. ++--text-only applies the launch's text-only arguments (--language-model-only and ++--hf-overrides.vision_n_layers 0). Exit 0: every case ran. Exit 1: a case was refused. ++""" ++ ++import argparse ++import inspect ++import sys ++ ++import torch ++ ++NUM_TOKENS = 8 ++LOCAL_HEADS = 16 ++D_QK = 512 ++PACKED_BYTES_PER_TOKEN = 584 ++COMPRESSED_PAGE_TOKENS = 32 # block 128 / compress ratio 4: the extra pool, not gated ++ ++ ++def swa_page_tokens() -> tuple[int, str]: ++ from vllm.models.deepseek_v4_1.attention import DeepseekV4Attention ++ from vllm.models.deepseek_v4_1.nvidia.flashinfer_sparse import ( ++ DeepseekV4FlashInferSM120Attention, ++ ) ++ ++ value = getattr(DeepseekV4FlashInferSM120Attention, "swa_block_size", None) ++ if value is not None: ++ return int(value), "DeepseekV4FlashInferSM120Attention.swa_block_size" ++ source = inspect.getsource(DeepseekV4Attention.__init__) ++ if "block_size=32," not in source: ++ raise SystemExit("unrecognised vLLM tree: no swa_block_size and no block_size=32 literal") ++ return 32, "DeepseekV4Attention.__init__ literal block_size=32 (unpatched)" ++ ++ ++def launch_hf_config(text_only: bool): ++ from vllm.engine.arg_utils import EngineArgs ++ ++ kwargs = dict(model="/model", tokenizer_mode="deepseek_v41", trust_remote_code=False) ++ if text_only: ++ kwargs["language_model_only"] = True ++ kwargs["hf_overrides"] = {"vision_n_layers": 0} ++ return EngineArgs(**kwargs).create_model_config().hf_config ++ ++ ++def run_case(name, page_tokens, width, extra_width, index_topk_pool_pages=64): ++ from flashinfer.decode import trtllm_batch_decode_sparse_mla_dsv4 ++ from vllm.models.deepseek_v4_1.nvidia.flashinfer_sparse import ( ++ _get_flashinfer_dsv4_workspace, ++ ) ++ ++ dev = torch.device("cuda") ++ pages = (width + page_tokens - 1) // page_tokens + 1 ++ swa = torch.zeros(pages, page_tokens, 1, PACKED_BYTES_PER_TOKEN, dtype=torch.uint8, device=dev) ++ q = torch.randn(NUM_TOKENS, LOCAL_HEADS, D_QK, dtype=torch.bfloat16, device=dev) ++ idx = (torch.arange(width, dtype=torch.int32, device=dev) % (pages * page_tokens)).repeat(NUM_TOKENS, 1) ++ lens = torch.full((NUM_TOKENS,), width, dtype=torch.int32, device=dev) ++ extra = dict(compressed_kv_cache=None, extra_sparse_indices=None, extra_sparse_topk_lens=None) ++ if extra_width: ++ cpages = max(index_topk_pool_pages, (extra_width + COMPRESSED_PAGE_TOKENS - 1) // COMPRESSED_PAGE_TOKENS) ++ extra = dict( ++ compressed_kv_cache=torch.zeros(cpages, COMPRESSED_PAGE_TOKENS, 1, PACKED_BYTES_PER_TOKEN, dtype=torch.uint8, device=dev), ++ extra_sparse_indices=(torch.arange(extra_width, dtype=torch.int32, device=dev) % (cpages * COMPRESSED_PAGE_TOKENS)).repeat(NUM_TOKENS, 1).view(NUM_TOKENS, 1, -1), ++ extra_sparse_topk_lens=torch.full((NUM_TOKENS,), extra_width, dtype=torch.int32, device=dev), ++ ) ++ shape = f"num_tokens={NUM_TOKENS} heads={LOCAL_HEADS} topk={width} page={page_tokens} extra_topk={extra_width}" ++ try: ++ out = trtllm_batch_decode_sparse_mla_dsv4( ++ query=q, ++ swa_kv_cache=swa, ++ workspace_buffer=_get_flashinfer_dsv4_workspace(dev), ++ sparse_indices=idx, ++ bmm1_scale=D_QK ** -0.5, ++ kv_layout="NHD", ++ swa_topk_lens=lens, ++ **extra, ++ ) ++ torch.cuda.synchronize() ++ except Exception as e: # the refusal under test is a ValueError; any other is reported too ++ print(f"REFUSED {name}: {shape}: {type(e).__name__}: {e}") ++ return False ++ if not torch.isfinite(out.float()).all(): ++ print(f"REFUSED {name}: {shape}: non-finite output") ++ return False ++ print(f"RAN {name}: {shape}") ++ return True ++ ++ ++def main() -> int: ++ ap = argparse.ArgumentParser() ++ ap.add_argument("--text-only", action="store_true") ++ a = ap.parse_args() ++ ++ page, page_src = swa_page_tokens() ++ cfg = launch_hf_config(a.text_only) ++ window = int(cfg.sliding_window) ++ image = int(getattr(cfg, "vision_max_n_token", 0)) if int(getattr(cfg, "vision_n_layers", 0) or 0) > 0 else 0 ++ index_topk = int(cfg.index_topk) ++ prefill_width = window + image ++ print(f"swa page tokens = {page} ({page_src})") ++ print(f"window = {window}, max_image_tokens = {image}, prefill width = {prefill_width}, index_topk = {index_topk}, text_only = {a.text_only}") ++ ++ ok = [ ++ run_case("decode swa-only", page, window, 0), ++ run_case("decode compressed", page, window, index_topk), ++ run_case("prefill-chunk swa-only", page, prefill_width, 0), ++ run_case("prefill-chunk compressed", page, prefill_width, index_topk), ++ ] ++ print("SELFTEST", "PASS" if all(ok) else "FAIL") ++ return 0 if all(ok) else 1 ++ ++ ++if __name__ == "__main__": ++ sys.exit(main()) diff --git a/dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch b/dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch new file mode 100644 index 00000000000..327c8fff0bc --- /dev/null +++ b/dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch @@ -0,0 +1,40 @@ +diff --git a/vllm/models/deepseek_v4_1/attention.py b/vllm/models/deepseek_v4_1/attention.py +index 9bb5352..39fed7a 100644 +--- a/vllm/models/deepseek_v4_1/attention.py ++++ b/vllm/models/deepseek_v4_1/attention.py +@@ -168,6 +168,10 @@ class DeepseekV4Attention(nn.Module, AttentionLayerBase, ABC): + # bf16 / per-tensor fp8 KV row. Backends can override the instance hook when + # a single attention class dispatches across arch-specific layouts. + use_fp8_ds_mla_layout: ClassVar[bool] = True ++ # Tokens per page of the SWA cache. The SWA backend admits any multiple of ++ # 32; a platform subclass whose kernel takes one fixed page size names it ++ # here (gunbc overlay: FlashInfer's SM120 DSv4 decode kernel takes only 64). ++ swa_block_size: ClassVar[int] = 32 + # Prefill is processed in fixed-size chunks; this bounds the bf16 kv-gather + # workspace allocated in _forward_prefill and is also read by the dummy-run + # path to pre-reserve that workspace. +@@ -453,7 +457,7 @@ class DeepseekV4Attention(nn.Module, AttentionLayerBase, ABC): + prefix=f"{prefix}.swa_cache", + cache_config=cache_config, + backend_cls=self.swa_backend_cls, +- block_size=32, ++ block_size=self.swa_block_size, + ) + + # The attention layer itself was already registered with the +diff --git a/vllm/models/deepseek_v4_1/nvidia/flashinfer_sparse.py b/vllm/models/deepseek_v4_1/nvidia/flashinfer_sparse.py +index 802dded..d69ea0e 100644 +--- a/vllm/models/deepseek_v4_1/nvidia/flashinfer_sparse.py ++++ b/vllm/models/deepseek_v4_1/nvidia/flashinfer_sparse.py +@@ -545,6 +545,11 @@ class DeepseekV4FlashInferSM120Attention(DeepseekV4Attention): + backend_cls = DeepseekV4FlashInferMLASparseBackend + swa_backend_cls = DeepseekSparseSWAFlashInferBackend + use_fp8_ds_mla_layout: ClassVar[bool] = True ++ # FlashInfer's SM120 DSv4 decode kernel is instantiated for 64-token pages ++ # only (flashinfer.mla._sparse_mla_sm120 _DECODE_DSV4_PAGE_BLOCK_SIZE), and ++ # every call with <= 64 query tokens must take it: the paged orchestrator is ++ # prefill-only. The SWA cache is the primary pool of every such call. ++ swa_block_size: ClassVar[int] = 64 + + @staticmethod + def _get_workspace(device: torch.device) -> torch.Tensor: diff --git a/dag/gunbc/spark/serving_arm.dag b/dag/gunbc/spark/serving_arm.dag index 31d90453132..3ffbd18f68d 100644 --- a/dag/gunbc/spark/serving_arm.dag +++ b/dag/gunbc/spark/serving_arm.dag @@ -10,6 +10,7 @@ import extdeps.vllm.engine_args { VllmPrefixCachingRequest, PrefixCachingOmitted, PrefixCachingExplicit, VllmKvCacheDtypeRequest, KvAuto, KvFp8, KvFp16, KvBf16, KvFp8DsMla, vllm_kv_cache_dtype_wire, VllmEngramConfigRequest, EngramConfigOmitted, EngramConfigExplicit, + VllmModalityRequest, ModalityOmitted, LanguageModelOnly, VllmHfOverride, } import extdeps.zhipu.glm_5_3_flash { glm_5_3_flash_native_manifest, glm_5_3_flash_vllm_integration_commit, glm_5_3_flash_mla_pe_dim, @@ -941,6 +942,7 @@ type ServingLaunchProfile sole_constructor { served_alias: NonEmptyStr prefix_caching: VllmPrefixCachingRequest engram: VllmEngramConfigRequest + modality: VllmModalityRequest } // THE PREFIX-CACHING REQUEST IS PART OF THE IDENTITY. Two launches differing only in whether they asked @@ -959,6 +961,7 @@ type ServingLaunchProfile sole_constructor { // draft-token count; "true" names neither. When this arm wants speculation it arrives as the modelled // policy that already exists -- extdeps.vllm.engine_args VllmSpeculativePolicy, and the argv shape // gunbc.spark.glm_canary_converge renders from SpeculationPolicy -- carrying the basis a Bool drops. +// THE MODALITY REQUEST ENTERS THE KEY ONLY WHEN IT IS MADE, for the Engram request's reason below. // THE ENGRAM REQUEST ENTERS THE KEY ONLY WHEN IT IS MADE. An explicit request changes what the engine // allocates, so two profiles differing in it are two subjects; an omitted one spells no flag, and // rendering it into the key would re-key every existing subject -- and every receipt recorded against @@ -981,6 +984,11 @@ fn serving_launch_profile_key(p: ServingLaunchProfile) -> String { EngramConfigExplicit { cpu_offload: c, dp_shared_memory: d } => join([" engram=cpu_offload:", if c { "on" } else { "off" }, ",dp_shared_memory:", if d { "on" } else { "off" }], "") }, + match p.modality { + ModalityOmitted => "" + LanguageModelOnly { hf_overrides: os } => + join([" modality=language-only", join(map(os, o => join([",", o.key as String, ":", o.json_value as String], "")), "")], "") + }, ], "") } @@ -1095,6 +1103,7 @@ data glm_native_initial_profile: ServingLaunchProfile = ServingLaunchProfile { served_alias: "glm-5.3-flash" as NonEmptyStr, prefix_caching: PrefixCachingExplicit { enabled: true }, engram: EngramConfigOmitted, + modality: ModalityOmitted, } // ── THE CONCURRENT PROFILE IS A DIFFERENT SUBJECT, AND THAT IS THE POINT ────────────────────── @@ -1151,6 +1160,7 @@ data glm_native_concurrent_profile: ServingLaunchProfile = ServingLaunchProfile served_alias: glm_native_initial_profile.served_alias, prefix_caching: PrefixCachingExplicit { enabled: true }, engram: EngramConfigOmitted, + modality: ModalityOmitted, } // ── THE V4.1 MEASUREMENT PROFILE: WHAT V4.1 DECIDES, OVER WHAT THE MEASUREMENT DECIDES ───────── @@ -1163,7 +1173,18 @@ data glm_native_concurrent_profile: ServingLaunchProfile = ServingLaunchProfile // - the KV cache is requested as fp8_ds_mla, the layout the measurement's startup receipt requires // to resolve (ResolvedFp8DsMla) before it admits a rank; // - the Engram tables are served through the prefetch path -- cpu_offload true, dp_shared_memory -// false -- which is what the file-backed overlay requires and refuses at model init without. +// false -- which is what the file-backed overlay requires and refuses at model init without; +// - THE LAUNCH SERVES TEXT ONLY, AND THE VISION CAPABILITY IS DROPPED ON PURPOSE. The checkpoint is +// the vision variant (vision_n_layers > 0), and at the pinned revision that widens every prefill +// sliding-window row from sliding_window (128) to sliding_window + vision_max_n_token (1152) for +// bidirectional image spans (vllm/v1/attention/backends/mla/sparse_swa.py prefill_index_width, +// vllm/models/deepseek_v4_1/attention.py max_image_tokens). FlashInfer's SM120 DSv4 kernel has no +// specialization at 1152, and every call with 64 or fewer query tokens must take one, so attempt +// 13 died in CUDA-graph capture on all four ranks. --language-model-only stubs the tower and +// admits no image; overriding vision_n_layers to 0 is what returns the width to 128, because +// max_image_tokens is read off the config and not off the modality limit. A launch that wants +// images back needs an SM120 kernel at the widened width first; it is a different product, not +// this one with a flag flipped. // The route's port and served name are the caller's because they belong to the route this launch // answers, not to the model. fn deepseek_v41_measurement_profile( @@ -1185,6 +1206,7 @@ fn deepseek_v41_measurement_profile( served_alias: served_alias, prefix_caching: prefix_caching, engram: EngramConfigExplicit { cpu_offload: true, dp_shared_memory: false }, + modality: LanguageModelOnly { hf_overrides: [VllmHfOverride { key: "vision_n_layers" as NonEmptyStr, json_value: "0" as NonEmptyStr }] }, } } @@ -1217,6 +1239,7 @@ data glm_exl3_canary_profile: ServingLaunchProfile = ServingLaunchProfile { served_alias: "glm-5.3-flash-exl3-canary" as NonEmptyStr, prefix_caching: PrefixCachingOmitted, engram: EngramConfigOmitted, + modality: ModalityOmitted, } diff --git a/dag/gunbc/spark/serving_arm_launch.dag b/dag/gunbc/spark/serving_arm_launch.dag index 02f76da9308..2c0a6c9af58 100644 --- a/dag/gunbc/spark/serving_arm_launch.dag +++ b/dag/gunbc/spark/serving_arm_launch.dag @@ -73,7 +73,7 @@ import std.measure { token_count_value, } import std.decl_ref { decl_ref } -import extdeps.vllm.engine_args { vllm_prefix_caching_argv, vllm_basis_point_fraction_wire, vllm_kv_cache_dtype_wire, vllm_engram_config_argv } +import extdeps.vllm.engine_args { vllm_prefix_caching_argv, vllm_basis_point_fraction_wire, vllm_kv_cache_dtype_wire, vllm_engram_config_argv, vllm_modality_argv } import gunbc.spark.model_snapshot_acquisition { SnapshotPlacement, @@ -483,7 +483,7 @@ fn arm_engine_argv(node: ArmNode, master: Ipv4Address, crow: ServingCheckpointRo ], concat( vllm_prefix_caching_argv(r: profile.prefix_caching), concat(checkpoint_quantization_argv(row: crow), concat(arm_headless_argv(node: node), - concat(arm_tokenizer_mode_argv(crow: crow), vllm_engram_config_argv(r: profile.engram)))))) + concat(arm_tokenizer_mode_argv(crow: crow), concat(vllm_engram_config_argv(r: profile.engram), vllm_modality_argv(r: profile.modality))))))) } // ── ONLY RANK ZERO HOSTS AN API SERVER, AND THE OTHER THREE MUST BE TOLD SO ──────────────────── diff --git a/dag/gunbc/spark/v41_runtime_candidate.dag b/dag/gunbc/spark/v41_runtime_candidate.dag index a65f4a9f710..74e5788594e 100644 --- a/dag/gunbc/spark/v41_runtime_candidate.dag +++ b/dag/gunbc/spark/v41_runtime_candidate.dag @@ -884,7 +884,44 @@ fn v41_storage_backed_patch_population() -> V41PatchPopulation { Present { value: hook_digest } => match v41_engine_patch(path: v41_engram_storage_hook_path, applies_to: vllm_d2d649e6_revision, sha256_hex: hook_digest.hex as String, purpose: "binding of the file-backed overlay in engram.py: FileBackedParallelEngramEmbedding for every Engram layer when the launch names a row-store manifest" as NonEmptyStr) { Absent => V41PatchesUnestablished { obligation: "DigestFile of the engram.py import member was not a sha256 digest" as NonEmptyStr } - Present { value: hook } => V41PatchesObserved { patches: [overlay, hook], partition: PartitionByLaunchConfiguration } + Present { value: hook } => v41_with_sm120_members(engram: [overlay, hook]) + } + } + } + } +} + +// THE SM120 MEMBERS. The population is the image's whole patch list, applied in order, and these two +// follow the Engram pair; neither touches a file the Engram members touch. +// 1. The SWA cache's page size. vllm/models/deepseek_v4_1/attention.py constructs DeepseekV4SWACache +// with the literal block_size=32 (V4 passes the SWA cache's own default, 64). FlashInfer's SM120 DSv4 +// decode kernel is instantiated for 64-token pages only (flashinfer.mla._sparse_mla_sm120 +// _DECODE_DSV4_PAGE_BLOCK_SIZE) and every call with 64 or fewer query tokens must take it, with the +// SWA cache as its primary pool -- so on this tree every decode step is refused, whatever its topk. +// No engine argument reaches that literal. The member lifts it to a class attribute, +// swa_block_size, default 32, and DeepseekV4FlashInferSM120Attention sets 64: the other platform +// classes keep the page they had. +// 2. The control. It creates vllm.models.deepseek_v4_1.nvidia.sm120_decode_shape_selftest, which calls +// FlashInfer's DSv4 entry point at the TP4 launch's decode and short-prefill shapes, reading the page +// size and the prefill width off the installed tree and the launch's own config arguments. Applied +// alone to the unpatched tree it exits 1 on attempt 13's exact refusal; with member 1 and the +// launch's text-only request (gunbc.spark.serving_arm deepseek_v41_measurement_profile) it exits 0. +data v41_sm120_swa_page_patch_path: NonEmptyStr = "dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch" as NonEmptyStr +data v41_sm120_decode_selftest_patch_path: NonEmptyStr = "dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch" as NonEmptyStr + +fn v41_with_sm120_members(engram: List) -> V41PatchPopulation { + match v41_digest_declared_patch(path: v41_sm120_swa_page_patch_path) { + Absent => V41PatchesUnestablished { obligation: "DigestFile could not read the SM120 SWA page member at dag/gunbc/spark/patches/d2d649e6/sm120_swa_page_block_size.patch" as NonEmptyStr } + Present { value: page_digest } => + match v41_engine_patch(path: v41_sm120_swa_page_patch_path, applies_to: vllm_d2d649e6_revision, sha256_hex: page_digest.hex as String, purpose: "SM120 SWA cache page of 64 tokens: DeepseekV4FlashInferSM120Attention.swa_block_size, the only page FlashInfer's SM120 DSv4 decode kernel takes" as NonEmptyStr) { + Absent => V41PatchesUnestablished { obligation: "DigestFile of the SM120 SWA page member was not a sha256 digest" as NonEmptyStr } + Present { value: page } => + match v41_digest_declared_patch(path: v41_sm120_decode_selftest_patch_path) { + Absent => V41PatchesUnestablished { obligation: "DigestFile could not read the SM120 decode-shape self-test member at dag/gunbc/spark/patches/d2d649e6/sm120_decode_shape_selftest.patch" as NonEmptyStr } + Present { value: test_digest } => + match v41_engine_patch(path: v41_sm120_decode_selftest_patch_path, applies_to: vllm_d2d649e6_revision, sha256_hex: test_digest.hex as String, purpose: "SM120 decode-shape self-test: FlashInfer's DSv4 entry point at the TP4 launch's decode and short-prefill shapes" as NonEmptyStr) { + Absent => V41PatchesUnestablished { obligation: "DigestFile of the SM120 self-test member was not a sha256 digest" as NonEmptyStr } + Present { value: control } => V41PatchesObserved { patches: concat(engram, [page, control]), partition: PartitionByLaunchConfiguration } } } } diff --git a/dag/test/claim/spark/v41_engram_runtime_witness_test.dag b/dag/test/claim/spark/v41_engram_runtime_witness_test.dag index b8fa9dcb86b..54757577808 100644 --- a/dag/test/claim/spark/v41_engram_runtime_witness_test.dag +++ b/dag/test/claim/spark/v41_engram_runtime_witness_test.dag @@ -441,7 +441,7 @@ test fn staging_verifies_readback_and_refuses_a_wrong_rank_identity() -> Bool { test fn the_live_storage_backed_patches_are_digestfile_of_committed_bytes() -> Bool { match v41_storage_backed_patch_population() { V41PatchesObserved { patches: ps, partition: _ } => - length(ps) == 2 + length(ps) == 4 && all(ps, p => (p.path as String).contains("d2d649e6")) V41PatchesAbsentByObservation => false V41PatchesUnestablished { obligation: _ } => false diff --git a/dag/test/claim/spark/v41_source_patch_converge_witness_test.dag b/dag/test/claim/spark/v41_source_patch_converge_witness_test.dag index f300840a45a..ad3eaf320b5 100644 --- a/dag/test/claim/spark/v41_source_patch_converge_witness_test.dag +++ b/dag/test/claim/spark/v41_source_patch_converge_witness_test.dag @@ -284,7 +284,7 @@ test fn noop_readback_grounds_and_apply_without_after_read_does_not() -> Bool { test fn digestfile_inhabits_the_committed_patch_population() -> Bool { match v41_storage_backed_patch_population() { V41PatchesObserved { patches: ps, partition: _ } => - length(ps) == 2 + length(ps) == 4 && match get(xs: ps, index: 1) { Absent => false Present { value: p } => (p.purpose as String).contains("FileBackedParallelEngramEmbedding for every Engram layer") From a64d70cacb86962e4847b63a0d868a63678224b3 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Thu, 1 Oct 2026 12:11:42 +0000 Subject: [PATCH 20/21] DEV: VllmHfOverride constructible from serving_arm (bug in #12906); move the claim-skip note out of the body (never merge) Co-Authored-By: Claude Opus 5.5 (1M context) --- dag/extdeps/vllm/engine_args.dag | 2 +- dag/gunbc/spark/v41_runtime_image_converge.dag | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/dag/extdeps/vllm/engine_args.dag b/dag/extdeps/vllm/engine_args.dag index f2cc7a2d379..c737c003d52 100644 --- a/dag/extdeps/vllm/engine_args.dag +++ b/dag/extdeps/vllm/engine_args.dag @@ -204,7 +204,7 @@ fn vllm_engram_config_argv(r: VllmEngramConfigRequest) -> List { // above is dotted. // WHICH CONFIG KEY A LAUNCH OVERRIDES, AND TO WHAT, IS THE LAUNCH'S POLICY, NOT A FACT ABOUT vLLM: the // key and its value arrive as data from the layer that decides them. Omitted renders no flag at all. -type VllmHfOverride sole_constructor { +type VllmHfOverride { key: NonEmptyStr json_value: NonEmptyStr } diff --git a/dag/gunbc/spark/v41_runtime_image_converge.dag b/dag/gunbc/spark/v41_runtime_image_converge.dag index 3bd5702ed15..6fe928928e6 100644 --- a/dag/gunbc/spark/v41_runtime_image_converge.dag +++ b/dag/gunbc/spark/v41_runtime_image_converge.dag @@ -446,7 +446,6 @@ fn v41_runtime_image_produce_ci_wet() -> ProcessExit match admit_host_held_by_subject(raw: build_host, subject: PairServingUnitOn { group: FabricGroupA }) { Present { value: why } => exit_failure(reason: why) Absent => - // DEV-ONLY (operator ruling 2026-09-30, never merged): Group A fences srv5-8 in the fabric store, so the host-effect claim is skipped for the dev rebuild. release_host_effect_live(handle: Absent, exit: { match with_spark_build_host( host: build_host, From 084201d13549cc049a94bedffa22e723028e86d2 Mon Sep 17 00:00:00 2001 From: Brian Searls Date: Thu, 1 Oct 2026 12:48:09 +0000 Subject: [PATCH 21/21] DEV: V4.1 production image = rebuild 36862980380 (Engram host tensors on cpu) --- dag/gunbc/spark/v41_runtime_image_converge.dag | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dag/gunbc/spark/v41_runtime_image_converge.dag b/dag/gunbc/spark/v41_runtime_image_converge.dag index 6fe928928e6..8d0efbe7194 100644 --- a/dag/gunbc/spark/v41_runtime_image_converge.dag +++ b/dag/gunbc/spark/v41_runtime_image_converge.dag @@ -539,11 +539,11 @@ data v41_arm_a_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { data v41_production_produced_image: V41ProducedImageHeld = V41ProducedImageHeld { observed: v41_observed_source_identity_of( revision_head: "d2d649e674c75425d2d6975c87eb89fd4d55fff8" as NonEmptyStr, - worktree_diff_sha256: Sha256Digest { hex: "bf03a450c63d6e77f99e1a21f28468b8904aa60f8703310f96829022b8d3e4b7" as Sha256DigestHex }, + worktree_diff_sha256: Sha256Digest { hex: "3c3474c5107aec1426ec96a4f6204d823e4dbe759ecb233bb1913ebe18a6f129" as Sha256DigestHex }, ), - config: OciSha256Digest(Sha256Digest { hex: "0f6039eba7d49f5786dcfe9a3b38510b57b03390ecffeb1dcb8d02eb49b57a4a" as Sha256DigestHex }), + config: OciSha256Digest(Sha256Digest { hex: "60dc7f83e05fae5c8c946f3c8856b951bd50121364119803fb2b68b90df7bbab" as Sha256DigestHex }), held_on: "srv8" as HostIdentity, - receipt_run: "36805010658" as NonEmptyStr, + receipt_run: "36862980380" as NonEmptyStr, } fn v41_produced_image_subject(p: V41ProducedImageHeld) -> VllmImageDistributionSubject {