Why Driver-Visible Paths Can Look Like Hardware Support on GB10, Even When Silicon Proof Is Missing
A field report on GB10 reverse engineering: how libcuda tables, helper cubins, and signed capability metadata can make tcgen05 look reachable from software while still falling short of proving that the underlying silicon really exposes the same path as B200 or GB100.

One of the easiest mistakes in GPU reverse engineering is to confuse a software-visible path with a hardware-proven capability.
That mistake is especially tempting on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story.
Once you start reading libcudaQuick term guidelibcudaThe user-space NVIDIA driver library that owns module load, metadata validation, and the helper-cubin patch lane in the GB10 experiments.GroundingHistory: libcuda patch lane Example: GB10 cubin patch repro Example: driver signal vs runtime proof sample, helper cubins, and capability metadata, the
datacenter Blackwell path can look uncannily close. You find architecture
tables. You find helper assets. You find capability descriptors. You can
sometimes patch one layer and watch the failure move to the next one. It is
very natural to tell yourself that you are one more patch away from exposing a
physically present feature.
Sometimes that is true. Sometimes it is the precise moment where a project starts lying to itself.
Our GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story smtest work is a good example of why this distinction matters.
For first touch, keep the target and metadata names narrow. sm_100Quick term guidesm_100Baseline Blackwell compiler target name in NVIDIA's architecture vocabulary, distinct from the architecture-specific and family-specific targets used elsewhere in the GB10 lane.GroundingAbout: GB10 tensor-path proof summary History: GB10 journey Example: GB10 cubin patch repro is the
baseline Blackwell target name in the compiler vocabulary, sm_100aQuick term guidesm_100aDatacenter Blackwell cubin target used by GB100/B200-class paths and by the source cubins in the public GB10 arch-patch repro.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Example: GB10 repro walkthrough is the
architecture-specific datacenter Blackwell target, and sm_121aQuick term guidesm_121aConsumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough is the
architecture-specific GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story target used by these receipts. tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough names the
SM100Quick term guidesm_100Baseline Blackwell compiler target name in NVIDIA's architecture vocabulary, distinct from the architecture-specific and family-specific targets used elsewhere in the GB10 lane.GroundingAbout: GB10 tensor-path proof summary History: GB10 journey Example: GB10 cubin patch repro tensor path we probed; .nv.capmercQuick term guide.nv.capmercObserved section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.GroundingHistory: libcuda patch lane Example: GB10 gate repro Example: GB10 gate walkthrough and .nv.merc.rela.*Quick term guide.nv.merc.rela.*Observed companion metadata section family that appears with .nv.capmerc in the deeper GB10 gate boundary.GroundingExample: GB10 gate walkthrough Reference: GB10 tensor-path proof summary are observed
capability-metadata stops in the cubin lane, not public API guarantees.
The misleading signal
The strongest misleading signal on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story is not one big thing. It is the accumulation of several smaller truths:
- a baseline
sm_100aQuick term guidesm_100aDatacenter Blackwell cubin target used by GB100/B200-class paths and by the source cubins in the public GB10 arch-patch repro.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Example: GB10 repro walkthrough cubin can run after an architecture-field rewrite; - the driver clearly contains multiple product-gating layers;
libcudaQuick term guidelibcudaThe user-space NVIDIA driver library that owns module load, metadata validation, and the helper-cubin patch lane in the GB10 experiments.GroundingHistory: libcuda patch lane Example: GB10 cubin patch repro Example: driver signal vs runtime proof sample ships capability-related machinery that is richer than a simple "unsupported GPU" branch;- helper selection and architecture routing can be nudged by patching metadata rather than by modifying hardware.
If you line those up in the most optimistic order, the story writes itself: the path is present, the driver is the only blocker, and deeper patching will eventually reveal the real feature.
The problem is that this is still a story built from indirect evidence.
The smallest checked-in decoder for that mistake is the driver signal versus
runtime proof sample.
For the claim boundary around the same receipts, pair this post with what our
GB10 experiments actually prove
and the sm_100a cubin patch repro.
What the baseline result really means
The baseline arithmetic probe is worth emphasizing because it is the cleanest fact in the entire chain.
After a low-bit e_flagsQuick term guidee_flagsThe ELF header field whose low architecture bits are rewritten from sm_100a to sm_121a in the public GB10 arch-patch lane.GroundingAbout: sm_100a cubin patch repro Example: baseline arch-patch proof sample Reference: ELF arch-field patcher rewrite from sm_100aQuick term guidesm_100aDatacenter Blackwell cubin target used by GB100/B200-class paths and by the source cubins in the public GB10 arch-patch repro.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Example: GB10 repro walkthrough to sm_121aQuick term guidesm_121aConsumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough, a trivial cubin
loaded, launched, synchronized, and produced the expected output on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story. That
means GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story is not rejecting all datacenter-targeted Blackwell SASS at the
decoder boundary. It also means at least one visible part of the architecture
split is enforced in software.
That is already interesting enough. It proves a real loader/runtime fact.
But it does not prove any of the following:
- that TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingAbout: GB10 tensor-path proof summary Reference: NVFP4 inference on GB10 Reference: GB10 stack parity is physically present on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story,
- that
tcgen05.mmaQuick term guidetcgen05.mmaThe Blackwell tcgen05 matrix-multiply-accumulate instruction family. On GB10, the public evidence still stops before a clean execution-grade proof.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails is physically present on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story, - that a deeper helper path would complete correctly if the driver accepted the image,
- or that a capability signal in the driver should be read as a silicon guarantee.
The baseline result is proof of a baseline result. Reverse-engineering gets messy exactly when we try to cash it out for more than that.
In glossary terms, CUDAQuick term guideCUDANVIDIA's GPU programming stack: compiler, runtime, driver, libraries, and kernel toolchain used by CUDA training and inference lanes.GroundingAbout: GB10 journey About: XLA vs CUDA stack decisions History: GB10 tensor-path proof summary is only the runtime submission layer here, TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingAbout: GB10 tensor-path proof summary Reference: NVFP4 inference on GB10 Reference: GB10 stack parity is the tensor-memory space named by the SM100Quick term guidesm_100Baseline Blackwell compiler target name in NVIDIA's architecture vocabulary, distinct from the architecture-specific and family-specific targets used elsewhere in the GB10 lane.GroundingAbout: GB10 tensor-path proof summary History: GB10 journey Example: GB10 cubin patch repro path, and tcgen05.mmaQuick term guidetcgen05.mmaThe Blackwell tcgen05 matrix-multiply-accumulate instruction family. On GB10, the public evidence still stops before a clean execution-grade proof.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails is the matrix instruction we still have not shown completing on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story.
The four software gates before execution
The GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story lane exposed four separate gating surfaces before a
tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough-oriented kernel reached anything close to normal execution:
- ELF architecture validation in
e_flagsQuick term guidee_flagsThe ELF header field whose low architecture bits are rewritten from sm_100a to sm_121a in the public GB10 arch-patch lane.GroundingAbout: sm_100a cubin patch repro Example: baseline arch-patch proof sample Reference: ELF arch-field patcher - reserved weak symbol handling
.nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingExample: GB10 gate repro Example: GB10 gate walkthrough Reference: kernel capability-record patcher per-kernel capability records.nv.capmercQuick term guide.nv.capmercObserved section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.GroundingHistory: libcuda patch lane Example: GB10 gate repro Example: GB10 gate walkthrough plus.nv.merc.rela.*Quick term guide.nv.merc.rela.*Observed companion metadata section family that appears with .nv.capmerc in the deeper GB10 gate boundary.GroundingExample: GB10 gate walkthrough Reference: GB10 tensor-path proof summary observed integrity-protected capability metadata
This gate structure is the backbone of the whole story because it explains why a driver-visible path can be so misleading.
In public-safe terms, gate 1 is the cubin's declared architecture, gate 2 is
reserved shared-memory bookkeeping accepted by the loader, gate 3 is mutable
.nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingExample: GB10 gate repro Example: GB10 gate walkthrough Reference: kernel capability-record patcher per-kernel capability records, and gate 4 is the deeper observed .nv.capmercQuick term guide.nv.capmercObserved section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.GroundingHistory: libcuda patch lane Example: GB10 gate repro Example: GB10 gate walkthrough
plus .nv.merc.rela.*Quick term guide.nv.merc.rela.*Observed companion metadata section family that appears with .nv.capmerc in the deeper GB10 gate boundary.GroundingExample: GB10 gate walkthrough Reference: GB10 tensor-path proof summary block where the checked-in lane still stops. The public
patchers that move the first three layers are the ELF arch-field
patcher, the reserved
shared-memory symbol
patcher, and the
kernel capability-record
patcher. The staged
walkthrough is GB10 gate
matrix together with
the GB10 repro
walkthrough.
Early gates are easy to misread. If you patch the architecture bytes and the error disappears, it is tempting to say "the driver was the problem." If you patch symbol handling and the error changes again, it is tempting to say "we are getting closer." If you strip or rewrite metadata and the failure moves again, it is tempting to say "the feature must be there."
What is really happening is subtler. You are moving through layers of software policy. You are learning how the driver packages and protects a capability. You are not yet learning whether the underlying hardware truly exposes the end state you care about.
Gate 4 is the most important example. Once the kernel image depends on integrity-protected capability metadata, the presence of that metadata tells you that NVIDIA cares about the capability boundary. It does not tell you whether the capability on the far side is guaranteed to work on this SKU.
The individual probe names matter for the same reason. tcgen05.allocQuick term guidetcgen05.allocDocumented tcgen05 allocation-side instruction family member used by the checked-in GB10 allocation probe.GroundingAbout: GB10 cubin patch repro Example: GB10 repro walkthrough Example: GB10 gate walkthrough is the
small alloc-side probe, tcgen05.ldQuick term guidetcgen05.ldDocumented Tensor Memory load-side instruction in the tcgen05 family, kept separate from alloc and mma in the checked-in GB10 probes.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough is the load-side probe, and TMA multicastQuick term guideTMA multicastThe cluster-scoped cp.async.bulk.tensor multicast form that attempts one tensor copy into shared memory of multiple CTAs in a cluster.GroundingAbout: GB10 tensor-path proof summary History: GB10 journey Example: GB10 repro walkthrough
is the clustered copy family that can appear in the same exploratory bundle.
Seeing those names together tells you the software stack knows how to describe
several advanced surfaces. It still does not prove that GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story cleanly executes
the whole family.
The same is true for reservedSmemQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingExample: GB10 gate repro Example: GB10 gate walkthrough Reference: reserved shared-memory symbol patcher: those weak undefined shared-memory symbols
are an earlier loader bookkeeping gate, not proof that the later tensor path is
already executing correctly on hardware.
Why helper cubins and routing tables can overstate support
This was the second misleading signal in our GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story work.
When a driver ships helper assets, architecture tables, or wrapper-selection logic that clearly knows about a more advanced path, engineers naturally treat that as evidence that the hardware path is nearby. Sometimes it is. But "the driver knows how to talk about the path" is still not the same claim as "this SKU can execute the path."
In plain English, helper cubins and routing tables mean the software stack knows how to name, package, or dispatch a richer capability. They do not mean the hardware has already been shown to complete that capability cleanly on this SKU. That is the whole trap.
Wrapper availability is the same kind of signal. If the driver can locate a helper cubin or wrapper path, that proves the submission machinery is richer than a flat "unsupported GPU" branch. It still does not prove the end-state instruction family completed correctly on hardware.
The cleanest way to think about it is to separate three layers:
- Routing knowledge: the software knows names, tables, helper assets, or dispatch rules for a capability.
- Submission knowledge: the software can package and submit something that looks like that capability.
- Runtime proof: the hardware actually completes the exact instruction family in a stable, intended way.
Most false positives in capability research happen when layer 1 or layer 2 gets mistaken for layer 3.
GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story is full of opportunities for that mistake because the surrounding software stack shares a family resemblance with datacenter Blackwell. That family resemblance is real. It is just not sufficient proof.
What the stronger patching path did and did not show
One of the more interesting follow-on directions was the deeper libcudaQuick term guidelibcudaThe user-space NVIDIA driver library that owns module load, metadata validation, and the helper-cubin patch lane in the GB10 experiments.GroundingHistory: libcuda patch lane Example: GB10 cubin patch repro Example: driver signal vs runtime proof sample patch
path: if we patch far enough to bypass validation or helper selection, can we
get the driver to submit the kernel in a way that tells us something
conclusive?
That is a valid research question. It is not the same thing as a finished proof.
Even the strongest version of that patching story only tells us that byte-level or table-level driver controls matter a lot. It may get us from "immediate rejection" to "deeper submission behavior." It may even get us to a launch or a hang. But a launch without a clean, stable, intended completion is still not the same as proving shipping silicon support.
That is why we are deliberately framing the deep patch path as a research lane, not a publication-grade capability receipt.
The copied-driver follow-on is documented separately in GB10 libcuda driver patch lane and why it still is not silicon proof and the driver patch lane walkthrough.
The practical rule for future bring-up
The rule we want future GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story work to follow is simple:
Separate these statements every time you write them down:
- the driver accepted the image,
- the driver routed the request through a richer helper path,
- the driver submitted something that looked closer to the desired capability,
- the exact instruction family completed and produced the expected result on hardware.
Only the last statement is silicon proof.
That rule sounds obvious. In practice it is the thing most likely to erode during a long debugging session because every incremental patch feels like momentum. The more exciting the path looks, the more aggressively you need to defend the distinction.
What this means for GB10 public claims
For public documentation we are intentionally using the stricter reading.
The driver-visible evidence on GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story is real and worth publishing. It shows that:
- product gating is layered;
- baseline
sm_100aQuick term guidesm_100aDatacenter Blackwell cubin target used by GB100/B200-class paths and by the source cubins in the public GB10 arch-patch repro.GroundingAbout: GB10 tensor-path proof summary Example: sm_100a cubin patch repro Example: GB10 repro walkthrough SASS is not categorically rejected; - capability metadata is protected much more heavily than a trivial unsupported-path check would require;
- software-visible signs of a datacenter path can persist even when a clean end-to-end proof is still missing.
What it does not justify is a claim that GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story has already been shown to
expose working B200/GB100 tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingAbout: GB10 tensor-path proof summary Example: GB10 cubin patch repro Example: GB10 repro walkthrough parity.
That is why the safer operational conclusion is still the one we use elsewhere on this site: treat GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story as a different kernel target and keep datacenter-only assumptions off until a real receipt exists. The same separation shows up in GB10 stack parity for MegaCpp and vLLM GB10 overlay and disabled paths, where software-visible paths are not allowed to masquerade as end-state support.
Why we are publishing examples instead of a triumphalist claim
The examples that accompany this post are intentionally boring in the best sense.
They publish the pieces we can actually defend:
- the baseline arch-patch probe in the baseline arch-patch proof sample;
- the driver-signal-vs-runtime-proof distinction in the driver signal versus runtime proof sample;
- the four-gate matrix in the compact gate-walk
mirror that shows
where the
tcgen05path still stops.
That is more valuable than a dramatic headline because it gives future engineers reusable decision tools instead of inheriting a claim they now have to unlearn.
The rule here is the same rule that should govern any GPU capability bring-up:
If you want to claim silicon support, publish the runtime proof. If you only have richer routing evidence, publish it as routing evidence.
GB10Quick term guideGB10Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.GroundingAbout: GB10 Stack Parity for MegaCpp: Torch 2.13 cu132, GCC 15, CUDA 13.2, and the Nightly Constraint About: GB10 tensor-path proof summary History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story is interesting enough without collapsing those two categories.
Frequently asked questions
Does a successful arch-patched cubin prove tcgen05 support on GB10?+
Why publish driver-signal examples if they are not silicon proof?+
What do the public patch scripts actually move?+
reservedSmemQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane. symbols. The kernel capability-record patcher trims selected mutable .nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary. per-kernel records at gate 3. The deeper copied-driver tooling belongs to the driver patch lane walkthrough and the companion GB10 libcuda driver patch lane and why it still is not silicon proof, not to the public-safe cubin patch lane.What counts as runtime proof here?+
Terms used in this article
Start here for quick definitions, then follow the linked posts for deeper context.
The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.
The user-space NVIDIA driver library that owns module load, metadata validation, and the helper-cubin patch lane in the GB10 experiments.
Observed weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.
Observed per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.
Consumer Grace Blackwell GB10 / DGX Spark bring-up lane used to separate driver-visible gates, patched cubin signals, and real execution proof.
The ELF header field whose low architecture bits are rewritten from sm_100a to sm_121a in the public GB10 arch-patch lane.
Baseline Blackwell compiler target name in NVIDIA's architecture vocabulary, distinct from the architecture-specific and family-specific targets used elsewhere in the GB10 lane.
Observed companion metadata section family that appears with .nv.capmerc in the deeper GB10 gate boundary.
Datacenter Blackwell cubin target used by GB100/B200-class paths and by the source cubins in the public GB10 arch-patch repro.
Consumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.
Documented tcgen05 allocation-side instruction family member used by the checked-in GB10 allocation probe.
Documented Tensor Memory load-side instruction in the tcgen05 family, kept separate from alloc and mma in the checked-in GB10 probes.
The Blackwell tcgen05 matrix-multiply-accumulate instruction family. On GB10, the public evidence still stops before a clean execution-grade proof.
The cluster-scoped cp.async.bulk.tensor multicast form that attempts one tensor copy into shared memory of multiple CTAs in a cluster.
Observed section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.
Blackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.
NVIDIA's GPU programming stack: compiler, runtime, driver, libraries, and kernel toolchain used by CUDA training and inference lanes.
How MegaCpp stabilized a GB10-oriented vLLM lane with an on-disk overlay, text-only model registration, and a deliberate keep-disabled list for…