What Our GB10 Experiments Actually Prove About Blackwell Consumer vs Datacenter Tensor Paths

The practical question behind these 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough experiments was simple: when a Blackwell-datacenter cubin fails 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough, is that because the silicon physically lacks the path, because the driver blocks it, or because we are mixing evidence from several different layers and telling ourselves a cleaner story than the data supports?

The answer is narrower than some early drafts made it sound.

Read this page in three layers. First, the narrow positive result is 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough can accept and execute at least some 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source-targeted SASS after a small architecture-field rewrite in the cubin. Second, tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity and TMA-adjacent probes hit later driver-side gates that the baseline arithmetic kernel never touches. Third, none of that is the same as proving 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough physically executes 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source the way B200 or GB100 does.

The target names matter before the deeper terms do. In NVIDIA's compiler vocabulary, 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.GroundingHistory: GB10 journey Example: GB10 cubin patch repro Reference: GB10 stack parity is the baseline Blackwell target name, 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source is the architecture-specific datacenter target used by B200-class cubins, sm_120fQuick term guidesm_120fFamily-specific consumer Blackwell compile target used when kernels should target family-common features without pinning to one exact device label such as sm_121a.GroundingAbout: GB10 stack parity History: GB10 journey Example: GB10 cubin patch repro is the family-specific consumer-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.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity is the exact 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough device target the driver reports on DGX Spark. These labels share the Blackwell brand, but they do not promise the same kernel contract. For the term map, use MegaCpp model glossary. For the staged 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough explanation, continue with Reproducing the sm_100a to sm_121a cubin patch on GB10 and the checked-in GB10 repro walkthrough.

For first touch, tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity means 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.GroundingHistory: GB10 journey Example: GB10 cubin patch repro Reference: GB10 stack parity/datacenter Tensor Core Generation 5 family, and TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample is the tensor-memory space those instructions are written around. That is why a receipt saying only "Blackwell" is not enough: the real question is whether the receipt is about a consumer warp-level path or about the datacenter TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample-coupled path. The public ladder here is older warp-level mma.syncQuick term guidemma.syncThe older warp-level matrix-multiply path that predates Hopper WGMMA and the SM100 tcgen05 family.GroundingHistory: GB10 journey Reference: NVFP4 inference on GB10 Reference: FA4 catalog on Blackwell, then Hopper wgmmaQuick term guideWGMMAHopper's warpgroup matrix-multiply path between the older mma.sync lane and Blackwell's tcgen05 family.GroundingHistory: GB10 journey Reference: upstream PR: TileLang and Megatron Reference: TileLang TMA bulk-copy sample, then Blackwell 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.GroundingHistory: GB10 journey Example: GB10 cubin patch repro Reference: GB10 stack parity tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity.

One reader-first decode order helps. Start with the target label (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.GroundingHistory: GB10 journey Example: GB10 cubin patch repro Reference: GB10 stack parity, 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source, sm_120fQuick term guidesm_120fFamily-specific consumer Blackwell compile target used when kernels should target family-common features without pinning to one exact device label such as sm_121a.GroundingAbout: GB10 stack parity History: GB10 journey Example: GB10 cubin patch repro, sm_121aQuick term guidesm_121aConsumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity). Then ask whether the probe is about the datacenter TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample-coupled family (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, 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.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity, 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source) or about the adjacent clustered-copy question (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.GroundingHistory: GB10 journey Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source). Only after that should you drop to metadata names like reservedSmemQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough, .nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough per-kernel records, and .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.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro. That is also the order used by GB10 stack parity for MegaCpp, Training the MegaCpp SLM Ensemble on GB10, and the checked-in GB10 repro walkthrough.

For engineering purposes, that means the safe rule is still the conservative one: treat tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough as unavailable unless you have a real end-to-end execution result for that exact instruction family. Driver-visible hints are not enough.

The narrow positive result

The strongest positive result in 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough test lane is also the cleanest one. A trivial arithmetic cubin compiled for 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source loaded and ran 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough after only the low architecture bits in ELF 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 were rewritten 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source to sm_121aQuick term guidesm_121aConsumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity. The smallest checked-in restatement of that exact fact is baseline arch-patch proof sample.

