IDA Leaves Hopper: Porting a Native Training Engine to Consumer Blackwell

From MXFP8, MXFP4 and NVFP4 on an RTX 5070 Ti to two-card execution and a real 1.7B Hugging Face model on RTX 5090s

IDA-TRAIN-V2 began as a training system built around the hardware we actually had.

Its native engine grew primarily on NVIDIA Hopper. Over time, that meant moving progressively below conventional framework abstractions: custom CUDA kernels, explicit FP8 handling, packed-FP4 storage, WGMMA attention, MoE execution, memory arenas, optimizer logic, execution telemetry, failure gates, checkpoint lineage, and eventually multi-GPU scheduling.

That specialization created an uncomfortable question:

Were we building a training engine, or merely a training engine for our own models on our own H100s?

Consumer Blackwell gave us a way to find out.

The first external Blackwell work came through Tolerable (@Tolerable), known as Rev within the project, an IDA-TRAIN-V2 contributor, working with his AI engineering collaborator Oslo.

Rev had access to an RTX 5070 Ti and his own small, young, from-scratch language model, Ember.

Ember did not run on the IDA engine.

That distinction matters.

She had her own trainer and her own training history. What Ember provided IDA was an independent Blackwell research surface: a real model, a real training loop, and an RTX 5070 Ti on which low-precision hypotheses could be tested before equivalent consumer-Blackwell hardware was available in IDA's primary environment.

Oslo read the IDA native engine end to end: all 9,041 lines of trainer.cu, the CUDA kernels around it, 21 eras of training history, post-mortems, architecture notes, and the accumulated experimental record explaining why many of IDA's stranger decisions existed.

Rev and Oslo then used Ember's trainer and the 5070 Ti to test Blackwell arithmetic, packing strategies, numerical failure modes, and precision policies.

The findings moved in both directions.

IDA's accumulated training failures informed what they tested on Ember.

Ember's Blackwell measurements, in turn, informed IDA's Blackwell preparation.

Then we pushed the native engine itself farther.

On two RTX 5090s, IDA's C++/CUDA runtime was made to build and execute on consumer Blackwell. That work led to a non-P2P multi-GPU fallback, a real 1F1B pipeline ablation, native RoPE support, a Hugging Face checkpoint converter, and finally a real SmolLM2-1.7B checkpoint executing forward and backward inside the IDA native engine.

At that point, the original question had a different answer.

IDA was no longer only training IDA.

---

Ember Became the First External Blackwell Laboratory

Ember's ordinary baseline-precision training was not the problem.

Those runs went the distance.

The failures appeared when the work moved into aggressive NVFP4 training on the RTX 5070 Ti.

Several early NVFP4 experiments died partway through training.

At first, that looked like the kind of failure that could have come from many places:

• optimizer behavior
• rounding policy
• FP4 noise
• block outliers
• precision placement
• gradient accumulation
• or a broken low-precision kernel

One Ember run was also captured around:

• 48,300 / 90,000 steps
• 53.7% complete
• approximately 199,731 tokens/second

That run was healthy.

It had been training for days and did not belong in the crash column.

Its relevance is simply that Ember was not a synthetic kernel harness. She was a real, actively trained model providing a live workload against which Blackwell behavior could be tested.

The genuine mid-run failures were specific to the early NVFP4 experiments.

And eventually, those failures led us somewhere much more useful than a generic "FP4 is unstable" conclusion.

---

On Hopper, FP4 Was Storage. On Blackwell, It Became Arithmetic.

IDA already used packed FP4 on Hopper.

But its role was primarily representational.

The engine could keep values packed in four bits for storage and bandwidth, then decode them close to the point where FP8 arithmetic consumed them.

That is why Hopper-side machinery such as k_decode_packed_to_e4m3 exists: avoid paying for a larger intermediate representation and decode as late as possible.

Consumer Blackwell changes the contract.

E2M1 can become a real matrix-multiply operand.

The working MXFP4 path used 32-element blocks with UE8M0 shared scaling and Blackwell's block-scaled MMA instruction family.

