collapse in torch is bit-exact with the hc_collapse kernel and a faster graph (512 rows: 0.15 vs 1.73 ms
standalone). post_torch (the kernel's operation order) is bit-equal standalone but rounds differently fused
in the layer graph, so it is opt-in. Tiny golden 9/9 with collapse. bind_cpu receives the physical core id;
_neuron_bind_cpus indexed CUDA_VISIBLE_DEVICES with it again (broke runs on cores other than 0..n).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- fused_sample: _host_flags/_greedy_verify flags persist while the batch's requests are unchanged
(each device read drained the queue and waited for the whole forward). Tiny golden 9/9; full model
step time unchanged (660 ms), driver dflash launch 71 -> 10.5 ms host.
- bind_cpu on Neuron: the device's NUMA node (DEEPINFER_NEURON_CPU_BIND=node|core|legacy); the generic
split gave each rank one CPU and ranks 0-15 the wrong node. Neutral on TPOT (228.2 ms).
- profiled followers exit through sys.exit so the runtime writes their trace.
- op_census: trace_syncs / census_sites (DEEPINFER_NEURON_TRACE_SYNCS=1|2|3).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- examples/neuron: tiny random V4.1 checkpoint in the real HF layout,
pure-torch replacements for the official reference kernels, golden
generation, a teacher-forced DeepInfer runner (LLM offline API) and a
comparison tool. CPU mode: 25/27 greedy tokens match, worst logits
cos 0.9985 / rel 0.056 against a reference self-noise floor of 0.032-0.060.
- model_loader: check parameters against the platform device, not is_cuda.
- deep_gemm paged MQA logits: do not dereference block-table slots past the
context.
- bind_cpu: skip in Neuron CPU mode (1 core per NeuronCore share would
serialize host compute).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