The important part is not the patch itself. The important part is what happened after the patch:

[cuModuleLoadDataEx]    CUDA_SUCCESS
[cuLaunchKernel]        CUDA_SUCCESS
[cuCtxSynchronize]      CUDA_SUCCESS
# out[0..7]: 00000001 00000003 00000005 00000007 00000009 0000000b 0000000d 0000000f

That is enough to establish three things.

First, the user-space driver really does contain software gating at the architecture-identification layer. Second, 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough's instruction path will accept at least some SASS originally emitted for the datacenter Blackwell line. Third, the absence of an immediate Xid or hard decoder fault means "consumer Blackwell" and "datacenter Blackwell" are not separated by one single binary hardware switch at the very first instruction boundary.

That is a meaningful result. It is just not the same as "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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough has working 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source." The same distinction is why NVFP4 inference, the FA4 catalog on Blackwell, and the driver signal versus runtime proof sample all keep driver-visible hints separate from execution-grade proof.

The four software gates we found

Once we moved from a baseline arithmetic kernel to tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity / TMA-oriented probes, the driver path turned out to be layered.

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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough smtest lane consistently exposed four gates before the cubin reached a usable execution state:

1. ELF 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 architecture validation
2. loader-side reserved shared-memory symbol bookkeeping such as .nv.reservedSmem.offset0Quick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough and .nv.reservedSmem.capQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough
3. mutable per-kernel capability metadata in .nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough records
4. the observed deeper integrity-protected capability block surfaced as .nv.capmerc.textQuick term guide.nv.capmercObserved section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro records 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.GroundingAbout: GB10 driver gates warning Example: GB10 gate walkthrough

The important engineering lesson is not just that there are four gates. It is that each gate answers a different question.

Gate 1 is the simple identity check: does the cubin even claim the right target? Gate 2 is loader-side reserved shared-memory bookkeeping: the driver is still rejecting missing reservedSmemQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough symbols before anything like normal execution begins. Gate 3 is mutable per-kernel capability metadata in .nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough records. Gate 4 is the observed deeper integrity-protected capability block surfaced as .nv.capmerc.textQuick term guide.nv.capmercObserved section-name family for the deeper integrity-protected metadata boundary where the public-safe GB10 gate walk still stops.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro records 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.GroundingAbout: GB10 driver gates warning Example: GB10 gate walkthrough. That is the point where naive cubin surgery stops being enough.

For the staged explanation, continue directly with Reproducing the sm_100a to sm_121a cubin patch on GB10, the checked-in GB10 repro walkthrough, and the compact GB10 gate matrix. The reserved shared-memory symbol patcher only proves that the weak undefined reservedSmemQuick term guidereservedSmemObserved weak undefined reserved shared-memory symbols that mark the earlier loader-side bookkeeping gate in the GB10 lane.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough layer is mutable enough to move the failure. The kernel capability-record patcher does the same for selected .nv.infoQuick term guide.nv.infoObserved per-kernel metadata records edited in the GB10 repro lane before the later integrity-protected boundary.GroundingAbout: GB10 driver gates warning Example: GB10 gate repro Example: GB10 gate walkthrough per-kernel records. Neither script turns gate 4 into runtime proof.

That wording is intentionally conservative. Public NVIDIA docs explain tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity, TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample, and clustered TMA instructions, but this article grounds .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.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro 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.GroundingAbout: GB10 driver gates warning Example: GB10 gate walkthrough only against the checked-in gate walk and the literal section names visible in the public bundle. We do not present .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.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro as an official NVIDIA API surface; the claim here is only that this observed signed capability boundary is where the public-safe repro still fails.

External public evidence reinforces that restraint rather than weakening it. We did not find a public, reproducible 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough receipt that carries a tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity / TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample probe through the later .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.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro-named boundary into a stable 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source execution result. Negative evidence is not proof of impossibility by itself, but it removes the usual shortcut where someone else's public repro would justify stronger language than our checked-in bundle supports.