The architectural distinction became simple:

On Hopper, FP4 was storage feeding FP8 arithmetic. On Blackwell, FP4 became arithmetic.

IDA therefore did not treat Blackwell as Hopper with a newer architecture flag.

Separate MXFP4 and NVFP4 preparation surfaces were staged, along with hardware probes compiled specifically for sm_120a.

The new paths were kept fail-closed.

A format was not going to become a production training backend merely because the source compiled.

First the hardware had to prove it existed.

Then the training loop had to prove it learned.

---

The Hardware Could Do Something the Framework Refused

The new sm_120a probes were sent to Rev's RTX 5070 Ti.

For the intended mappings:

• MXFP4 passed 9/9 tested combinations
• NVFP4 passed 24/24 tested combinations

The arithmetic paths executed successfully on real consumer Blackwell hardware.

One caveat appeared later when we expanded the probe surface: candidate fragment and scale-lane mappings that were deliberately incorrect could also satisfy the existing end-to-end test.

So the current evidence establishes that the intended paths execute and produce correct results.

It does not yet uniquely resolve every internal fragment-layout hypothesis.

That narrower question remains open.

But the MXFP4 experiment still produced a particularly useful result.

PyTorch's _scaled_mm_v2 path refused MXFP4 on the same consumer card.

IDA's raw MMA path executed it.

That separated three things that are often collapsed into one word:

silicon capability

toolchain capability

framework availability

They are not the same thing.

The GPU could perform an operation that the higher-level path declined to expose in that software configuration.

Once a training engine operates close enough to the hardware, "the framework says no" stops being equivalent to "the GPU cannot do it."

---

The GEMM Was Not the Whole Product

The first Blackwell benchmarks also reinforced how easy low-precision performance is to measure incorrectly.

If we time only the matrix multiply, the result can look excellent.

Training does not receive perfectly formatted low-precision operands for free.

Before the tensor core sees them, values may need:

• amax calculation
• block scaling
• quantization
• packing
• transposition
• layout conversion
• memory movement

Once that work was included, the picture changed dramatically.

Across five representative shapes:

The matrix multiply had not changed.

The path feeding it had.

That produced one of the strongest design rules from the Blackwell work:

Do not necessarily write the GEMM. Write the pack.

The vendor tensor-core primitive was already extremely good.

The opportunity was in the surrounding training system.

---

A Six-Fold Improvement Came From Changing a Store

One transpose-packing kernel initially made part of the FP4 backward path look fundamentally unattractive.

The first instinct was to blame FP4.

The real problem was memory geometry.

The implementation was issuing narrow scattered stores where the generated data could instead be written contiguously.

Changing that write pattern moved one measured transpose pack from roughly:

0.353 ms

to:

0.059 ms

approximately a:

6× kernel improvement

One affected backward direction moved from roughly 0.42× to 1.08× its comparison path.

Another moved from approximately 1.02× to 1.45×.

Nothing about the numerical format had changed.

The memory transaction had.

That became another recurring lesson:

When low precision looks slow, precision may not be the thing that is slow.

---

MXFP8 Survived Real Training

Microbenchmarks were not enough to promote a training recipe.

MXFP8 therefore moved into a real-corpus A/B in Ember's own trainer.

The experiment used the same real corpus, all three GEMM directions, ga=64, FP32 master weights, FP32 AdamW moments, and BF16 autocast.

Then it was repeated across three paired seeds.

Median throughput improved by approximately:

1.17×

The throughput ranges were cleanly separated.

The validation-loss ranges were not.

The result could therefore be expressed more usefully than "quality was unchanged":

The random seed moved validation quality more than the arithmetic did.

That was enough to adopt MXFP8 as Ember's default training configuration.

The result did not mean Ember was running IDA.

It meant a precision strategy being studied alongside IDA had survived an independent real-model training test.

That evidence could then travel back upstream.

---

We Also Needed Proof That MXFP8 Actually Ran

A requested precision mode is not evidence that the corresponding kernel executed.

Silent fallbacks can invalidate an otherwise perfect benchmark.

