Triton Backed Supported (Beta) - #313
Open
student-Jungmin wants to merge 81 commits into
Open
Conversation
triton_src_dir parsed HEXAGON_MLIR_ROOT out of tnpu's setup/versions.env. That key is gone: tnpu moved its triton and triton_shared checkouts out of a shared /workspace/hexagon-mlir parent and into siblings, and the variable became TRITON_ROOT. The parse silently missed, fell through to the hardcoded /workspace/hexagon-mlir default, and returned a path that no longer exists, so ensure_triton_importable could not borrow tnpu's checkout at all. The two keys are not interchangeable. HEXAGON_MLIR_ROOT named the PARENT, so a "triton" segment had to be appended; TRITON_ROOT names the checkout itself and must not have one. Appending to the new key would have produced a path just as wrong as the old default, which is why this is not a rename. Both keys are read, newest first, because the two repos are versioned independently and a tnpu from before the rename still publishes only the old one. The fallback moves to /workspace/triton-src/python to match tnpu's current default. Verified against a migrated workspace: triton_src_dir returns /workspace/triton-src/python, the directory exists, ensure_triton_importable returns True and triton 3.6.0 imports in the driver interpreter. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
strip_for_tnpu raised SpecIncomplete on any libdevice.* call, on the grounds
that those members are extern stubs with no triton_shared implementation. That
was true when it was written and is not true now: triton_shared binds them in
get_module_map as of df7452a, and triton-npu's own libdevice kernels lower
through it -- ops_math_derived, ops_math_special, ops_math_bits, activations
and lrn all pass end to end, and the one that does not, ops_math_elementwise,
is blocked on a register spill its docstring already records, not on libdevice.
So the missing piece was never lowering. It was the same one tl_math has three
lines above: Inductor reaches libdevice through torch, strip_for_tnpu drops the
torch import, and nothing rebinds the name. Bind it to triton's own copy and
the calls emit tt.extern_elementwise, which the converters downstream already
key on.
Unblocks 2 of the 5 tests the coverage report bucketed as libdevice, and the
other 3 stop being libdevice problems:
test_transcendental PASS tanh, exp, erf, sin, cos
test_exponent PASS
test_pointwise now IndexError inside the tnpu pipeline
test_layernorm now IndexError inside the tnpu pipeline
test_floormod_axis_split
now triton_helpers.maximum, the known helper gap
The allowlist takes the two that pass. The report's libdevice bucket is stale
as of this commit; the other three want a fresh sweep to be rebucketed rather
than a hand-edited count.
Verified: the two new entries pass through the route, and test_add and
test_triton_codegen still pass, so the strip path did not regress.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Two blockers found by running ResNet-18 on the Triton route. Both were guards or omissions on our side of the seam, not lowering gaps. estimate_peak. ExtensionWrapperCodegen.generate overrides upstream's generate and calls memory_plan_reuse directly, but upstream sets self.estimate_peak in run_wrapper_ir_passes -- the method the override replaces -- so nothing set it. AllocateLine.should_reuse_buffer reaches it only when a reuse candidate is more than one scheduler node back, which is why add never noticed and ResNet died with AttributeError instead of a planning miss. Now set the same way upstream does, under the same config.allow_buffer_reuse guard. triton_helpers. strip_for_tnpu refused every kernel naming one, because the module lives in torch and the tnpu venv has none. But maximum, minimum, min2, max2 and the two they call are pure triton -- @triton.jit over tl.* and nothing else -- so a copy runs there unchanged. They are emitted into the spec instead of imported: there is no torch to import from, and a module dropped beside the spec would depend on how tnpu's stage-1 worker sets sys.path. The block must be self-sufficient because it is prepended, ahead of the body's own imports. The NaN handling is copied deliberately. `mask |= a != a` makes a NaN operand win, which is what torch.maximum promises and tl.maximum does not do, so this is not simplifiable to the builtin. A helper outside that set still raises, and the message now says what makes one vendorable rather than asking for the whole module. ResNet gets past both and stops at the third blocker, a multi-axis grid (ynumel=64, YBLOCK=None), which is the known block-size policy gap. Verified: test_activation passes end to end (ReLU, Sigmoid, SiLU, SwiGLU) and joins the allowlist. test_add, test_triton_codegen, test_prologue_fusion, test_batchnorm and test_transcendental all still pass. test_softmax, test_matmul_activation and test_mlp clear the helper guard and now fail deeper in the tnpu pipeline -- test_softmax on 'vector.load' op most minor memref dim must have unit stride -- so they are rebucketed, not fixed. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The workflow section said how to open a PR and nothing about when work is finished, so finished work sat uncommitted until someone asked for it. It is finished when it is pushed; that is now the rule, and the rule is the permission so there is nothing to ask about. Naming the remote is the load-bearing half. Reading "push" as "push to origin" puts our work on PSAL-POSTECH/PyTorchSim, the upstream everyone shares, when it belongs on the student-Jungmin fork. That is the one mistake here that cannot be undone quietly, so the fork is written down rather than left to be remembered. The rest is what this session learned the hard way. Clearing outputs/ between runs, because a cached artifact replays and made a real fix look like it changed nothing -- which produced one wrong conclusion before the cache was noticed. Pinning TNPU_DIR to a known-good worktree, because stages 1-5 live in a repo someone else may be mid-refactor in, and a pass in that state fails in ways that read as ours: "no lane axis", NameErrors that vanish on re-run. And saying plainly when a guard was cleared so a test now fails deeper, since a rebucketed failure counted as a pass is how coverage numbers stop meaning anything. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Three things, all reached by running ResNet-18 down to its third blocker.
THE AXES WERE ALWAYS EMPTY. fixed_config_for passed kernel.numels, keyed by bare
prefix, to parallel_axes, which looks for "ynumel" and found "y". So axes came
back [] for every kernel on this route. Single-axis ones survived by accident on
the cfg.setdefault("XBLOCK", lanes) at the end; multi-axis ones lost YBLOCK and
surfaced much later as "YBLOCK=None" out of grid_of. The multi-axis warning
right above never fired either, which is the tell nobody had one working. Both
callers now read the dict collect_meta already built.
THE GRID WAS BACKWARDS. grid_of orders axes outermost-first (z, y, x) to match
Inductor, while tnpu's spec.grid is positional and X-first -- wrapper._grid3
reads g[0], g[1], g[2] as gx, gy, gz. For a 1D kernel the two are the same
tuple, so every kernel on this route agreed until a second axis appeared. Handed
the outermost-first tuple, tnpu read gridY as gridX, the inner axis never
enumerated, and a transpose came back with its first XBLOCK columns correct and
the rest still zero. Wrong output, no crash. grid_xyz now builds it by axis name.
The timing path was already correct -- it maps axes through _PID_SLOT -- so only
the functional spec was affected.
BLOCKS COME FROM THE HARDWARE CONFIG. Tile dim 0 gets vpu_num_lanes: the MVIN
DMA scatters that axis across the lanes and it is also the systolic array's side
(sa_dim = n_vu), so one value has to serve both, which is why matmul needs this
read rather than assumed. The innermost axis gets vpu_vector_length_bits divided
by the widest element, so each lane holds exactly one vector register -- 8 fp32,
16 fp16, 32 int8 -- instead of the 1 it was pinned to, which made every
work-item move a strided column.
Which letter is dim 0 depends on rank, so it is indexed, not named. Confirmed in
04-custom.mlir: the 2D transpose stages memref<128x8xf32, 1> with
vlane_split_axis = 0 under YBLOCK=128 XBLOCK=8, and 1D add stages
memref<128xf32, 1>, same axis, under XBLOCK=128.
Verified: all 14 allowlist tests pass, so nothing regressed. A standalone
transpose is exact on 64x156, 128x128 and 128x8 -- mismatched 0 of every element
-- where before the kernel could not be described at all.
MULTI-AXIS IS NOT DONE, and the allowlist is deliberately unchanged.
test_transpose2D, test_transpose3D, test_conv_fusion and test_conv_view_input no
longer raise SpecIncomplete; they compile, run, and come back numerically wrong.
That is a rebucketing, not a fix. The next bug is the fused case: for
transpose(0,1) + b the addend is applied correctly and the transposed load
returns the wrong element, so it is indexing, not the grid.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Three port-layer gaps, none of them needing a compiler change. DTYPES THE TABLE HAD STOPPED SHORT OF. _DTYPE knew fp32/fp16/bf16, i64/i32/i8 and i1. tnpu grew int16, the unsigned types and float64 in cf6c683 -- with its own coverage kernel -- and this table was never told, so a kernel touching int16 or uint8 lost the dtype for its output buffer and died as "no dtype/numel for out_ptr0", naming a metadata hole rather than the missing entry. The two tables now agree. test_widen_dtype passes. triton_helpers.any IS PURE TRITON. It is tl.reduce over `a | b`, so it belongs in _HELPERS_SRC with the others and never needed refusing. test_gqa passes, the first attention test through this route. R0_BLOCK WAS LOOKED UP UNDER THE WRONG KEY. The reduction prefix carries a trailing underscore -- Inductor names the parameter r0_numel and the block R0_BLOCK -- so "r0numel" missed every time and the block stayed None on kernels whose extent was known all along. That is the same shape of mistake as the kernel.numels one: a key spelled from memory instead of read. reduction_axes now derives the prefixes from the keys, and _block_name already produced the right name from them. The value is not free to choose. The scratchpad is lane-banked with no lane-crossing primitive, so the reduced axis has to sit inside ONE lane, which makes covering the whole extent the only executable layout -- there is no accumulator across grid steps to make a partial block work. So R0_BLOCK is the extent rounded up to a power of two, and stays None when it would not fit vpu_spad_size_kb_per_lane, which fails loudly instead of picking a layout the hardware cannot run. Verified: all 14 allowlist tests pass. test_widen_dtype and test_gqa join it, taking the route to 16. WHAT DID NOT GET FIXED, said plainly. R0_BLOCK now lets five kernels past the guard and every one of them fails inside tnpu instead -- bmm_reduction, matmul_reduction, attention_fusion, transformer_fusion and test_transformer, four of those aborting. test_softmax gets further still and returns wrong values. That is rebucketing to the compiler side, not progress on those tests, and tnpu's own kernels/reduce.py already records why: linalg.reduce is left alone and falls through to convert-linalg-to-loops. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
welford, welford_combine and welford_reduce are pure triton -- tl.where, tl.zeros_like, tl.full and tl.reduce, nothing reaching into torch -- so they meet the same bar as the helpers already in _HELPERS_SRC and were being refused for the module they live in rather than for what they do. That is the numerically stable variance every LayerNorm and BatchNorm kernel Inductor emits goes through, so the guard was closing a common path on a wrong test. THIS BUYS NO NEW PASSES, and the allowlist is unchanged. All three affected tests move to a later failure and none of them reach a correct answer: test_layernorm runs, returns wrong values test_vit rejected by tnpu's select_lane_axis test_convnextv2 rejected by tnpu's select_lane_axis Worth landing anyway, because where they now stop is the real bottleneck rather than a guard of ours: the vit kernel is a single fused add_cat_convolution_expand_native_layer_norm_permute_view, and select_lane_axis declines to pick a lane axis for it. That pass is where the answer has to come from, and these two are ready to re-measure the day it changes. Verified: all 16 allowlist tests pass. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Running tests/models/DeepSeek/test_deepseek_v3_base.py against a current triton-npu turned up four, each hiding the next. STAGE NUMBERS ARE NOT AN INTERFACE. tnpu renumbers its stages whenever one is added or split, and the post-vcix IR moved from 04-custom.mlir to 05-custom.mlir. Three call sites hardcoded the number, so a pipeline that had already SUCCEEDED reported FileNotFoundError from the launcher. stage_artifact() matches on the suffix, which is the stable half of the name, and leaves the number to tnpu. R0_BLOCK IS PER LANE, AND IT WAS SIZED AS IF IT WERE NOT. The tile is [XBLOCK, R0_BLOCK] with XBLOCK on the lanes, so R0_BLOCK is what ONE lane has to hold, multiplied by however many buffers are live. Sizing it to cover the whole reduction -- 1536 rounded to 2048 -- put eight 4KB tiles in every lane and blew the double-buffer budget at link time, with an error that talks about the scratchpad and never mentions block sizes. It never had to cover the extent: Inductor already emits `for r0_offset in range(0, r0_numel, R0_BLOCK)` and carries the partial sums itself, so the block is a tile, not the axis. The reduced axis still lives inside a lane, which is the constraint that matters -- a chunk of it is as in-lane as all of it. The budget for it also came from the wrong place. extension_config says 128 KB per lane; tnpu enforces 64 KB and launches spike with --scratchpad-size=65536. Sizing against the YAML overshoots by exactly 2x. tnpu is the one that refuses, so its number is the one to size against, and TNPU_SPAD_SIZE overrides both. The two configs disagreeing is a separate bug and is not fixed here. dep_analysis REFUSED TWO OPS IT UNDERSTANDS. Loop terminators -- scf.yield and friends -- forward a memref to their parent and read nothing; the reads are the loads inside the body, classified on their own. Every chunked reduction ends its body that way, so shrinking R0_BLOCK walked straight into it. And togsim.transfer carries its direction in dma_kind, not in operand order: operands[0] is the DRAM side and operands[2] the SRAM side both ways round, so classifying by position calls an MVOUT an MVIN and gets every dependency backwards. Verified by advance, not by a green suite: the test now clears all four and fails further on, in wrapper_codegen, which another session is already working in. A full allowlist run was not made because the working tree currently holds that session's uncommitted changes and would have measured a mixture; these five files were committed by name for the same reason. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…nder a template call inductor_templates.py exists on origin/feature/triton-helpers (01190f3) and not here: this repository is a shallow clone and the branch that carries the route today has a separate 17-commit history, so the file was simply absent rather than removed. Without it use_triton_template asks is_gpu, npu is neither gpu nor cpu, and every mm/bmm/addmm/conv goes to extern_kernels -- where conv raises convolution_overrideable not implemented and stops resnet18 at its first layer. Taken at the branch tip, so it brings _short_circuit_degenerate_gemms and the pick_config that ranks extern last with it, plus the device_guard fix from the same commit ("pass" is not something the caller can write `with` in front of). The one thing that is new here: a template kernel does not go through TritonNPUKernel.call_kernel, so the sympy rendering that lived there covered only pointwise kernels and a template call reached ", ".join as an Integer. Moved to TritonNPUWrapperCodegen.wrap_kernel_call, which every kernel call in this route passes through, and the pointwise copy deleted. Measured on resnet18: 20 extern convolutions and one extern addmm become triton_convolution2d_* and triton_mm_*, and the whole graph reaches codegen. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Inductor's mm and conv templates bound the last tile with min(M - pid * BLOCK_M, BLOCK_M), so arith.minsi reaches the trace producer. convert-arith-to-emitc marks minsi/maxsi/minui/maxui illegal and offers no pattern, and arith-expand only covers the ops that need extra arithmetic, so nothing lowers them and the pipeline fails to legalize: error: failed to legalize operation 'arith.minsi' that was explicitly marked illegal The kernel itself was fine -- tnpu lowered it and Spike wrote the right values -- so this was the timing path alone, and it presented as a working kernel whose trace producer would not build. min(a, b) is select(a < b, a, b) and both of those do convert; a MAX differs only in the predicate, so the true arm is `a` either way. Collected before rewriting, because walk_ops recurses into an op's regions after yielding it and erasing during the walk invalidates the handle it is about to ask. Measured: tests/ops/fusion/test_addmm_residual.py passes, with TOGSim simulating the template kernel (3470 cycles) instead of the run stopping at lower_to_emitc. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
get_config_heuristics has no registry lookup -- it is an if/elif over cuda, xpu, cpu and mtia and then BaseConfigHeuristic() -- so npu took the generic set, whose first conv entry is ConvConfig(64, 256, 16, 2, 4). With 128 lanes that is a [64, 256] tile banked on the N axis at two columns per lane, and bank_vectorize refuses it: operand [64, 256] banked on axis 1 with strides [1, 64] is not reached by indexing its own row-major type in the nest [2, 64] this op's maps ask for resnet18 reaches it at its eleventh conv, the first whose out_chan is 256; below that preprocess_mm_configs clamps BLOCK_N to the channel count and the tile fits by accident. N is the lane axis so it takes the lane count exactly, the same choice fixed_config_for makes for XBLOCK. M and K are per-lane depth and cost scratchpad rather than lanes, so they are offered small first. THE REFUSAL IS STILL A DEFECT. A tile deeper than one element per lane is a shape this backend has to handle; this only stops resnet from being what reports it. The heuristic registered above has carried the same TODO since it was written -- a block size is a statement about the machine, and a table written for a GPU is not one. Measured on resnet18: 10 kernels compiled before, 21 now, all 15 convolution kernels among them. The run gets past codegen entirely and stops in Spike on a scratchpad address, which is the next thing. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
A pointwise kernel's grid is ceil(numel / BLOCK) because that is how it walks its output. A template kernel (mm, conv) does not walk its output that way: it tiles with its own BLOCK_M/BLOCK_N and reads tl.program_id(0..2) directly, so the grid is the template's to state and Inductor passes it as three extra launcher arguments (FixedGrid). The numels such a kernel reports describe the output tensor, not its iteration space. The route applied the pointwise formula to both. For ResNet's first conv that produced 6272 programs where the template asks for 196 -- 32x -- and the surplus programs indexed tiles past the end of every operand, which Spike reported as "User load segfault @ 0xd0f66270". Worse and quieter: every template whose second axis is real had gridY pinned to 1, so tl.program_id(1) never moved and only the first BLOCK_N columns were computed. Two of ResNet's convs are in that shape ((1,4,1) and (4,2,1)). collect_meta now records the grid by asking grid_fn the same way Inductor asks it for its own benchmark harness, and refuses with SpecIncomplete if it cannot -- inventing a number is what sent this trail 32x wrong. launch_axes/launch_extents pair the axes with the extents in one place, and all four consumers read it: the spec's grid, grid_of, the WorkItem's program-id arguments, and the trace producer's shape file. write_shape additionally stops reading a template's trailing call arguments as numels; for a template those arguments ARE the grid. launch_axes declares all three slots for a template even when an extent is 1. The source still says tl.program_id(1) -- conv 0 multiplies pidY by BLOCK_N -- and an undeclared axis is a kernel argument nothing accounts for, which stops the trace producer with "kernel arg still used after build_skeleton". A declared axis of extent 1 runs its loop once. Measured: conv 0 spike FAIL (359.3s, segfault) -> ok (206.8s), full 802816-element output written. Allowlist 16/16, no regressions. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
A template kernel (mm, conv) was being launched with the pointwise grid,
xnumel / XBLOCK. That number has nothing to do with what the template asks
for. Its source is a jinja template with BLOCK_M/N/K baked in as literals
and its grid is over output tiles, which torch states directly in
select_algorithm.TritonTemplateKernel.call_kernel:
call_args.append(f"*{fn_name}({...call_sizes...}, {meta})")
Measured on the DeepSeek run, triton_npu_fused_mm_t_view_5 with BLOCK_M=32,
BLOCK_N=32 was launched at grid 144 = 18432 / 128, the pointwise formula,
where the template wanted cdiv(M, 32) * cdiv(N, 32). Five template kernels
in that graph all had pointwise grids. Wrong program count over tiles the
kernel never asked for, and no shape disagrees loudly enough to notice.
So collect_meta now records the template's own grid and grid_xyz prefers
it. Verified on tests/ops/gemm/test_matmul.py: every emitted mm matches
cdiv(M, BLOCK_M) * cdiv(N, BLOCK_N), including the non-multiple case
M=128 N=56 which rounds up to 4 * 2 = 8 where the pointwise formula gave
56.
An ordinary kernel has numels and no grid_fn, so it returns None and the
existing path is untouched. A template whose grid function will not
evaluate raises SpecIncomplete rather than falling back to the pointwise
formula, because falling back is the bug being fixed.
A kernel indexes with the tensor's STRIDES, not with its shape. Inductor reads those strides at compile time and writes them into the source as constants -- resnet18's first conv carries stride_xc = 1, because the graph put its input in channels-last -- so the bytes handed to Spike have to be the storage in address order. `.contiguous()` produces logical order. For a channels-last tensor that is a different permutation of the same values, so the kernel read every element from the wrong place while computing perfectly correctly on what it found there. read_outputs had the same defect mirrored: `view_as(t)` reads the flat file as logical order and scatters it home wrongly. HOW IT WAS PINNED. The kernel's output matched a torch reference taken over the file read channels-last to 1e-5, and differed from the real answer in 602115 of 802816 elements at 1.38288. PyTorchSim's own per-kernel check reported 602108 and the same 1.38288 -- the same divergence seen from the other side, which is what says the kernel is right and the marshalling is not. Measured: resnet18's first divergent buffer moves from buf1, the first kernel of the graph, to buf32, and from 602108/802816 elements to 3823/25088. Triton route allowlist 16/16. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
… on changed
A graph is re-run to see the NEXT layer, not the ones already settled, and Spike
is the whole cost -- resnet18 spends minutes there per conv and the first twenty
are unchanged between runs. So a launch whose inputs are the ones it had last
time replays that answer instead of simulating it.
The key is everything the answer depends on, which is why this is safe to leave
on by default:
- the ELF's own bytes, so a fix anywhere in tnpu misses. The workdir is keyed
by the TRITON SOURCE alone and would not catch that.
- the bytes of every input, so a different tensor misses. That also makes the
chain self-repairing: change an early layer and every later one misses,
because its input is a different tensor.
Nothing else reaches the kernel, and the Triton source is already in the workdir
path.
Not keyed by layer number on purpose. Cutting at "layer N onward" replays a
stale answer without noticing when something before N moved.
The cache sits BESIDE the workdirs, in outputs/.triton_replay, not inside one.
Picking up a tnpu-side fix means deleting outputs/triton_*, which is this repo's
own instruction; a cache kept inside would go with it exactly when it was most
wanted. The ELF is in the key, so it does not need the deletion.
A replay is logged as one ("[Spike] <kernel> replayed <key>") rather than as a
run, because "it passed" means something different when nothing was simulated.
TORCHSIM_TRITON_REPLAY=0 turns it off.
Verified on tests/ops/conv/test_conv2d.py, cold then warm: 1 replay then 7, same
"Conv2d Forward Test Passed" both times.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
A full run is the thing being trusted, and a result that came out of a file is not a result the simulator produced today. The key argues it would have been the same -- the ELF's bytes and every input's bytes are in it -- but an argument is not a measurement, and a report that says a model passed should mean the model was simulated. So TORCHSIM_TRITON_REPLAY=1 asks for it and nothing else turns it on. The inner loop is where it earns its place: the same graph re-run to reach the kernel actually being worked on, where the first twenty layers are settled and only the next one is in question. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
`triton_route_sweep.py --all` already ran tests/models/test_resnet.py, in the
"Full sweep (report)" step, which is continue-on-error and gates nothing. The
allowlist is the step that gates, and the model was not in it.
UNIT KERNELS PIN THE PIECES; ONLY THE GRAPH PINS THEM MEETING. triton-npu
carries seven coverage kernels reduced from resnet18 -- its sixteenth and
seventeenth convolutions, the global average pool, and the first conv's bounds
mask -- and every one of them passed on both sides of the merge that broke the
model. What broke was two rank-1 operands landing on opposite iteration dims of
one elementwise op, which needs both halves present in one graph and cannot be
reduced to a kernel that would have caught it in advance.
measured the allowlist step with this line, on the merged backend:
17/17 ok, test_resnet.py at 587.0s. It runs beside the other
jobs (needs: build-app), so the wall clock of the workflow is
unchanged.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…iming be switched off
Three changes to the Triton route, all about work the backend is asked to do and
does not need to.
1. THE GRID. Inductor's conv grid asks for the whole channel count on the axis
its kernel indexes per group:
cdiv(n * h * w, BLOCK_M), cdiv(c, BLOCK_N), GROUPS
while the template masks `idx_y_c` against GROUP_OUT_C = OUT_C // GROUPS.
Axis 1 therefore runs cdiv(OUT_C, BLOCK_N) blocks where only
cdiv(GROUP_OUT_C, BLOCK_N) have a live column; the rest are launched, DMA
their tiles, mask everything away and write nothing.
2. THE BLOCK SIZE. The clamp that narrows BLOCK_N to the channel extent already
exists -- that is why a 32-channel convolution gets BLOCK_N 32 -- and it is
handed out_chan, every group at once. So a depthwise layer, whose every group
has ONE output channel, was configured with BLOCK_N = 128 and masked 127 of
them away on every program. `groups` is known only inside the lowering, so
that frame records it and the heuristic reads it.
3. THE TIMING HALF. `__call__` says "the two halves are independent" and only
the functional one was switchable: a run checking VALUES paid for a cycle
simulation of every kernel it touched. It reads `pytorchsim_timing_mode` now,
the same switch extension_codecache.py already reads on the MLIR route.
measured mobilenet_v2's depthwise convolutions, before -> after:
GROUPS=960 (1, 8, 960) BLOCK_N=128 -> (1, 1, 960) BLOCK_N=16
GROUPS=576 (4, 5, 576) BLOCK_N=128 -> (4, 1, 576) BLOCK_N=16
GROUPS=384 (4, 3, 384) BLOCK_N=128 -> (4, 1, 384) BLOCK_N=16
GROUPS=144 (49, 2, 144) BLOCK_N=128 -> (49, 1, 144) BLOCK_N=16
Its ten grouped kernels launch 28544 programs where they launched
about 78000, and GROUPS=1 is untouched -- there GROUP_OUT_C IS
OUT_C and the two expressions are the same number.
measured e2e resnet18 through the Triton route with all three in:
"resnet18 inference Test Passed", max diff 2.86e-06, 21 kernels.
That is the control -- it is all GROUPS=1, so it is the run that
says these changes do not disturb an ordinary convolution.
WHAT IS NOT VERIFIED YET, and it is the reason this says so rather than claiming
more: e2e mobilenet_v2, the model the grouped path is FOR, was still running its
value check when this was committed (22 of 57 kernels launched, no errors). The
grids and block sizes above are measured; its output is not yet compared.
AND THIS REMOVES WASTE RATHER THAN CREATING WORK. A depthwise convolution still
launches one program per group -- 28544 of them here -- because the template puts
`group` on a grid axis and indexes one group's weights per program. Folding
several groups onto the N axis is what would use the array, and that needs a
template of our own rather than a fix to this one.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…reads The launcher's own docstring says "the two halves are independent" and only the functional one was: the timing half ran whatever the config said, so a graph being checked for VALUES paid for a cycle simulation of every kernel it touched. An e2e model is where that is the whole cost. `pytorchsim_timing_mode` is the switch the MLIR route already reads (extension_codecache.py), so this route reads the same one rather than inventing a second name. Validation runs go with it off; the cycles are a separate claim and want their own run. Not measured here beyond the switch taking effect -- test_transformer.py is still stopping in tnpu, so there is no full-model number to put against it yet. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
DeepSeek-V3's MoE router picks its experts with top-k, and the kernel Inductor
emits for it calls triton_helpers.sort_with_index. That name lives in torch, the
tnpu venv has no torch, so the spec writer refused it:
SpecIncomplete: kernel uses triton_helpers.{sort_with_index}, which lives in
torch and the tnpu venv has no torch.
It meets the test this block already states for what may be copied here -- pure
triton, @triton.jit over tl.* and nothing else. The closure is three functions,
sort_with_index -> _bitonic_merge_with_index -> _compare_and_swap_with_index,
and the two names they reach outside it are triton's own, not torch's: _log2 is
triton.language.standard._log2, which torch itself imports from there through
triton_compat, and is_floating is the copy already in this block.
THE NaN HANDLING IS THE ALGORITHM, not noise to tidy. Comparisons with NaN are
always false, and sort has to treat NaN as the larger value, so the left_isnan /
right_isnan terms are what make it agree with torch.sort -- the same reason the
note above _tnpu_maximum gives for its `mask |= a != a`. Copied verbatim.
VERIFIED, AND SAY WHAT THAT MEANS. The block imports and jits under the
torch-free venv at /workspace/mlir-env, with sort_with_index bound to a
JITFunction and torch absent. On the DeepSeek run with a cleared cache,
triton_npu_fused_copy__sort_44 no longer raises SpecIncomplete and now reaches
the tnpu pipeline, where it fails further on:
'memref.expand_shape' op expected expanded type to be 'memref<128x2x16xi8>'
but found 'memref<128x2x16xi1>'
That is a triton-npu gap on the i1 mask this sort reshapes, not this one. The
kernel is not passing; it is failing one stage later, which is the claim being
made here and the whole of it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The sweep that gates this route compares every test's output against a torch
reference and reads no cycle count anywhere, but it ran with the default config,
whose pytorchsim_timing_mode is on. So a gate about VALUES was also paying for a
cycle simulation of every kernel each test compiled.
That is what kept mobilenet_v2 out of the allowlist. It is the model the grouped
convolution path is for, its values pass, and it does not finish inside the
1800s timeout with timing on -- its depthwise layers launch one program per
group, 28544 of them across the model.
measured mobilenet_v2 through the route, timing off: about 21 minutes,
"MobileNet Output Test Passed". With timing on it exceeded the
1800s timeout. resnet18 takes 587s with timing on.
The new config is the mirror of _timing_only, which switches the functional half
off for the same reason from the other side; it differs from the default by that
one line. The sweep uses it as a DEFAULT, so a caller who wants cycles sets
TOGSIM_CONFIG and this leaves it alone.
WHAT WAS NOT RUN, and it is the reason this says so: the gate sweep itself. It
was started twice and both runs are void -- another session was linking
triton-shared-opt at the time, so every kernel died at stage 1 with
PermissionError on that binary, which is neither this change nor the backend.
What IS measured is above: mobilenet's runtime with timing off, and the previous
sweep in which the other 17 allowlisted tests passed and only mobilenet timed
out. Re-run scripts/ci/triton_route_sweep.py once the toolchain is quiet.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ers it
Inductor decides persistent-vs-looped from the reduction's extent against a
threshold -- 64, or 1024 when the hint is INNER -- and this backend then pins the
block itself, in kernel_spec.fixed_config_for, to "cover the extent and no more,
then shrink to the budget". The two disagree in the middle:
r0_numel = 96 > 64 Inductor: looped
R0_BLOCK = 128 >= 96 ours: the loop body runs ONCE
A looped reduction is not merely redundant there. Its partial results have to be
COMBINED across blocks, and for a variance that combination is Welford's: a
tl.reduce over a (mean, m2, weight) triple with a six-argument body.
triton_shared's ReduceConverter takes a body of exactly one op from a fixed list
-- "Only support lowering reduction with body containing 1 max(i/f), addf, ori,
or mulf" -- so the triple survives as tt.reduce into a pass with no lowering for
it. Before p10_bufferize learned to refuse a Triton op by dialect, that was a
segfault inside one-shot-bufferize's inliner.
With persistent chosen, Inductor emits the two-pass form itself and says why:
"For persistent reductions, don't bother with welford's algorithm since it uses
more registers, and taking two reductions doesn't increase memory usage." That
is a sum, a mean, and a sum of squared deviations -- each one addf, each one
convertible.
AND THE BLOCK CANNOT DISAGREE, which is what would make this dangerous. For a
persistent reduction Inductor writes the block into the kernel itself
(R0_BLOCK: tl.constexpr = _get_persistent_RBLOCK(numel), codegen_static_numels),
so fixed_config is not consulted for that axis and the value is the next power of
two above the extent -- covering it by construction. Where it does not fit a
lane's scratchpad this returns False and the decision is exactly what it was.
measured e2e convnextv2: 5 kernels compiled -> 16. It does NOT pass; it
now stops in p09 on "transfer's masked_axes cannot take the move",
a deliberate refusal about a merged run being half-masked. That is
a rebucketed failure, not a fix, and it is a different problem.
measured the gating allowlist, 16 of 18 run at the time of this commit:
15 ok, and the one FAIL is the check turned around below.
tests/ops/reduce/test_batchnorm.py is the one that exercises this
directly and it passes. NOT RUN: test_resnet and test_mobilenet,
which were still going -- resnet18 and mobilenet_v2 both passed
e2e earlier today on the parent commit, and mobilenet calls
welford in 0 of its 57 kernels, so this change cannot reach it.
tests/system/test_triton_codegen.py asserted the OPPOSITE -- that a reduction
must be refused, "tnpu has no lane-aware reduction ... When the lane path lands,
this is the test to delete". It landed sideways: nothing in tnpu changed, but a
reduction that finishes inside one tile never crosses a lane. Rather than delete
the check, it now asks what it was really guarding -- whether the numbers are
real. t.sum(dim=1) over [128, 64] comes back at 1.91e-06.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Splitting a reduction buys parallelism this backend cannot collect yet, and pays
for it in a form it cannot compile.
The option says what it buys: "For reductions with a small output size (usually
1, e.g. x.sum()) there is not enough parallelism to saturate the GPU ...
split_reductions: uses multiple kernels to gain more parallelism". The extra
kernel is a win only where the two halves RUN AT ONCE, and today they do not --
tnpu compiles one binary per kernel and the C wrapper walks the grid as a
sequential loop. That is the CURRENT LAUNCHER, not a property of the machine: the
hardware has cores and the config has num_cores, so when launches run
concurrently this trade changes sign and this line wants revisiting rather than
assuming.
What it costs meanwhile: stage one leaves partial (mean, m2, weight) triples and
stage two combines them, which for a variance is Welford's -- tl.reduce over a
triple with a six-argument body. Inductor has a two-pass fallback for
welford_reduce and NONE for welford_combine, and triton_shared's ReduceConverter
takes a body of exactly one op from a fixed list, so the triple survives as
tt.reduce into a pass with no lowering for it. The split is not merely unpaid
for; it produces a kernel that does not build.
THE ORDER TO LIFT THIS IN, when concurrent launches land: give welford_combine a
lowering (or a two-pass fallback of its own) FIRST, then turn the split back on
and measure. Doing it the other way round only brings the failure back.
measured e2e convnextv2's convolution_native_layer_norm_permute_17:
r0_numel 2, so persistent by any threshold, and still a
(3 x tensor<128x2xf32>) -> (3 x tensor<128xf32>) Welford because
the reduction was split in two. With the split off, convnextv2
goes 17 kernels compiled to 35. It does NOT pass -- the next stop
is a Welford from a reduction that is genuinely looped (r0_numel
768 against a 512 block), where the triple is the right answer and
the missing piece is its lowering.
measured the gating allowlist, 17 of 18: 16 ok and one FAIL that reproduces
clean on its own (exit 0, every check ok) -- that run overlapped
another session linking triton-shared-opt, which took every test
down at stage 1. NOT RUN: test_mobilenet, stopped for time. It
passed e2e earlier today and calls welford in 0 of its 57 kernels.
The second change is unrelated in mechanism and found by this run: the warning
this route logs when timing is off said "[TOGSim]", and the sweep buckets a
failure by matching output against `TOGSim|trace\.so|SIGSEGV|...`. Every test
that failed in this mode was therefore filed under togsim whatever went wrong --
test_triton_codegen came back "[togsim]" for a stage-1 toolchain problem. Tagged
"[timing]" instead.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The whole encoder block passes: twelve kernels compiled by tnpu, run on Spike,
and every one of its 84 intermediate tensors inside rtol = atol = 1e-4 of the
CPU graph -- checked one buffer at a time by
`pytorchsim_functional_verify_per_kernel`, not just at the output.
Encoder Block Forwrad Test Passed
WHAT IT TOOK, all of it upstream or in tnpu and all pinned by kernels there:
a clamp that could never fire blocking a unit-axis collapse (layernorm); a bool
store dragging every access in its function onto the unstructured path, whose
address step was zero for i1; and two wrap defects in triton_shared's
PtrAnalysis -- a pointer advance folded into a dimension that wraps, and a wrap
swapped onto an axis it cannot be expressed on (`fix/wrap-advance-fold` and
`fix/wrap-axis-with-layout`, both on PSAL-POSTECH).
measured compile 12 kernels in 10s; Spike is the whole cost after that
(softmax_6 36s, each projection ~20s, the two bmms 15s and 14s).
Cycles are not in this: timing mode stays off for a values run.
THE LAST TWO MEASUREMENTS OF THIS MODEL WERE VOID and it is worth knowing why.
`outputs/triton_<hash>/` is keyed on the INDUCTOR SOURCE, so a fix below it --
in triton-shared or in a tnpu pass -- does not change the hash, and the launcher
reuses the ELF it already has. clear_codegen_cache.sh does not remove those
dirs. Two runs reported the same divergence at buf9 while the kernel, given the
model's own recorded inputs, passed standalone at 2.7e-07; the artifacts were
dated the previous day. Delete `outputs/triton_*` when the compiler moves.
This needs the fixed triton-shared: until the pin moves, TNPU_TRITON_SHARED_OPT
has to point at a build of it.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… are missed
`clear_codegen_cache.sh` removed `.torchinductor` and the MLIR route's
per-source-hash dirs, and left `outputs/triton_<hash>/` standing. That hash is
of the INDUCTOR SOURCE, so a fix anywhere BELOW it -- a tnpu pass, triton-shared
-- does not change it, and the launcher reuses the ELF it already has. The
script is exactly what someone runs after such a fix, and it was the one thing
that did not clear.
measured two runs of test_transformer.py reported the same divergence at
buf9 while that kernel, given the model's own recorded inputs,
passed standalone at 2.7e-07. The artifacts were dated the
previous day. A compiler that moved and a cache that did not is a
wrong measurement, not a slow one.
Matched by prefix rather than by length so it does not depend on the hash width,
and the existing 11-char rule is untouched. Verified: a triton_* dir and an
11-char dir go, an unrelated dir under outputs/ stays.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…oices hook
A NON-PERSISTENT REDUCTION IS WHAT MAKES INDUCTOR EMIT WELFORD. Its Triton
backend takes the fallback -- sum, mean, sum of dx^2 -- only on the persistent
path ("don't bother with welford's algorithm since it uses more registers"), and
the two forms are the same numbers. What differs is what reaches this backend.
The Welford form does not reach it at all. Its combiner folds THREE tensors with
a six-argument body, which triton-shared does not convert, so it arrives as a
`tt.reduce` over three tensors -- an unregistered op holding tensors, which
bufferize's inliner walks and dies in C++ on (triton-npu 3cd5f33 turned that
crash into a refusal by dialect). There is nothing to lower it TO either: the
reduce this machine has is one accumulator folded by a known combiner inside a
lane's vector.
So this is a statement about the target rather than a tuning knob, which is why
it goes in the choices handler. Inductor's own rule is 1024 elements for
ReductionHint.INNER and 64 for anything else, and a LayerNorm over 768 lands on
either side depending on what it was fused with -- INNER for test_transformer's,
OUTER for test_vit's, where the same normalisation sits with a patch convolution
and a permute. 2048 covers both; past it the tile stops fitting a lane's
scratchpad and the refusal moves to fit_to_hardware, which names it.
`V.set_choices_handler` with an `InductorChoices` subclass is the documented way
to say this -- no private symbol, no monkeypatch (rule 16) -- and every other
decision defers to the base class.
measured test_vit: welford gone from every kernel it generates, and the
model advances from a bufferize refusal to a vector-width one in
fit_to_hardware, three kernels further in.
test_transformer: still passes, all 84 tensors, and generates no
welford either.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Three changes, one of them a deletion, all from taking test_vit past the wall it
had been stopped at.
A BUFFER IS AS LONG AS ITS STRIDES REACH. The kernel addresses with the
buffer's strides -- Inductor bakes them into the source as constants -- so a
layout whose strides leave gaps reaches past its element count:
buf18 = empty_strided((1, 12, 197, 197), (465792, 38816, 197, 1))
197*197 = 38809 elements per head at a pitch of 38816. Sized by the shape, the
last head runs 77 elements past the buffer the harness allocated and those 77
are lost on the way back, while every head after the first is read at the wrong
offset. It surfaced as ViT's first functional-verify failure, 416959 of 465708
elements of a softmax over tolerance -- and the kernel had computed that softmax
correctly, 0.0318 max error over the padded reading with all of it in the
truncated tail. kernel_spec._buffer_numel now returns the last addressable
element, and functional's round trip says the same thing in one statement:
`flat.as_strided(t.shape, t.stride())`, which subsumes the permute-to-memory-
order it used to do (a channels-last tensor lands the same way) and also covers
the case a permutation cannot express. buf18 verifies; the divergence moves on
to buf20.
R0_BLOCK IS CORRECTED, NOT GUESSED. fixed_config_for sizes it against a budget
divided by _REDUCTION_LIVE_TILES, a constant standing in for how many
block-sized tiles the kernel keeps live -- a property of the LOWERING, which
does not exist until the block size has been chosen. ViT's first LayerNorm,
fused with a patch convolution, an addmm and a transpose, keeps 41: 77504
bytes/lane at 512, 38680 at 256. No constant serves both that kernel and an
ordinary reduction, so codecache now compiles, reads tnpu's machine-readable
refusal (tnpu-spad-overflow: usage= budget=, triton-npu dab7026), divides the
block by usage/budget rounded up to a power of two, and recompiles. Measured on
that kernel: one retry, four seconds. It is read off TnpuError.output rather
than str(exc) on purpose -- the message keeps only lines that look like a
diagnostic, and this one is addressed to a function.
reduction_choices IS DELETED. It existed because a non-persistent reduction
makes Inductor emit Welford, whose three-tensor combiner had no lowering on this
target. It has one now (triton_shared 46d70d7), so the file was a workaround for
something that is fixed, and forcing persistence was costing the tile size the
scratchpad then failed on.
Verified: test_transformer.py passes with all three in and reduction_choices
gone; test_vit.py compiles all 21 kernels (it stopped at the fourth) and reaches
buf20.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…it passes
`.to("cpu")` does not carry arbitrary strides across -- it materialises a
contiguous copy -- so a padded tensor arrives packed, and `cpu.stride()` then
reproduces the very defect the previous commit was written to fix. Measured on
ViT's buf18: the file kernel 12 was handed diverged from the file kernel 10 had
written at flat index 38809, which is 197*197, the head size WITHOUT the
seven-element hole. Head 0 was exact and every head after it was shifted by
seven, which is why it read as a bmm returning plausible-but-wrong numbers
(first bad index [1, 0, 0], 138316 of 151296 elements) rather than as garbage.
The copy into a strided view is by logical index, so the contiguous cpu tensor
lands in the padded slots correctly; only the LAYOUT has to come from the tensor
the kernel was compiled against.
tests/models/test_vit.py now passes end to end on the Triton route -- 20
kernels, 17 buffers verified against their CPU goldens, final output 1.07e-06 --
and is added to the allowlist. test_transformer.py and test_resnet.py re-run
with the fix and pass (2.38e-06).
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…s one
Two block sizes here were justified by a layout claim that measurement does
not support. kernel_spec.fixed_config_for said tile dim 0 is the axis
bank_vectorize spreads over the lanes, so the outermost block takes the lane
count for a per-lane depth of 1, "the shape every tnpu baseline runs". The
conv heuristic in inductor_templates said "N IS THE LANE AXIS" and cited the
same reason.
Nothing in this repo picks the lane axis. tnpu's select_lane_axis does,
several passes into the other repo, from what the ops demand: a matmul wants
its last axis on the lanes, a reduction wants the axis it folds off them,
elementwise operands have to agree, conflicts resolve by flipping the
matmul, and dim 0 is only the default when nothing asked. Counted over the
tnpu dumps, 26 of 36 stamped kernels are axis 0, gemm_fp16_kernel is axis 1
throughout, and gemm and bmm carry both on different operands of one op.
The per-lane depth of 1 was not required either, and that is why the first
claim looked right: pinning the outermost block to the lane count made dim 0
the degenerate answer often enough that neither half was ever tested. Both
are tested now, by two kernels added to triton-npu on develop-select-grid:
tile_deeper_than_one_per_lane block 256 on 128 lanes, axis 0, exact
tile_gemm_lane_axis_deeper BN 256 on the axis a matmul demanded,
9.54e-06 against a 1.53e-05 control
The second kernel's 04-adapted.mlir carries lane_axis 1 three times plus
array<i64: 1, 1, 1>, and nine spad buffers of memref<128x256xf32, 1> where
the control has thirteen of memref<128x128xf32, 1>, so the deeper tile is
built and addressed rather than quietly reshaped.
THE NUMBERS DO NOT CHANGE. Losing a justification is not a measurement that
some other size is better, and a block size change needs its own
verification -- raising the mm tile to the machine size earlier today broke
tests/ops/fusion/test_matmul_vector.py and was reverted. What changes is
that the lane count is now written down as a default width rather than as
the only legal answer, so the next person can weigh changing it.
Also measured while here: of 13 kernels generated across four allowlist
tests, none has a ynumel and none has an r0_numel. The multi-axis branches
in fixed_config_for are barely exercised.
Verified: scripts/ci/triton_route_passing.txt plus test_matmul, 15 of 17,
the same two failures as the baseline (test_gqa, test_prologue_fusion, both
bmm template kernels dying in lower_to_emitc._rewrite_signature, both
confirmed failing without any change of mine).
…lpers
CLIP PASSES ALREADY and nobody had run it. 21 kernels, 1.35e-05 -- its vision
tower is the same shapes ViT's is, so everything that landed for ViT carried it
without a line of its own. test_single_perceptron passes too (4 kernels, weight
and bias update). Both are in the allowlist now, which is the only thing that
makes "passes" survive the next change.
div_floor_integer AND remainder_integer are vendored into _HELPERS_SRC. They are
pure triton -- tl.where over `//` and `%` -- so they meet the bar that block
already states. `a // b` in triton truncates toward zero while torch's rounds
toward minus infinity, and the difference is the whole reason the helpers exist:
SwinV2's window partition indexes with floor division, and a negative operand
appears there the moment a cyclic shift is applied. remainder_integer is its
pair and comes along because the two are one convention, not because a kernel
asked yet.
Measured, on the models nobody had gated:
test_clip 21 kernels PASSES 1.35e-05
test_single_perceptron 4 kernels PASSES
test_swinv2 22 -> 26 kernels, now a CppCompileError
test_convnextv2 51 kernels, Spike segfault in kernel 40
(a load outside the scratchpad)
test_mlp 2 kernels, buf2 diverges (addmm, 512/512)
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The per-kernel check went in after the first kernel to name a buffer, on the
stated assumption that the producer precedes every consumer in topo order. That
holds for a buffer one kernel fills. It does not hold when ONE fx op is split
across several kernels, and then the check compares a half-built buffer against
a finished golden and reports a divergence that is not one.
DeepSeek-V3's MoE router is that shape. aten.scatter.value comes out as two
kernels sharing one origin node:
triton_npu_fused_scatter_zeros_like_38(buf9, 256)
_fverify.verify_check(buf9, "buf9", "scatter", "aten.scatter.value")
triton_npu_fused_scatter_zeros_like_39(buf8, buf9, 128)
38 writes the zeros, 39 scatters the ones, and the check sat between them. It
reported
first divergence at buffer 'buf9' (op aten.scatter.value)
max abs diff 1, 128/256 elements over tol
(0, 2): npu=0 cpu=1 (0, 3): npu=0 cpu=1 ...
every scattered one "missing" -- because it had not been written yet. Kernel 39
run standalone against a torch reference of the same store gives max_abs_err 0.
So the lines are walked twice: once to record where each buffer is last written,
once to emit. Same buffers, same one check each; only the position moves.
With this, DeepSeek-V3 passes end to end with the checks on: 55 verify_check
calls emitted, zero divergences, and the graph-level comparison passes too.
Before it, the run stopped at buf9 on a buffer that was correct.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The cap's justification was "wide tiles come back wrong", which is not true: with the wrap gone, 512x512x512 passes on the maximal (512, 512, 512) tile and a bare 32x768x1536 matmul at BLOCK_N = 512 comes back at 6.96e-05. What is left is one thing. Running each suspect in test_gqa alone at BLOCK_N = 512, M = 32, K = 768, N = 1536: matmul(a, b) 6.96e-05 matmul(a, bt.t()) 5.53e-05 addmm(bias, a, b) 5.84 addmm(bias, a, bt.t()) 5.84 An addmm epilogue reads its bias as a full [BLOCK_M, BLOCK_N] tile with a zero DRAM row stride -- tl.load(in_ptr0 + tl.broadcast_to(idx_n, [BLOCK_M, BLOCK_N])) -- and that replicating load breaks as soon as the lane axis carries more than one element: 128 passes, 256 and 512 do not. The rank-1 spelling of the same epilogue passes at 512, so it is the replication rather than the width. Reproducer: triton-npu kernels/coverage/tile/ tile_bias_row_deeper_than_one_per_lane.py (45f4542). Comment only; the cap and the code are unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
tests/models/GPT2/test_gpt2.py joins scripts/ci/triton_route_passing.txt, so GPT-2 now gates the Triton route alongside resnet18 and mobilenet_v2. The default preset drops from small to tiny, because triton_route_sweep.py runs each test with NO ARGUMENTS -- the default is the gate. Measured with exactly that invocation: GPT-2 lm (tiny) passes, max diff 4.7684e-07, 14 kernels, and with pytorchsim_functional_verify_per_kernel on there are zero divergent buffers at rtol=atol=1e-4, so every intermediate matches CPU and not just the logits. --preset small (the real 768-wide block) does NOT pass and the comment at the default says so: kernel 0 stops because a gathered MVIN deeper than one element per lane replicates its first element instead of fetching the second. That is pinned by triton-npu kernels/coverage/gather/gather_masked_deep_in_loop.py (82871d0). Raise this default to small the day that kernel goes green. GPT-2 needs triton-npu develop-e2e-gpt2 ceedbd7 or later: 0cdcf38 for the causal mask kernel and the develop-select-grid merge for the attention bmm. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… are gone The cap's justification was "wide tiles come back wrong". Three separate things hid behind that sentence. Two are now fixed and the cap is over the third. FIXED, the wrap: on the same 512x512x512 kernel in triton-npu, with only the rm % M changed, 110.93 off becomes 5.34e-05 at the same (256, 256, 512) tile and the same 2x2 grid, and the array instructions come out as the same 2 x 4 x 2 loop nest either way. clamp_instead_of_wrap removes it. FIXED, the bias: an addmm epilogue loads its bias through a broadcast pointer tensor, and triton_shared's PtrAnalysis rewrote every op on the way to it except the broadcast, so the load found no descriptor at [BLOCK_M, BLOCK_N] and became a gather -- which then wrote each lane's one fetched value into both of the slots that lane owns. Fixed upstream by PtrAnalysis::rewriteBroadcastOp (triton_shared 9f4b16f), pinned by triton-npu tile_bias_row_deeper_than_one_per_lane.py (8fc3f6f). Isolated by running each suspect alone at BLOCK_N = 512, M = 32, K = 768, N = 1536: matmul(a, b) 6.96e-05 matmul(a, bt.t()) 5.53e-05 addmm(bias, a, b) 5.84 addmm(bias, a, bt.t()) 5.84 With both in, test_gqa passes with the cap lifted. NOT FIXED and not yet named: with the cap off, tests/ops/fusion/test_prologue_fusion.py is 127.39 off. Its kernel is a bmm at BLOCK_M = BLOCK_N = 512, BLOCK_K = 64, one program per batch, and a bare 512x512x64 matmul at the SAME tile passes standalone on triton-npu at 7.63e-06 -- so it is the bmm form or the fusion, not the width, and saying which needs a measurement nobody has taken. So the cap stays, over one case instead of three. Allowlist 16/16. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…t cap clamp_instead_of_wrap looked for `tl.load(A + ` and gave the whole rewrite back when it found none. That is the mm template's spelling; the bmm template accumulates the address into the pointer and writes `tl.load(A)` or `tl.load(A, mask=..., other=0.)`. So every bmm kept its wrap, kept its gather, and said nothing -- the log printed a line per mm and none per bmm, which was visible and got walked past. With the loads matched on `tl.load(<ptr>` instead, tests/ops/fusion/test_prologue_fusion.py goes from 127.39 off to passing with the tile cap lifted. Its kernel is a bmm at BLOCK_M = BLOCK_N = 512, BLOCK_K = 64, and a bare 512x512x64 matmul at the SAME tile already passed standalone on triton-npu -- so this was never a third defect, it was the first one not being reached. Two smaller things in the same function: A dividing block needs no mask, so the anchor and the load patterns are now looked for only when some axis asks for one. A template this does not recognise still gets its dead wrap removed. The existing mask is spliced by index, not by regex. `mask=([^,)]+)` stops at the first comma or paren, and `mask=rk[None, :] < k` has a comma inside a subscript, so widening it produced mask=(rk[None) & _tnpu_row_mask, :] < k which parses and means nothing. The end of the expression is the ` other=` the templates always follow it with, or the load's own closing paren. AND THE CAP GOES. It was over two things, both now fixed -- this, and the addmm bias, whose broadcast pointer tensor became a gather until triton_shared's PtrAnalysis::rewriteBroadcastOp (9f4b16f). The docstring keeps both measurements and points at the two reproducers in triton-npu. Allowlist 16/16 with no cap. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
67 kernels, the whole model, and the graph-level comparison passes. Measured
under the sweep's own conditions -- its functional-only config, one job -- at
948.9s, which is 53% of the 1800s per-test timeout and in the same range as
mobilenet_v2, already on the list.
It also passes with the per-kernel golden check on: 55 buffers verified against
CPU, zero divergences, twice from a cleared cache. That is a stronger statement
than the graph comparison alone, which can be right while an intermediate is
wrong.
WHAT IT NEEDS, and neither is in this repository:
triton-npu develop at 5c580e1 -- eleven defects this model found, six of
them in the lane-axis pass alone
triton_shared torchsim at c3e1019 -- three widths PtrAnalysis was losing.
The prebuilt triton-shared-opt predates that merge, so a
checkout using the pinned binary will still stop in the MoE
gate until it is rebuilt.
Listed rather than left passing quietly: a test nobody runs is a test that goes
red without anyone noticing, and this one crossed eight refusals to get here.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…dropped Every tile gemm_combination_mapping returns is a divisor of the padded extent times the lane count, so a power-of-two extent gives power-of-two tiles. Asked about the real shape it returns blocks Triton cannot take -- tl.arange needs a power of two -- and the whole shape then falls back to torch's generic table. Measured, mapped tiles kept after the filter: 129x61x56 2 -> 0 100x100x100 1 -> 0 384x384x384 8 -> 1 32x32x96 1 -> 0 which is test_gqa's bmm Rounding each extent up to the next power of two first: 129x61x56 2 tiles, best (256, 64, 64) 100x100x100 1 tile, (128, 128, 128) 384x384x384 21 tiles, best (256, 256, 512) 32x32x96 1 tile, (32, 32, 128) and nothing is dropped by the filter at all. ROUNDING UP IS SAFE BECAUSE THE TAILS ARE MASKED, which was not true before this branch. M and N are bounded by kernel_spec.clamp_instead_of_wrap -- that is what it is for -- and K by the template's own EVEN_K masks. The budget is computed on the rounded shape, so it over-reserves rather than under, and the power-of-two filter stays as the check that this held rather than as policy. Allowlist 16/16. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
… stopgap The comment said the budget divisor was a policy over an unknown and that the real fix was the re-codegen loop the MLIR route has and this route does not. Building that loop was the next item; it is not needed, and the measurement says so. Trying deliberately to overflow: a 512-cubed matmul with FIVE fused elementwise epilogue nodes comes back at 1.98e-04, and a 1024-cubed one with TEN at 1.71e-03. Both link. Halving the budget makes the mapping pick a smaller tile, and a smaller tile makes the epilogue's own tiles smaller with it, so the reservation scales with what it is reserving for. A retry path with no case that needs it is machinery nobody can check -- rule 2 wants a kernel that fails without the fix and there is none -- so this stays a measured reservation. If a kernel ever does overflow, that is the reproducer and the loop can be built against it. Comment only. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…ed to
The file was tests/models/Mixtral8x7B/test_attention.py and the test inside it
calls itself "Mistral" -- `test_result("Mistral", res, cpu_res)` -- and the MLIR
route's CI job that runs it is already named "Run test_mistral.py". Only the
path disagreed, so the rename brings it in line with all three; the workflow's
path is updated with it.
It passes on the Triton route: 43 kernels over a 32-token prompt and three
decoded tokens, all three verified. It used to stop at kernel 6. What it needed
was the bool-mask chain -- triton_shared af3aa55 walks the pointer bitcast
triton puts in front of every torch.bool buffer and moves the BASE with it,
triton-npu 2855c27 erases the one lower_tts_to_transfer consumed, and
develop-e2e-gpt2's 0cdcf38 wired the gather's `other` slot before either.
Not re-verified for this commit: the rename and the allowlist line are the whole
of it, and the passing run is the one recorded above.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…0) works
t.sum(dim=0) stopped at stage 4 while t.sum(dim=1) on the same tensor
passed, so the axis looked like the difference. It was not. Inductor
hands both down correctly:
sum(dim=0) xnumel 64, r0_numel 128
sum(dim=1) xnumel 128, r0_numel 64
and fixed_config_for gave BOTH XBLOCK = 128, the lane count, without
looking at the numel. For dim 1 that is exact. For dim 0 it is a tile
twice the iteration space, and the surplus is carried by a mask whose
rank-1 index rows are stranded ONE_LANE while the data is banked across
the lanes:
NotVectorisable: one operand of this elementwise op is in a single
bank (ONE_LANE) while another is banked across the lanes on iteration
dim 0 of extent 128
ins(%3, %3, %2 : memref<128xi32, 1>, memref<128xi32, 1>,
memref<128x128xf32, 1>)
memref<128x128> for a [128, 64] tensor is the whole story in one type.
Blocks are now clamped to the smallest power of two that covers the
numel -- rounded up because tl.arange needs one, so a numel of 100 still
gets 128 and still needs its mask.
sum(dim=0) refused -> 3.815e-06
sum(dim=1) 1.907e-06 -> 3.815e-06 (the r0 block moved with it)
NO COVERAGE KERNEL FOR THE BACKEND'S SIDE, and that is measured rather
than skipped. A hand-written tnpu kernel reducing over the banked axis
with a block deliberately wider than the extent -- [128, 64] data under
a [128, 128] block, so the mask genuinely fires -- COMPILES AND IS
EXACT. With the block equal to the extent it also passes, because a
trivially true mask folds away before stage 4 sees it. So the refusal is
narrower than its message and a kernel for it would pass either way,
which pins nothing.
test_triton_codegen now checks both axes; its docstring said dim 0 was a
gap and that is no longer true.
Allowlist 16/16.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
# Conflicts: # PyTorchSimFrontend/triton_backend/codecache.py # PyTorchSimFrontend/triton_backend/kernel_spec.py # PyTorchSimFrontend/triton_backend/timing.py # tests/system/test_triton_codegen.py
docs/worktrees.md:52 says ".envrc is local to each worktree and not committed", but it was not in .gitignore, so every worktree that follows that instruction shows a permanent untracked entry and `git status` stops being a useful signal there. Unrelated and deliberately left alone: the line above ends "experiments/artifact/logs/*# tutorial generated artifacts" with no newline, so that pattern cannot match anything. Fixing it would start ignoring files that are not ignored today, which is a behaviour change and not this commit's business. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
e6e9e0c set the default preset to tiny and said to raise it the day the wide-tile gather went green. It has: triton-npu develop-select-grid 2007619 fixed a gathered MVIN that filled only lane-count elements per lane and then repeated them, which took out this model's very first kernel at any width above the lane count. The default is the gate -- triton_route_sweep.py runs each test with no arguments -- so this is measured with exactly that invocation: preset=small n_layer=2 n_embd=768 n_head=12 vocab=1024 seq=32 26 kernels through Spike, max diff 2.0862e-06 zero divergent buffers over 219 goldens with pytorchsim_functional_verify_per_kernel on, rtol=atol=1e-4 so every intermediate matches CPU, not just the logits. That is two layers of the real GPT-2 block rather than a scaled-down stand-in. Needs triton-npu develop-e2e-gpt2 61c1b94 or later. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…e chosen The scratchpad retry's lever is the reduction block, and a persistent reduction takes it away. Inductor writes R0_BLOCK: tl.constexpr = <next power of two above the extent> INTO the generated source, so fixed_config_for's answer is never read and _shrink_reduction_blocks halves a number the kernel does not consult. Measured on BERT-small kernel 0 (three embedding gathers plus the first LayerNorm, r0_numel 768, 21 scratchpad globals). Inductor calls 768 persistent, being under the INNER threshold of 1024, and bakes R0_BLOCK 1024. The retry then recompiles EIGHT times and the measurement does not move by one byte: 70944 bytes/lane over a budget of 65536, retrying with R0_BLOCK 128 70944 ... 64 70944 ... 32, 16, 8, 4, 2, 1 then blocks is empty and it re-raises. With persistence declined the same kernel compiles on the FIRST try, no retry, at a looped R0_BLOCK of 256. The rule is one fact, not a heuristic: persist iff the persistent block IS the block we would have pinned. Equal and nothing is given up. Bigger and persisting trades our only correction for a tile we already know does not fit. kernel_spec.reduction_block_for is lifted out so both callers read one number rather than deriving it twice. Wired through torch._inductor.config.inductor_choices_class, the documented out-of-tree hook -- virtualized.py: "We virtualize InductorChoices to allow changing inductor heuristics from out of tree." An earlier version of inductor_templates forced persistence by REBINDING the same method and its removal comment called that out as the thing rule 16 forbids; this is the inverse decision on different wiring. The class is module level because Inductor pickles it into the FX cache key -- a local class raises "Can't pickle local object" and silently disables that cache. VERIFICATION, STATED AS IT STANDS. BERT tiny still passes at 4.77e-07, and BERT-small now reaches kernel 9 in the DEFAULT config rather than dying at kernel 0 under a probe flag. The Triton-route allowlist was swept with timing off: 19 ok and FIVE failures -- test_transformer, test_clip, test_vit, test_convnextv2, test_triton_codegen. Only the last is isolated: pre-existing, failing identically without this change and under the sweep's own config, because a timing-off run emits no trace.so for it to find. THE OTHER FOUR ARE NOT ISOLATED. They may be this change, the loop-fold guard in triton-npu 8997973, the develop-select-grid merge at 42b30a9, or pre-existing. Committed on a scratch branch at the user's direction with that question open.
… and DeepSeek The clamp added in 58d3fb7 takes a parallel block down to the smallest power of two covering its numel. For xnumel 64 that removed a mask that stranded rows and made sum(dim=0) work. For xnumel 1 it takes XBLOCK to 1, and a parallel block of 1 leaves no axis for the lanes -- the tile banks on the REDUCTION axis instead, which is the one arrangement stage 4 refuses, and it refuses it with the scalar the reduction produced: NotVectorisable: operand [1, 1] spans iteration dims [] of a tile [512, 1] banked on dim 0, and it was staged one element per lane -- every lane needs the whole row, which takes a replicating transfer this operand did not get Measured on Mistral's RMSNorm (xnumel 1, r0_numel 512): XBLOCK went 128 -> 1 and the tile went [512, 128] banked on x to [512, 1] banked on r0. It passed before the clamp and stopped after it. So did tests/models/DeepSeek/test_deepseek_v3_base.py, which reaches the same shape. The two directions have measured counterexamples and the boundary is exactly a block of 1: xnumel 64 (sum(dim=0)) 128 fails, 64 passes xnumel 1 (RMSNorm) 128 passes, 1 fails so the clamp now declines a numel of 1 and keeps the lane count there. Mistral passes, sum(dim=0) is 3.815e-06 and sum(dim=1) 1.907e-06. TWO WRONG READINGS ON THE WAY, both the same mistake. The sweep bucketed these as stage 3 (triton-shared) failures and I repeated that, because stage.log's last line is the stage 3 command -- the failure is in stage 4 and the log is truncated before it. Re-running the pipeline on the spec is what shows it. I had already been caught by this today on sum(dim=0) and was caught again, which is why the fix is to run the stage rather than to read the tail. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
# Conflicts: # scripts/ci/triton_route_passing.txt
…s env fork/develop-npu carries a TRACKED .envrc from the transformer line (870a00e). Merging it overwrote this worktree's untracked copy with one that names another session's absolute scratchpad paths, points TNPU_DIR at /workspace/wt-transformer, and does not set TORCHSIM_TRITON_CODEGEN. MEASURED CONSEQUENCE, and it is the kind rule 13a exists to prevent: the GPT-2 e2e run right after the merge went down the MLIR ROUTE against somebody else's triton-npu worktree and died with TypeError: MLIRScheduling.can_fuse_with_exceptions() takes 3 positional arguments but 4 were given which is a known MLIR-route breakage on torch 2.10, has nothing to do with this branch, and reads exactly like our regression. docs/worktrees.md:52 already says ".envrc is local to each worktree and not committed", and 1ba5e7f put it in .gitignore on that authority. A tracked copy contradicts both and cannot work for anyone but its author, since the paths inside it exist only in one session's scratchpad. Untracked here; the local file is restored to this worktree's own settings. This does not remove it from fork/develop-npu. Whoever owns that line should untrack it there too, or every worktree that merges gets its environment silently swapped. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…ratch # Conflicts: # .envrc
…once BERT passes end to end on the Triton route at --preset small: 2 layers, hidden 768, 12 heads, 15,014,400 parameters, 25 kernels, max abs diff 3.4720e-06 in 134s with timing off and per-kernel verify on. tiny is 4.7684e-07. Run again as CI will run it, with no arguments, byte-identical at 3.4720e-06. So the allowlist entry is the default preset, which is small. Gating on tiny would be writing down a pass this route does not have. Both walls that held it are gone and only one of them is this branch's work: the scratchpad retry now has a lever. 08ba6cd declines persistence at a block we would not have chosen, so kernel 0 loops and the retry MOVES the number -- 68448 bytes/lane over 65536, one recompile at R0_BLOCK 256, done. Before, the same kernel sat at 70944 through eight recompiles because a persistent R0_BLOCK is baked into the source. the cross-lane fold in kernel 9 is not reached any more. That is develop-npu's block-sizing and lane-cap work, not ours; triton-npu's guard for it is still in and does not fire. reduction_block_for now takes lane_bytes from the caller. The develop-npu merge moved fixed_config_for onto machine["spad_size"] while the helper still read TNPU_SPAD_SIZE, which is two sources for one number -- the exact seam the helper was lifted out to close. tnpu_bridge.machine() writes the same value into that variable, so the default is the same number rather than a second opinion. WHAT WAS NOT MEASURED, and it is a real gap: the rest of the Triton-route allowlist was not swept on this tree. Four of its tests -- test_clip, test_transformer, test_vit, test_convnextv2 -- were failing on the previous tree and were never isolated, and this merge rewrote the block sizing under all four. Their state here is unknown.
…re the block is not ours
Measured with per-kernel verify on, so every intermediate buffer is checked against a CPU golden rather than only the logits: preset layers n_embd vocab kernels max diff goldens time tiny 1 128 256 14 4.7684e-07 121 small 2 768 1024 26 2.9206e-06 219 medium 4 768 4096 50 4.1723e-06 415 235s full 12 768 50257 146 5.1260e-06 1199 717s full is GPT-2 base itself, 123,702,528 parameters. Zero buffers diverge in any of the four, and the error grows only gently with scale. vocab 50257 was the shape expected to find something -- the embedding gather and the 32x50257 lm_head projection are fifty times wider than small's -- and it found nothing. The five walls cleared today were all reached by widening a shape, so once they were covered, size stopped being a variable. The gate stays at small on purpose, and the docstring says why: full's 146 kernels are about 30 distinct ones repeated twelve times, so it would cost the sweep twelve minutes to cover what small already covers. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…t names it
The per-kernel verify checked every bare-identifier argument of every
kernel call, inputs included, while both halves of it said "writes":
_fverify_last_writer is named for the last kernel that writes a buffer
and recorded the last that mentions one, and _fverify_emit_checks says
"this kernel's output buffers" and iterated all of them.
The two part company under buffer REUSE. The wrapper renames storage
(buf20 = buf9; del buf9 # reuse), so one buffer's contents live under
another buffer's name, and that name's origin_node describes what
Inductor MEANT to put there. Check it after a kernel that only reads it
and the comparison is against a value nothing computed.
WHAT MADE IT REACHABLE. Inductor's inplace_buffers can name an
in_out_ptr the kernel never stores to. Measured on Stable Diffusion
v1.5's UNet: twelve kernels take an in_out_ptr0 and TWO of them --
..._native_group_norm_silu_unsqueeze_12 and ..._41 -- only load it, fold
it into a value and store that value elsewhere. mutated_arg_names still
lists in_out_ptr0. Those two are called eight times between them, on
buf20, buf88, buf104, buf170, buf189, buf208, buf230 and buf298, and
those eight were EXACTLY the eight divergences the model reported, each
against an add_N node no kernel materialises. Inductor's removed_buffers
does not name them either -- 94 entries, none of these.
SO ASK THE KERNEL, NOT THE TABLE. kernel_spec.stored_args reads the
argument names the generated source actually stores to (tl.store and the
atomics, which write too), demote_unwritten_inout turns an inout with no
store back into a plain in, and record_roles keeps the per-kernel roles
in call order for the wrapper. Position, not name: the same kernel is
called with different buffers, so the meta's own buffer fields name the
first call only. An unrecorded kernel keeps every argument checkable, so
this narrows only where there is a measurement to narrow it with.
Demoting the role also stops the runtime copying that buffer back as an
output, which was harmless only because the file still held what had
been written into it.
THIS IS A REBUCKETING, NOT A PASS. SD1.5 does not pass; its wall moved
and the checks behind it became real:
before first divergence buf20, 5 buffers verified OK
after first divergence buf351, 169 buffers verified OK
buf351 is the output of kernel 83, whose GATHERS are exact -- probed
with one input set to arange at a time so the result names the mean
index read: in_ptr0 off by 0, in_ptr2 by 9.8e-04 (float32 on 1023.5),
in_ptr3 by 2.4e-07, and a wrong element count would move those too. Its
error of 0.026 comes from its inputs, buf350, buf297 and buf301, none of
which is verified anywhere: they are outputs of conv/template kernels,
and that is the next hole, not this one.
Verified: tests/ops/fusion/test_addmm_residual.py, tests/models/
test_resnet.py and tests/models/test_transformer.py all pass, each with
outputs/ and the Inductor cache cleared first. The full allowlist was
not run.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
… route
Llama 3 had no test in this repo -- tests/models/Llama held only test_llama.py,
which is Llama 2 (hidden 4096, intermediate 11008, 32 query heads against 32 kv
heads). Nothing on any branch covered 3.
The new test carries the three things Llama 3 changes, each of which lands
somewhere different in the backend:
grouped-query attention 8 kv heads against 32 query heads, so repeat_kv
broadcasts on the head axis before the bmm
rope_theta 500000 a different inv_freq over the same access pattern
wider SwiGLU intermediate is 3.5x hidden rather than 2.7x
Sizes are scaled down from 8B so a layer fits a test, but the 4:1 head ratio is
8B's -- GQA is the one difference that is structure rather than a number, and
shrinking it away would leave the test covering nothing 2 did not already cover.
Both files now pass through the Triton codegen route and are added to the
allowlist so CI holds them:
test_llama.py DecoderLayer 8.34e-07 Model 3.93e-06
test_llama3.py DecoderLayer 9.54e-07 Model 4.47e-06
Llama 2 needed the wrap fix in triton_shared PtrAnalysis (grid offset parked on
a singleton dim escaped the modulo, so RoPE's cat([freqs, freqs]) came back as
angle 0 and the layer was off by 0.5288). Llama 3 passed first try on top of it;
GQA needed no backend change.
Verified with outputs/triton_* and outputs/.torchinductor cleared before each
run, TNPU_DIR and TNPU_TRITON_SHARED_OPT both pinned.
tests/models/Diffusion/test_diffusion.py, the UNet2DConditionModel at
reduced channel counts: 87 kernels compiled, 220 executed, per-kernel
functional verify clean, max diff 1.79e-06. Two consecutive runs with
outputs/ and the Inductor cache cleared first.
Three fixes had to land before it ran, one per repository, and the
allowlist note records which:
buf3 GroupNorm's Welford mean, 32 of 32 channels equal to the
mean over the whole tile. triton-npu 27fd9e6.
kernel 83 a segfault on the second trip round a reduction loop.
triton_shared 6a96e7a.
buf20 eight verify reports against nodes nothing materialises.
6c22011, this repository.
Runtime is about 11 minutes, inside the sweep's 1800s default. Not
re-verified for this commit -- it adds a line to the allowlist and
changes no code.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
# Conflicts: # scripts/ci/triton_route_passing.txt
105 files: CMake's whole build tree under PyTorchSimDevice/build/ (object files, CMakeCache.txt, compiler-probe binaries) and the four shared objects the extension links -- _C.cpython-311-x86_64-linux-gnu.so and torch_openreg/lib/*.so. They went in by accident with 870a00e, a CI commit about gating test_transformer, and nothing has needed them since: they are what "python -m pip install --no-build-isolation -e ." produces. CLAUDE.md already says they are build artifacts and not to be committed, and .gitignore covered TOGSim/build/ but not this one. The cost was not only size. A checkout of this branch aborts three times over on any machine that has built the extension locally -- "untracked working tree files would be overwritten" -- and the way through is to delete your own build and take the committed one, which is somebody else's compiler and somebody else's paths. Removed from the index only; the files stay on disk, so an already-built tree keeps working. .gitignore now names the three paths. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
YWHyuk
reviewed
Aug 14, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.