This matters because public discussions about 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough support often collapse all of these layers into one sentence like "the driver blocks it" or "the hardware lacks it." That hides the real structure of the problem. Some things are byte-patchable. Some things are not. And the fact that you can move through the earlier gates does not tell you what would happen if the exact tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity kernel reached a fully valid submission state.

What the tcgen05 probes actually hit

The conservative 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough evidence stops at gate 4.

The checked-in probes keep the hard terms separate on purpose because they are not one bucket. 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 asks whether the TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample allocation side can even advance past the early metadata gates. 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.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity asks about the load side of that same family. 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source asks about the matrix instruction itself. The clustered TMA multicast probeQuick 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.GroundingHistory: GB10 journey Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source is adjacent rather than identical: it is one local copy-path probe inside the broader 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.GroundingHistory: GB10 journey Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source family, not the name of the whole family and not a successful 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source receipt in disguise. None of those surfaces produced a clean public-safe 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough execution receipt. The isolated 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 lane, for example, can be moved through the earlier gates and still ends in CUDA_ERROR_INVALID_IMAGE once the 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.GroundingAbout: GB10 driver gates warning History: libcuda patch lane Example: GB10 gate repro-named integrity layer becomes decisive.

For non-kernel readers, the split is simpler than the mnemonics make it sound. 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 smallest "can the TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample-side setup even begin?" question. 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.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity is the "can data be pulled through that setup?" follow-on. 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source is the actual matrix-operation question. 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.GroundingHistory: GB10 journey Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source is the neighboring clustered-copy question rather than the matrix path itself. The checked-in bundle separates them because moving one question forward does not automatically answer the others.

The important split is that minimal TMEM allocation probe source isolates 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, while full GB10 tensor-path probe source keeps 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.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity, 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source, and the clustered TMA multicast probeQuick 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.GroundingHistory: GB10 journey Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source separate in the broader bundle. For the staged command flow, continue with the checked-in GB10 repro walkthrough and GB10 gate matrix.

That is the key distinction that older drafts blurred.

There is a big difference between:

• proving the driver can be pushed farther than its default routing policy,
• proving a helper or wrapper path exists inside libcuda,
• proving a kernel can be submitted and hang,
• and proving 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source is a stable, physically present, usable execution path 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough silicon.

Only the last one would justify a public claim of B200-style tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity availability. We do not have that proof.

Why older stronger claims are stale

Some of our earlier exploratory notes leaned too hard on the most exciting interpretation of the data. That is normal in a live reverse-engineering session and unacceptable in a publication.

The strongest overreach looked like this:

• a helper or routing patch can be made to reach deeper driver paths,
• therefore the silicon probably carries the full capability,
• therefore 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough effectively "has" the datacenter path if we patch enough bytes.

That leap is too large.

A more honest reading is:

• the software stack clearly contains layered product gating,
• at least some datacenter-targeted SASS decodes and executes 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough,
• the driver ships enough capability-related machinery to make the path look tantalizingly close,
• but the publication-grade tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity proof is still missing.

That is why we are treating the earlier stronger wording as stale. It was useful as a research hypothesis. It is not the standard we want attached to a public article or a customer-facing example repo.

Consumer Blackwell vs datacenter Blackwell

What should an engineer conclude from this if they just want to ship kernels?

The practical conclusion is not mysterious:

• 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough is not a small B200.
• Driver-visible datacenter artifacts do not make it one.
• If a path depends on TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample, tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity, or other datacenter-only assumptions, you should 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough as a separate target with its own kernel contract.

That conclusion matches the rest of 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough bring-up story across this site. In inference, we already 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough's OMMA-based FP4 lane as real while keeping TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample-coupled paths off. In the FA4 catalog, we already gate 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough separately from the 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source line. The tensor-path experiments fit the same pattern: shared branding, partial decode overlap, different operational contract.

The split is wider than one instruction mnemonic. The same research lane that keeps tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity in the unsupported bucket also keeps clustered multi-CTA copy assumptions and CTA-pair / 2-SM tensor cooperation on the datacenter side of the fence. In reader terms: if a kernel design depends on clustered delivery or two-SM tensor cooperation as part of its normal contract, treat that as another sign you are outside the safe 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough subset rather than one patch away from a supported path.