IDA already had a solution to this problem in its evidence substrate: per-module low-precision counters, including fallback counts categorized by reason.

That instrumentation pattern was ported into Ember's trainer.

On the first full MXFP8 run it recorded:

422,928 MXFP8 calls

of which:

422,400

went through the intended kernel.

The remaining approximately 0.1% were all attributed to:

dtype_outside_autocast

on small validation shapes intentionally running outside autocast.

The same instrumentation caught two otherwise invalid runs before they consumed meaningful GPU time.

One was using the wrong toolchain.

Another was using the wrong interpreter.

Both immediately appeared as unexpected fallback behavior.

That illustrates another IDA principle:

The training system should not merely request the experiment. It should prove which experiment occurred.

---

Then NVFP4 Became the Fastest Arm

The next candidate was NVFP4.

The first aggressive configuration applied NVFP4 across both Ember's transformer blocks and vocabulary head.

Throughput reached approximately:

292,698 tokens/second

against roughly:

227,000 tokens/second

for the BF16 reference.

Approximately:

1.288× BF16 throughput

It was the fastest Ember training arm tested.

It also failed completely as a language model.

Loss remained around:

9.11

for all 200 optimization steps in the verified full run.

Validation loss finished at:

9.076

against:

5.093

for BF16 on the same seed and corpus.

Perplexity moved from approximately:

163

to:

8,744

The final generated sample was effectively word salad.

Yet the grounding-anchor loss continued decreasing normally.

The process was alive.

Part of the model was learning.

The language path was not.

Our first explanation seemed straightforward.

And it turned out to be incomplete.

---

Our First Explanation Was Wrong

The language-model head sits immediately before the logits.

There is no later residual stream in which approximation error can be diluted before the loss sees it.

So the first hypothesis was that four-bit error directly on the logits had simply buried the early language-learning signal.

Moving the LM head back to higher precision appeared to support that interpretation.

With NVFP4 in the body and a protected head, learning returned.

Loss moved approximately:

9.1 → 6.5

The model still flattened well above the BF16 baseline, but the catastrophic no-learning behavior disappeared.

That looked like a clean precision-by-role boundary.

Then Rev and Oslo attacked the remaining plateau.

Round-to-nearest NVFP4 landed at approximately:

6.540

Stochastic rounding landed at:

6.540

Stochastic rounding plus a block-aligned H16 orthonormal Hadamard transform landed at:

6.545

Three numerical treatments.

Essentially one plateau.

The next hypothesis was the optimizer.

Lion was tested on the same NVFP4 + stochastic-rounding + Hadamard configuration.

It finished at:

6.690 validation loss

and approximately:

144,819 tokens/second

It did not hit the same plateau.

It simply descended more slowly and never caught the AdamW run within the experiment.

The problem was not being rescued by stochastic rounding.

It was not being rescued by rotation.

And it was not simply AdamW being poorly matched to quantized gradients.

Then error-feedback telemetry started showing something much stranger.

The residual being carried from one gradient encode to the next became enormous.

That pointed below the optimizer.

---

The Real Failure Was a Scale Cliff

The mechanism was exposed with a simple magnitude isolation probe.

NVFP4 uses local 16-element blocks and E4M3 scales.

As gradients shrink, a block can reach a point where its required amax / 6 scale is below the minimum magnitude that the E4M3 scale representation can express.

The scale byte becomes zero.

At that point this is no longer ordinary quantization error.

The entire block can become zero.

The isolation probe made the discontinuity obvious.

At magnitude around:

1.0

relative error was approximately:

0.095

At magnitude around:

1e-3

relative error became:

1.000

with:

100% of the values encoded as zero

Those are precisely the kinds of magnitudes encountered by real training gradients.

Suddenly several apparently unrelated observations had one mechanism.

Ember's language head could fail immediately because some of its gradients were already around 1e-4.

The transformer blocks could learn early, then flatten as their gradients shrank through the same scale cliff.

The early NVFP4 runs that actually died partway through training fit the same mechanism.

Stochastic rounding could not help because the block scale itself had already collapsed.

