# GB10 Long Context: Read the Failure Before Changing Flags A long prompt can expose a kernel path that a short chat never exercises. On a pair of GB10 Sparks, first identify the failed layer: dtype handling, architecture support, per-block shared memory, correctness or container transport. Preserve the exact build and request shape before changing a setting. This note interprets historical records from [WeZZard’s dgx-spark-bench](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/README.md), pinned at `25e8dbf1b6113dd0b3f12d27a670fb284944d260`. The relevant incidents are from September 2026; the documents were checked on 18 October. They do not establish the behavior of today’s upstream releases. Run AI Here did not reproduce these model runs. ![Conceptual incident review from observation through isolation to verification, with distinct evidence states.](/assets/a08-batch3-local-spark.webp) Figure: a conceptual incident-review diagram generated as SVG by the local Spark model and reviewed for publication. It contains no measured timing or hardware topology. ## Four evidence states, four different conclusions | Case | What the pinned source supports | What you should preserve | |---|---|---| | Qwen FP8 chunked prefill | A single-user failure was reproduced by the source | Probe version, KV dtype, chunk size and prompt shape | | QSA long-context corruption warning | A warning and launcher guard; the proposed ladder remains unperformed | Unknown threshold and missing correctness results | | GLM backend and tilelang failures | Different attempted paths failed; a tuned bf16 path served | The successful image, kernel tuning and later follow-up | | Hidden container RDMA devices | A source-verified transport incident | Running-container transport evidence, separate from host bandwidth | ## Case 1: a single user can trigger the FP8 KV crash The Qwen probe used `RadixArk/Qwen3.8-Flash-Next-NVFP4`, revision `7b719225242aacd3dbd3f9407468c2ee9a9d2594`, with `sglang-qwen38fn:sm121-ours`. Its log identifies SGLang `0.0.0.dev1+gd91c3682b`, `float8_e4m3fn` KV and a 310,656-token allocated pool. A short single-user control completed 64 output tokens and retrieved its needle. A 30,301-token single-user prompt then hit `Unsupported rhs dtype fp8e4nv` and terminated the scheduler. [Probe and diagnosis](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L142). The source traces this to nonzero prefix lengths in chunked prefill. With an 8,192-token chunk size, a longer request reaches a path reading the FP8 paged pool instead of fresh bf16 K/V. Concurrency alone is therefore the wrong explanation. The document proposes a compatible build or dtype conversion; it supplies no verified fix commit. Raising the chunk size moves the boundary. It does not demonstrate a corrected kernel. The failed probe’s output count and the image digest remain unknown. Do not borrow the 332,480-token pool from an earlier campaign: it is a different observation in the same document. ## Case 2: the 120,000-token guard is not a correctness result The QSA document carries a warning about corruption beyond roughly 120K and contrasts it with a recipe requesting 262,144 tokens. The author inspected the image, did not find the named successor attention path, and capped the launcher at 120,000. The proposed 131K/262K retrieval ladder is explicitly still pending in this snapshot. [Warning, guard and unperformed experiment](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L3). Preserve the distinction between an operating guard and a measured maximum safe context. No completed ladder, corruption threshold or upstream fix verification is available here. The named `pr36845.diff` is a successor reference, not proof of a merged or validated fix. A future retrieval pass would be useful evidence for that task; it would not establish universal long-context correctness. ## Case 3: unified RAM and kernel shared memory are different limits The GLM records concern `LibertAIDAI/GLM-5.3-Flash-NVFP4`, revision `aa28e1f54130286c95fee10d0705c74ce8743734`. FP8 attempts encountered an architecture allowlist rejection and a separate TensorRT-LLM unsupported-architecture error. A stock bf16 tilelang path requested 169,984 bytes of dynamic shared memory per block against a reported 101,376-byte device limit. This is not the machine’s 121 GiB unified-memory pool. [Attempt sequence](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/glm-dsa.md#L130). The source reports success with bf16 KV in `sglang-glm53:gb10-tilelang`, using `block_I=32`, `num_stages=1`, `threads=128`. The early run allocated 800,448 KV tokens and served 26,535-token prompts with eight users. The [pinned community build patch](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docker/patch-tilelang-gb10.py#L40) identifies the modified v1 path; it is not an upstream release-fix claim. Read the chronology. Although the early GLM document says DFlash remains untested, the [later speculation document](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/speculation.md#L248) reports clean eight-token draft-verify capture for batch sizes one through eight on the tuned image. Keep that follow-up separate from the early no-speculation pool. Do not present the earlier warning as the final outcome. ## Case 4: verify transport inside the running container The source’s September 4 Qwen investigation found working host fabric and visible container HCA entries in sysfs, while `/dev/infiniband` was absent. NCCL used sockets. Mapping the required verbs devices, setting IPC_LOCK/memlock and selecting the connected HCA led to channel logs showing `NET/IB`. [Container evidence and remediation](https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/interconnect.md#L8). A host link test cannot confirm the serving container’s transport. Inspect the running service after startup; device visibility alone is also insufficient. This note does not promote the earlier sweep’s speed ratios to a current benchmark, and the original author still identifies remaining unexplained latency. ## A useful incident record Before tuning, save the first relevant error, artifact revision, engine version, image digest when known, attention backend, KV dtype and allocated pool. Add prompt/output lengths, concurrency, chunk size, cache state and speculative mode. An omitted runtime identifier stays unknown even when the documentation is pinned. Then write one falsifiable explanation and one change that tests it. Keep a short control and a representative longer correctness task. A server that starts proves startup; it does not prove long-context correctness. Preserve failed runs, successful follow-ups and unperformed proposals as separate records. The adjacent JSON exposes four `case_reports` with explicit `evidence_status`, per-case `conditions`, source pointers, remediation state and unresolved fields. Unknown values remain `null`; the warning is never exported as a completed experiment. Use the [two-Spark conditions guide](/guides/two-spark-deployment-conditions/) and [For AI reference](/for-ai/) alongside these cases. ## Original sources and conditions - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/README.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: b815d32807147f9638b1abbeadc02af257c55bd1f7651a123e5a8090bebe7469 - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/REPORT.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: d08b9489656537f7b4a5d11cf4f07ed9f58c9b86ba6c7badd20ca52260a8224e - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8 - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/glm-dsa.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: a4b9448f5265f687504130bc1791ec3e7c1777d16c648851a94c3a3b4db99e32 - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/interconnect.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: 9e7904db6d48fc44691c9f170c03f402c77480dcbf54f66bc107ef57ef0ea94f - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/speculation.md - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: 46c64755b468e6e996e84fd1d19727160a9130bb3565731e7b5f26185c1dfe7f - WeZZard: https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docker/patch-tilelang-gb10.py - Document revision: 25e8dbf1b6113dd0b3f12d27a670fb284944d260 - SHA-256: 88455f3811883262abc5ea292ca5db0dda91fd8fd6f6db7d12ce6d6a01b817ff Checked 2026-10-18. AI-assisted writing, source checked by Run AI Here. No independent model tests. Conditions and selected source facts are in the adjacent JSON export. ```json { "facts": [], "reported_results": [], "conditions": { "community_snapshot_revision": "25e8dbf1b6113dd0b3f12d27a670fb284944d260", "snapshot_date": "2026-09-06", "source_checked_on": "2026-10-18", "historical_cases_not_current_release_claims": true, "independent_site_model_runs": 0 }, "case_reports": [ { "id": "qwen-fp8-chunked-prefill", "title": "FP8 KV kills a single-request chunked prefill", "evidence_status": "source_reproduced_failure", "conditions": { "hardware": "2 x DGX Spark GB10 sm_121", "node_count": 2, "memory_per_node_gib": 121, "physical_interconnect_gbps": 200, "parallelism": "TP2", "engine": "SGLang", "image_digest": null, "runtime_execution_revision": null, "temperature": null, "top_p": null, "site_reproduced": false, "model_artifact": "RadixArk/Qwen3.8-Flash-Next-NVFP4", "checkpoint_revision": "7b719225242aacd3dbd3f9407468c2ee9a9d2594", "weight_precision": "NVFP4", "image_tag": "sglang-qwen38fn:sm121-ours", "experiment_date": "2026-09-04", "executed_engine_version": "0.0.0.dev1+gd91c3682b", "kv_dtype": "float8_e4m3fn", "kv_pool_tokens": 310656, "configured_max_total_tokens": 600000, "context_limit_tokens": 120000, "chunked_prefill_size": 8192, "concurrency": 1, "reported_long_prompt_tokens": 30301, "failed_probe_output_tokens": null, "short_control_workload": "agentic-4k", "short_control_prompt_tokens": null, "short_control_output_tokens": 64, "short_control_failed_requests": 0, "short_control_needle_hits": 1, "short_control_needle_trials": 1, "speculative_decoding": "same source launcher; exact probe flag tuple not fully logged" }, "symptom": "Unsupported rhs dtype fp8e4nv; scheduler exits during the longer probe.", "diagnosis": "The source traces nonzero prefix lengths to the _ck path, which consumes the FP8 paged KV pool; the no-prefix path consumes fresh bfloat16 tensors. Chunk splitting, rather than concurrent users alone, explains the single-user reproduction.", "remediation": { "state": "proposed_not_verified_in_this_probe", "proposal": "A compatible image/Triton path or a dtype conversion at the _ck call site; the later campaign used bfloat16 KV. Increasing chunk size only moves the failing boundary.", "verified_fix_commit": null }, "unresolved": [ "Failed-request generated-token count and image digest are not supplied.", "The probe does not establish the behavior of a current SGLang release." ], "observations": [ { "source_document_id": "spark-docs-qwen-qsa", "source_document_sha256": "b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8", "source_line": 152, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L152" }, { "source_document_id": "spark-docs-qwen-qsa", "source_document_sha256": "b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8", "source_line": 146, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L146" }, { "source_document_id": "spark-docs-qwen-qsa", "source_document_sha256": "b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8", "source_line": 136, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L136" } ] }, { "id": "qsa-long-context-warning", "title": "The approximately 120K corruption warning is not a completed correctness test", "evidence_status": "source_warning_unperformed", "conditions": { "hardware": "2 x DGX Spark GB10 sm_121", "node_count": 2, "memory_per_node_gib": 121, "physical_interconnect_gbps": 200, "parallelism": "TP2", "engine": "SGLang", "image_digest": null, "runtime_execution_revision": null, "temperature": null, "top_p": null, "site_reproduced": false, "model_artifact": "RadixArk/Qwen3.8-Flash-Next-NVFP4", "checkpoint_revision": "7b719225242aacd3dbd3f9407468c2ee9a9d2594", "weight_precision": "NVFP4", "image_tag": "sglang-qwen38fn:sm121-ours", "experiment_date": null, "executed_engine_version": null, "kv_dtype": null, "kv_pool_tokens": null, "concurrency": null, "reported_prompt_tokens": null, "output_tokens": null, "source_recipe_context_limit_tokens": 262144, "launcher_context_limit_tokens": 120000, "proposed_test_contexts_tokens": [ 131072, 262144 ] }, "symptom": "A carried warning about QSA output corruption conflicts with a longer published context setting.", "diagnosis": "The author inspected the image for the successor attention path and chose a conservative launcher cap; the planned long-context retrieval ladder is explicitly still unperformed in this document.", "remediation": { "state": "guard_applied_correctness_ladder_unperformed", "proposal": "Pin and validate a matching successor build before claiming longer correct context. The document names pr36845.diff; its merged/fixed status is not established here.", "verified_fix_commit": null, "named_successor_artifact": "pr36845.diff" }, "unresolved": [ "120000 is a launcher guard, not a measured maximum safe context.", "No completed ladder result, corruption threshold or verified upstream fix is available in this snapshot.", "A retrieval check is useful but does not prove correctness for all tasks." ], "observations": [ { "source_document_id": "spark-docs-qwen-qsa", "source_document_sha256": "b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8", "source_line": 3, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L3" }, { "source_document_id": "spark-docs-qwen-qsa", "source_document_sha256": "b642874472f6727e05dc259ab38ff6bc7d78a80691436c65a9c34417abc201e8", "source_line": 76, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/qwen-qsa.md#L76" } ] }, { "id": "glm-gb10-backend-shared-memory", "title": "Separate architecture rejection from per-block shared-memory failure", "evidence_status": "source_reproduced_remediation", "conditions": { "hardware": "2 x DGX Spark GB10 sm_121", "node_count": 2, "memory_per_node_gib": 121, "physical_interconnect_gbps": 200, "parallelism": "TP2", "engine": "SGLang", "image_digest": null, "runtime_execution_revision": null, "temperature": null, "top_p": null, "site_reproduced": false, "model_artifact": "LibertAIDAI/GLM-5.3-Flash-NVFP4", "checkpoint_revision": "aa28e1f54130286c95fee10d0705c74ce8743734", "weight_precision": "NVFP4", "image_tag": "sglang-glm53:gb10-tilelang", "executed_engine_version": null, "experiment_date": "2026-09-04", "kv_dtype": "bfloat16 on working tilelang path; early rejected attempts used fp8_e4m3", "kv_pool_tokens": 800448, "reported_prompt_tokens": 26535, "concurrency": 8, "output_tokens": null, "speculative_decoding": "off in the early success; later tuned DFlash verify is a separate follow-up", "reported_transport": "NCCL over RoCE", "stock_dynamic_shared_memory_bytes_per_block": 169984, "device_shared_memory_limit_bytes_per_block": 101376, "working_tuning": { "block_I": 32, "num_stages": 1, "threads": 128 } }, "symptom": "An architecture allowlist rejection, an sm121 unsupported-architecture failure, and a tilelang dynamic shared-memory failure occur in different attempted paths.", "diagnosis": "The first two are backend/build compatibility failures. The third exceeds the GPU per-block shared-memory limit; this is separate from unified system RAM and is not proof that long context alone caused an out-of-memory condition.", "remediation": { "state": "source_verified_for_named_tuned_image", "proposal": "The source selects bfloat16 KV and tuned tilelang. A pinned community build patch applies the smaller v1 kernel configuration when the device limit is below the stock request.", "verified_upstream_fix_commit": null, "community_patch": { "source_document_id": "spark-docker-patch-tilelang-gb10", "source_document_sha256": "88455f3811883262abc5ea292ca5db0dda91fd8fd6f6db7d12ce6d6a01b817ff", "source_line": 45, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docker/patch-tilelang-gb10.py#L45" } }, "follow_up": { "state": "source_verified_separate_follow_up", "observation": "The later speculation document reports clean DFlash draft-verify graph capture with eight tokens per request for batch sizes one through eight on the already tuned image; do not carry forward the early untested warning as the final state.", "conditions": { "image_tag": "sglang-glm53:gb10-tilelang", "kv_dtype": "bfloat16", "draft_tokens": 8, "graph_capture_batch_sizes_min": 1, "graph_capture_batch_sizes_max": 8, "earlier_kv_pool_tokens_not_reused_as_final_profile": true }, "source_document_id": "spark-docs-speculation", "source_document_sha256": "46c64755b468e6e996e84fd1d19727160a9130bb3565731e7b5f26185c1dfe7f", "source_line": 248, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/speculation.md#L248" }, "unresolved": [ "No immutable image digest or executed engine version is supplied for the early success.", "This community patch is not evidence of an upstream release fix." ], "observations": [ { "source_document_id": "spark-docs-glm-dsa", "source_document_sha256": "a4b9448f5265f687504130bc1791ec3e7c1777d16c648851a94c3a3b4db99e32", "source_line": 95, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/glm-dsa.md#L95" }, { "source_document_id": "spark-docs-glm-dsa", "source_document_sha256": "a4b9448f5265f687504130bc1791ec3e7c1777d16c648851a94c3a3b4db99e32", "source_line": 92, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/glm-dsa.md#L92" } ] }, { "id": "container-rdma-hidden", "title": "A fast host link does not prove container RoCE", "evidence_status": "source_reproduced_transport_incident", "conditions": { "hardware": "2 x DGX Spark GB10 sm_121", "node_count": 2, "memory_per_node_gib": 121, "physical_interconnect_gbps": 200, "parallelism": "TP2", "engine": "SGLang", "image_digest": null, "runtime_execution_revision": null, "temperature": null, "top_p": null, "site_reproduced": false, "model_artifact": "RadixArk/Qwen3.8-Flash-Next-NVFP4", "checkpoint_revision": "7b719225242aacd3dbd3f9407468c2ee9a9d2594", "weight_precision": "NVFP4", "image_tag": "sglang-qwen38fn:sm121-ours", "experiment_date": "2026-09-04", "executed_engine_version": null, "kv_dtype": null, "kv_pool_tokens": null, "concurrency": [ 1, 6 ], "reported_prompt_range_tokens": "74 to approximately 30300 across the source sweep", "output_tokens": null }, "symptom": "The host fabric works and container sysfs lists HCAs, but verbs devices are absent and NCCL uses sockets.", "diagnosis": "Visible /sys/class/infiniband entries do not establish access to /dev/infiniband. Host bandwidth is not an observation of the running container transport.", "remediation": { "state": "source_verified_transport_change", "proposal": "The source maps the required verbs devices, sets IPC_LOCK/memlock and selects the connected HCA; actual NCCL channel logs then show NET/IB. Recheck transport after the service starts.", "verified_fix_commit": null }, "unresolved": [ "No speed ratio from the earlier sweep is promoted to a current benchmark here.", "This incident does not explain every remaining source latency." ], "observations": [ { "source_document_id": "spark-docs-interconnect", "source_document_sha256": "9e7904db6d48fc44691c9f170c03f402c77480dcbf54f66bc107ef57ef0ea94f", "source_line": 17, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/interconnect.md#L17" }, { "source_document_id": "spark-docs-interconnect", "source_document_sha256": "9e7904db6d48fc44691c9f170c03f402c77480dcbf54f66bc107ef57ef0ea94f", "source_line": 69, "evidence_url": "https://github.com/WeZZard/dgx-spark-bench/blob/25e8dbf1b6113dd0b3f12d27a670fb284944d260/docs/interconnect.md#L69" } ] } ] } ```