That is also why public B200 or GB100 capability tables are only partial context here. They can be perfectly accurate descriptions of the datacenter contract and still tell you nothing about whether 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough inherits TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample, clustered TMA delivery, or 2-SM tensor cooperation. For 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough, the trustworthy rule is still receipt first: use the datacenter docs to decode the vocabulary, then look for a 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough execution-grade result before treating any of those paths as available.

Another way to keep the boundary straight is TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample versus the narrower consumer low-precision lane. On B200-class 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.GroundingHistory: GB10 journey Example: GB10 cubin patch repro Reference: GB10 stack parity, NVIDIA's public docs place tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity on the TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample-coupled datacenter path. On 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough side of this site, the low-precision receipts we do trust live in a different contract: the consumer sm_120fQuick term guidesm_120fFamily-specific consumer Blackwell compile target used when kernels should target family-common features without pinning to one exact device label such as sm_121a.GroundingAbout: GB10 stack parity History: GB10 journey Example: GB10 cubin patch repro / sm_121aQuick term guidesm_121aConsumer Blackwell cubin target used by GB10/DGX Spark and the patched destination in the public arch-field repro.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity lane sometimes described in neighboring posts with local shorthand like OMMA, plus the NVFP4 and FA4 examples that stay inside that consumer staging model. Those receipts matter, but they answer a different question than tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity, which is why NVFP4 Inference for the MegaCpp SLM Ensemble, The FA4 Catalog on Blackwell, and GB10 stack parity for MegaCpp can all be true without upgrading this article into TMEMQuick term guideTMEMBlackwell tensor-memory scratch storage used by datacenter-oriented tensor paths; the public GB10 evidence treats it as unavailable.GroundingNVFP4 inference on GB10 GB10 stack parity TileLang TMA bulk-copy companion sample parity.

The main correction here is about proof discipline. A capability table, helper cubin, or partially patched submission path is not a shipping contract.

The conservative public rule

The right public rule is stricter than the most optimistic early research note.

If you are writing documentation, examples, or runtime policy for 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough:

• you may say that some 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source SASS executes 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough after an arch-field rewrite;
• you may say that multiple driver-side gates exist before tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity probes can run;
• you may say that driver-visible capability machinery can make unsupported paths look deceptively close;
• you should not say 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough has proven working 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.GroundingExample: sm_100a cubin patch repro Reference: GB10 claim-scope guardrails Reference: full GB10 tensor-path probe source parity with B200 or GB100.

That is also why the new public examples for this topic focus on the baseline arch-patch probe, the gate matrix, and the difference between a software-visible signal and runtime proof. Those are the parts we can defend cleanly.

What we are publishing instead

For this topic we are publishing three things and drawing one line.

The three things:

• a compact baseline probe showing what the positive 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.GroundingExample: sm_100a cubin patch repro Example: GB10 repro walkthrough Reference: full GB10 tensor-path probe source result really means;
• a compact example showing why driver-visible support is not runtime proof;
• a near-copy gate-matrix example showing exactly where the tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity path still stops.

The line:

We are not publishing "6-byte patch unlocks tcgen05Quick term guidetcgen05The Blackwell tensor-generation instruction family that covers alloc, load, and mma paths beyond the older dense consumer path.GroundingExample: GB10 cubin patch repro Example: GB10 repro walkthrough Reference: GB10 stack parity 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 History: Training the MegaCpp SLM Ensemble on GB10: a Grace Blackwell war story Example: GB10 repro walkthrough" as a settled statement.

That line is worth keeping. It saves future engineers from inheriting an evidence problem disguised as a success story.

This narrower claim lines up with Training the MegaCpp SLM Ensemble on GB10, NVFP4 Inference for the MegaCpp SLM Ensemble, and The FA4 Catalog on Blackwell: consumer Blackwell is a real target, but only on the exact kernel and precision paths that have execution-grade evidence rather than suggestive driver metadata.