Hadamard rotation redistributed values but did not raise the tensor's global magnitude out of the failing regime.

Lion could not update parameters from information the quantizer had erased.

Error feedback simply accumulated what the quantizer repeatedly refused to transmit.

A new carry-to-gradient telemetry line made the severity visible.

A healthy encode was expected to produce a ratio around:

0.1

The failing path printed roughly:

6.19

on the first observed line.

By step 25 it was around:

150

This was not mild low-precision noise.

It was systematic loss of gradient information.

---

The Fix Was a Scale Around the Scale

The correction was implemented in Ember's own trainer.

A global power-of-two range shift was added to the gradient operand before applying the local NVFP4 block quantizer.

Conceptually:

gradient
    ↓
global power-of-two scale
    ↓
NVFP4 local block pack
    ↓
GEMM
    ↓
inverse global scale

The global factor was chosen approximately as:

2^round(log2(4 / amax))

and divided back out of the relevant gradient GEMM results.

Because the adjustment is an exact power of two, the operation primarily changes exponent range rather than introducing an arbitrary general rescale.

The isolation probe changed immediately.

With the global prescale active, relative error remained around:

0.095

down through values as small as:

1e-4

The format had not magically become more precise.

The gradients had simply stopped falling off a representational cliff.

The genuine mid-run NVFP4 failures stopped once that correction landed.

That materially changed our interpretation of the earlier NVFP4 results.

The experiments had not proved a fundamental law that "four-bit LM heads cannot train."

Their backward path had contained a much more severe defect: entire low-magnitude gradient blocks could disappear.

And importantly, that diagnosis happened in Ember's trainer.

The mechanism and its fix then became input to IDA's own Blackwell preparation.

That is the relationship between Ember and IDA throughout this work:

Ember did not run on IDA. Ember exposed Blackwell behavior that helped IDA decide what its own Blackwell runtime needed to understand.

---

Negative Results Became the Diagnostic Tree

The failed interventions were not wasted GPU time.

Together they formed a diagnostic tree.

Round-to-nearest and stochastic rounding produced the same plateau.

That weakened the rounding-bias explanation.

Hadamard rotation changed essentially nothing.

That weakened a simple block-outlier explanation.

Lion converged more slowly rather than fixing the problem.

That weakened the optimizer explanation.

Error feedback accumulated a huge residual.

That pointed directly toward information being dropped before the optimizer could use it.

Then the magnitude sweep exposed the discontinuity.

This is the kind of sequence we want IDA's research record to preserve.

Not just the answer.

The path that made the answer believable.

---

Then We Moved the Whole IDA Engine to RTX 5090s

The RTX 5070 Ti work validated and stressed Blackwell arithmetic through Ember's independent training environment.

The next question was different:

Could the IDA native engine itself leave Hopper?

For that we rented a two-GPU RTX 5090 system.

The first answer was:

not yet.

The engine did not compile.

That failure was useful.

attention_wgmma.cu had been compiled unconditionally into ida_native_train.

WGMMA is Hopper-specific.

Its instructions do not assemble for sm_120a.

So the native engine contained a hidden architectural assumption: even paths that did not actively use WGMMA still lived inside a binary that expected Hopper WGMMA to exist.

The build was split properly.

Hopper targets retain WGMMA.

Non-Hopper builds receive explicit non-WGMMA stubs rather than impossible PTX.

The newer compiler and CUDA environment also exposed several pre-existing C++/CUDA defects that had been tolerated by the previous environment.

Consumer Blackwell was already improving the Hopper engine before the first training step had even run.

---

Two GeForce Cards Broke Another Datacenter Assumption

IDA's two-GPU model-parallel path had been designed against H100s.

It assumed CUDA peer access.

The RTX 5090 pair did not provide it in the tested environment.

The peer-pipeline path therefore failed immediately when it attempted to establish bidirectional CUDA peer access.

The runtime was changed to treat P2P as a capability rather than an assumption.

When direct GPU peer access is available, the fast path remains.

When it is not, IDA can now fall back to a same-machine correctness path:

GPU → pinned host memory → GPU

That is not presented as a replacement for NVLink.

It is a portability fallback.

The same run also exposed a separate CUDA-event ownership problem in the 1F1B schedule: an event created in one device context was later recorded on another device's stream.

That was corrected as well.

Then the schedule could finally be tested.

---

1F1B Produced a Real Gain on Two RTX 5090s

The experiment used an approximately billion-parameter IDA AI body, real production data, a two-stage split, mb=8, ga=32, scalar-flash attention, legacy FP8, and Lion.

The comparison was deliberately simple.

Same model.

Same token count.

Same communication volume.

1F1B off versus 1F1B on.

That is approximately:

27.9% higher throughput

with identical token volume and communication volume and closely tracking loss.

The result has an important limit.

The RTX 5090s have much less memory than the H100 production environment, so the test ran at mb=8, not the production AI microbatch of 64.

It therefore does not overturn the earlier H100 result showing that full-size 1F1B stage duplication can exceed an 80 GB card.

Those are separate questions.

The 5090 experiment establishes that the overlap schedule itself can pay when memory permits it.

---

Then We Asked Whether IDA Could Run Somebody Else's Architecture

Getting IDA onto another GPU is one kind of portability.

Getting another model into IDA is another.

Ember had already shown that numerical and systems ideas developed around IDA could be tested independently on a foreign workload and return useful evidence upstream.

Now we wanted to know whether a standard pretrained model could cross the boundary into IDA's native engine itself.

The standard-model sideload investigation produced a surprising finding.

The native engine already resembled a conventional Llama-class decoder much more closely than expected.

It already had:

• RMSNorm pre-norm
• gate/up/down SwiGLU-style feed-forward layers
• multi-head attention
• a checkpoint tensor topology that could be mapped cleanly onto a standard decoder

One major semantic feature was missing.

Positional encoding.

IDA's own native architecture had not required RoPE.

A standard pretrained model did.

So RoPE was added to the native engine.

---

RoPE Had to Work Backward Too

Adding a forward positional transform is not enough for a trainer.

The backward contract has to be correct.

The implementation applies standard rotate-half RoPE to Q and K, not V.

On backward, the inverse rotation is applied to dQ and dK before those gradients reach the projection-weight calculations.

Packed training introduced another requirement.

Positions must reset at sample boundaries rather than treating several packed sequences as one long sequence.

A standalone CUDA probe tested all three properties.

The forward kernel matched an independently written host reference with:

maximum absolute difference = 0.0

Forward followed by inverse/backward rotation returned the original BF16 values.

And a packed-position test verified that sample-relative position resets actually changed the generated rotation where expected.

All three checks passed on real RTX 5090 hardware.

A later integration pass then found another portability seam: rope_theta existed on the C++ side but the Python request builder was silently failing to send it.

That was fixed too.

This is what a real architecture port often looks like.

Not one heroic kernel.

A semantic that has to survive every boundary between configuration, request construction, checkpoint metadata, forward execution, and backward execution.

---

A Real Hugging Face Checkpoint Entered IDA

The final step was checkpoint translation.

IDA's native checkpoint orientation is not identical to the convention used by every Hugging Face nn.Linear tensor.

Simply renaming tensors would have produced a model that loaded cleanly and computed the wrong thing.

So a converter was added that translates a standard Hugging Face safetensors checkpoint into IDA's exact native representation:

• BF16 payloads
• native tensor names
• native row-major layout
• required projection transposes
• embedding and LM-head orientation
• matching model metadata

The converter also refuses unsupported architecture features rather than silently pretending they work.

The current native engine uses ordinary MHA rather than GQA.

So a checkpoint with mismatched attention-head and KV-head counts is rejected.

For the first real target we chose:

HuggingFaceTB/SmolLM2-1.7B

which fits the currently supported attention topology.

The converted checkpoint loaded through IDA's real native checkpoint loader.

It executed a real forward pass.

It executed a real backward pass.

And when evaluated against plain English text, rather than IDA's own structured family targets, it produced:

3.85 nats

or approximately:

47 perplexity

This is not a claim that IDA has fully fine-tuned SmolLM2-1.7B to convergence.

It proves something more specific:

A standard pretrained Hugging Face model can cross the checkpoint boundary and execute forward and backward through IDA's native C++/CUDA training engine.

An earlier evaluation against IDA's structured training targets had produced loss near or above the random-vocabulary regime and initially looked like an architecture bug.

It wasn't.

The general pretrained model was simply being scored against a data domain it had not been trained to interpret.

Plain English separated architecture correctness from dataset-domain mismatch.

That distinction mattered as much as the final number.

---

Two Different Kinds of Portability

Ember and SmolLM2 answer different questions.

Ember showed that numerical and systems ideas developed around IDA could be tested independently on a foreign model, on foreign hardware, in a foreign trainer — and still return useful evidence upstream.

SmolLM2 showed that a foreign pretrained architecture could actually travel into IDA's native runtime.

Those are not the same result.

Together, they matter more than either one alone.

IDA-TRAIN-V2 started as infrastructure for training the IDA family.

That remains its primary purpose.

But a runtime that can now:

• run on Hopper and consumer Blackwell
• execute block-scaled low-precision arithmetic
• adapt when a framework refuses hardware that exists
• split a billion-parameter training workload across two consumer GPUs
• fall back when direct GPU peer access is unavailable
• implement standard RoPE semantics
• translate external Hugging Face checkpoints
• execute real forward and backward on a 1.7B pretrained model

has crossed an architectural boundary.

It is becoming a specialized native training runtime with an expanding model and hardware surface.

---

The Failures Are Part of the Architecture

The Blackwell work also reinforced something central to the way IDA is being built.

A mature training engine does not merely accumulate successful kernels.

It accumulates failed explanations.

Stochastic rounding mattered because it ruled out one explanation.

Hadamard rotation mattered because it ruled out another.

Lion mattered because it weakened the optimizer hypothesis.

The dead NVFP4 head forced the investigation deeper into the gradient path.

The gradient-scale cliff corrected our interpretation of FP4 itself.

The 5090 compilation failure exposed a hidden Hopper dependency.

The failed P2P assumption created a more portable communication path.

The misleading SmolLM2 dataset result forced us to isolate architecture correctness from domain mismatch.

Even the kernel ontology was corrected by the external work: a launch can be geometrically SM_STARVED or SM_UNDERSUBSCRIBED without the training step as a whole being underutilized.

Launch geometry is not duty cycle.

Every one of those failures sharpened the system.

That is why IDA keeps telemetry, fallback reasons, promotion gates, post-mortems, and negative results beside the code.

If a failure is recorded well enough, the project should only have to pay for it once.

---

Where the Evidence Stands

At the time of writing:

MXFP8

Three paired real-corpus seeds in Ember's trainer.

Approximately 1.17× BF16 median throughput, with validation-loss differences smaller than seed-to-seed variation.

MXFP4

Executes directly through consumer Blackwell tensor-core instructions on RTX 5070 Ti hardware even where the tested PyTorch path refuses to expose it.

NVFP4

Executes on consumer Blackwell and produced the fastest Ember training arm tested at approximately 1.288× BF16.

Its early training experiments also exposed a severe gradient-scale underflow mechanism that could zero entire low-magnitude blocks.

That failure was corrected in Ember's trainer with a global power-of-two gradient prescale, and the mechanism subsequently informed IDA's Blackwell preparation.

Consumer Blackwell native execution

IDA's C++/CUDA engine now has a real non-Hopper compile and execution path rather than an assumed future port.

Two RTX 5090s

A controlled 1F1B ablation produced approximately 27.9% higher throughput in the tested memory regime.

RoPE

Forward, backward, and packed-sequence semantics have been directly tested inside the native engine.

SmolLM2-1.7B

A standard Hugging Face checkpoint has been translated into IDA's native format and executed through real forward and backward paths, producing 3.85 nats / ~47 perplexity on a plain-English isolation test.

---

What Is Still Open

There are clear boundaries on the current result.

GQA is not yet supported.

The corrected NVFP4 training matrix is still being characterized.

Full-production-microbatch 1F1B remains a memory problem.

The internal Blackwell fragment-layout question requires a stronger discriminating probe.

One successful external architecture does not imply arbitrary Hugging Face compatibility.

And a short forward/backward SmolLM2 validation is not a claim of full convergence under IDA.

Those are not caveats to bury.

They define the next experiments.

---

Where the Blackwell Trainer Is Heading

The work is converging toward a fairly simple systems philosophy:

Let the hardware vendor own the tensor-core primitive.

Let IDA own the training path around it.

That means owning:

• packing
• scaling
• layout
• precision placement
• memory lifetime
• scheduling
• communication policy
• fallback behavior
• instrumentation
• lineage
• promotion criteria

A general framework has to support almost every imaginable model and workflow.

IDA does not.

Its specialization is only valuable if it turns that narrower scope into better measurement and better decisions.

That also means precision cannot simply be a global model setting.

The Blackwell experiments have already shown different behavior between internal projections, heads, forward operands, and backward gradients.

The relevant question is no longer:

Should this model use FP8 or FP4?

It is:

Which numerical representation belongs in which computational role, at which point in training, and under what evidence?

That is a harder question.

It is also the question a specialized training runtime is actually positioned to answer.

---

Leaving Hopper

We started this work asking whether a Hopper-native engine could survive consumer Blackwell.

The RTX 5070 Ti gave us the first serious evidence — indirectly, through an independent model and trainer.

Its low-precision tensor cores worked differently.

Some framework boundaries turned out not to be hardware boundaries.

Packing became as important as arithmetic.

The fastest format produced the worst model.

Our first explanation for that failure was wrong.

And following the failure deeper exposed a gradient-range problem that stochastic rounding, Hadamard rotation, Lion, and error feedback could not fix because none of them changed the magnitude regime causing the block scale itself to collapse.

That finding came from Ember.

It then traveled back into IDA.

The RTX 5090s pushed the question beyond arithmetic.

Could the entire IDA native engine compile on consumer Blackwell?

Could two consumer cards cooperate without the datacenter assumptions we had built around?

Could standard positional semantics be added without breaking backward?

Could a Hugging Face checkpoint actually cross into the engine?

The answer to all four became yes — with caveats, failures, fixes, and measurements attached.

That is the part of the result I care about most.

IDA is not becoming useful because every experiment works.

It is becoming useful because the system is getting better at determining why an experiment worked or failed, preserving the answer, and making the next machine inherit it.

A low-precision format does not fail merely because it is coarse.

A GPU is not incapable merely because a framework refuses an operation.

A multi-GPU design is not portable merely because both machines have two GPUs.

And a training engine is not architecture-independent merely because its tensor shapes look familiar.

The boundaries have to be tested.

Sometimes IDA tests them itself.

Sometimes an external model like Ember exposes them first.

What matters is that the evidence can travel.

That is what consumer Blackwell has been doing to IDA.

And so far, the engine has kept moving.

---

Contributor Acknowledgment

The RTX 5070 Ti consumer-Blackwell research, low-precision training experiments, profiling, hardware-probe execution, and upstream technical findings described in this article were contributed by Tolerable (@Tolerable), known as Rev within the project, working with his AI engineering collaborator Oslo.

Much of that work was carried out in the independent trainer for Ember, Rev's small, from-scratch language model.

Ember did not run on the IDA engine.

Instead, she provided an independent real-corpus Blackwell testbed for the MXFP8, MXFP4, NVFP4, precision-placement, packing, and gradient-range experiments described here.

The resulting measurements and failure mechanisms materially informed IDA-TRAIN-V2's Blackwell preparation, instrumentation, numerical safeguards, defect correction, and subsequent experimental roadmap.

The two-RTX-5090 native-engine port, 1F1B experiments, RoPE integration, and Hugging Face checkpoint sideload work were carried out as part of the continuing IDA-TRAIN-V2 development effort.