MODIFICATIONS -- every internalised file, and how many guards it carries
========================================================================

Generated from the tree, and kept honest by the property it asserts.  Apply
this transform to any file below and you get FlashInfer v0.6.15, byte for
byte -- steps 1 and 2 mechanically, step 3 from upstream and not from here:

    1. strip every `// PIE:` marker, its continuation comments, the
       `#ifndef __CUDACC_RTC__` under it and that directive's matching
       `#endif` -- keeping the guarded BODY, which is upstream's, and
       DISCARDING the `#else` branch, if any, which is ours;
    2. rewrite every `#include <pie/...>` back to the spelling upstream
       wrote -- the map that held those rewrites was in
       `tests/upstream_manifest.rs` and was empty; see THE TEST IS GONE
       below;
    3. put back everything a `// PIE: REMOVED` marker names, the files listed
       under HOST C++ below, and the EDITS listed under THE EDIT THAT IS NOT
       A REMOVAL, THE SECOND EDIT and THE THIRD EDIT below it.  A strip cannot
       perform step 3 on its own: removed text leaves nothing to strip, only a
       marker saying what stood there, and an ADDITION leaves the added code
       standing after its marker is gone.  Step 3 needs the upstream
       tarball; steps 1 and 2 do not.

THE DIRECTION OF STEP 3 REVERSED, and this document said the opposite until
it did.  It used to read "restore `#include "exception.h"` to `page.cuh` and
to `utils.cuh`, and restore the four deleted files ... it applies to two
files".  It now applies to FIFTEEN, because the guarded host C++ those two
sentences were built around is gone rather than hidden.  What follows is the
count, not a summary of it.

THE TREE HAS MOVED TWICE AND STEP 2 IS STILL THE IDENTITY.  This paragraph
used to read "no vendored file has moved out of `csrc/vendor/flashinfer/`
yet".  That directory no longer exists: `csrc/` was narrowed to device text
only, which left a provenance-named directory nowhere to be, and these 24
files went to `csrc/src/attn/flashinfer/`.  Vendoring became INTERNALISING.

They are now at `kernels/flashinfer/`.  `csrc/` was retired in turn -- there
is no C source in this crate to name, no build script and no archive, only
text NVRTC compiles -- and the tree it held is `kernels/`, where a file's path
below that directory IS the string an `#include` writes.  The second move also
took this tree out from under `attn/`, which it never belonged to: `comm/`
here holds the two all-reduce headers that are `kernels_cuda::comm`'s root.

Not one byte of upstream's changed in either move.  What changed is what a
directive here is answered by: the entries in `kernels-cuda`'s `src/source.rs`
that carried this tree under `attn/flashinfer/...` now carry it under
`flashinfer/...`, and the six `#include`s in `xqa/` that named `attn/xqa/...`
name `xqa/...`.  Both are PIE spellings already -- see the `// PIE:` markers
above them -- and neither is a byte of upstream's own text.

THE TEST IS GONE, AND NOTHING CHECKS THIS FILE ANY MORE.  Every paragraph
below that says `tests/upstream_manifest.rs` will catch something is now
saying it about a file that does not exist: `1a08b179a` deleted it, 2,064
lines, along with twelve other test files whose subject was the declaration
machinery that sweep removed.  It held the empty `<pie/...>` map step 2 refers
to and it re-derived all four columns of the table at the bottom of this file
and compared them to the rows.  That was the ONE check here that ran without a
network, and it is the check that caught `attention/mla.cuh` -- see ONE ROW WAS
ALREADY WRONG below, which is an argument for a test that no longer exists.

So the property this document asserts is now entirely manual, in both halves:
the columns against `wc -l` and `wc -c`, and the stripped text against a CPM
checkout of v0.6.15.  The columns were re-derived by hand for this revision and
all twenty-three rows equal their files.  Recording that is not the same as
enforcing it, and the difference is the whole reason the deleted test was
written.

THE EDIT THAT IS NOT A REMOVAL, AND IT CARRIED NO MARKER.  Every modification
this document had described until now was a deletion, a rename, a guard or an
`#include` spelling.  There is one that is none of those, it is in XQA rather
than here, and it went in after the strip without a `// PIE:` marker and
without any count being re-run:

    xqa/mha.cuh:1455    result.fill(AccType{});
    upstream mha.cu     result.fill(AccType{0, 0});

`2af552a24`, in `smemFp16ArraySum`.  `AccType` is
`mha::conditional_t<useFp32Acc, float2, InputElem2>`; for the bf16 lattice
`INPUT_ELEM2` is `__nv_bfloat162` (`xqa/defines.h:41`), which
`shim/cuda_bf16.h:262` aliases to the prelude's `device::bf16x2` --
`struct { bf16 x; bf16 y; }` whose element type declares its constructors
`explicit`, so a braced `{0, 0}` copy-initialising each member cannot see one.
NVIDIA's `__nv_bfloat162` takes it; ours refuses it.  `AccType{}`
value-initialises and is zero for the `float2` arm too.  The cause is OUR shim,
not upstream, which is what makes this a modification to record rather than a
bug to report upstream.

It costs 4 bytes and no lines, and it is the reason the XQA byte total in
`NOTICE` read 265,143 against a tree of 265,139: a line count cannot see it, a
byte count can, and neither was re-run.  It is step 3 work on a rebase -- put
`AccType{0, 0}` back and the file is upstream's, and it will not compile here
until it is changed again.

IT CARRIES NO MARKER, AND THAT IS A DECISION RATHER THAN AN OMISSION.  A
marker is a line of its own -- step 1 of the recipe above strips whole marker
lines, so a trailing one on the code line could not be removed without taking
the code with it -- and `mha.cuh`'s header note pins the arithmetic that makes
its citations readable: TWO insertions, both above the cited lines, so upstream
`:409` is `:435` here and `:2757` is `:2783`.  A third insertion at 1454 shifts
every citation below it and silently falsifies both of those worked examples.
The cost of the marker is larger than the cost of the record.

So the record is here, and `mha.cuh`'s own header says so at the top of the
file where a rebaser will read it before touching anything.  What a strip must
not do is produce a file that claims to be upstream's and is not; naming the
line in two places it will be read is what prevents that.

THE SECOND EDIT IS UPSTREAM'S BUG RATHER THAN OUR SHIM'S, AND IT IS THE ONLY
ONE OF THESE THAT A ROOT COULD NOT COMPILE WITHOUT:

    attention/mla.cuh:762    o_smem.stmatrix_m8n8x4(o_smem_offset_w, o_frag_f16);
    upstream mla.cuh:762     o_smem.template stmatrix_m8n8x4(o_smem_offset_w, o_frag_f16);

`smem_t::stmatrix_m8n8x4` (`permuted_smem.cuh:148`) is not a template.  The
`template` disambiguator asserts that the name which follows IS one, so the
compiler demands the argument list that a non-template cannot take, and NVRTC
says *"argument list for template ... is missing"*.  Upstream's own TEN other
call sites -- `mla.cuh:531` and nine in `prefill.cuh`, from `:892` to `:1900`
-- all spell it without the keyword; 762 is the only one that does not.  This
is not our shim making an upstream line unbuildable, which is what the XQA
edit above was; it is a line that no compiler accepts, and reporting it
upstream is the right disposition.

IT WAS INVISIBLE UNTIL THE ARCH REACHED IT.  The line sits under
`#ifdef FLASHINFER_STMATRIX_M8N8X4_ENABLED`, which `mma.cuh:38` defines only
for `__CUDA_ARCH__ >= 900`.  Every arch below that compiles the `#else` branch
and never parses the faulty one, so the bug ships dead and stays dead until
something targets Hopper or newer.  `every_instantiation_compiles` is what
reached it: `attn/attention_mla_fa2` was 1 of 127 roots failing, and the whole
kernel was unreachable at run time on this machine because a JIT root that
does not compile is a kernel that does not exist.

IT CARRIES NO MARKER FOR THE SAME ARITHMETIC AS THE FIRST, AND THE CITATIONS
THAT WOULD BREAK ARE NAMED.  `attn/attention_mla_fa2.cuh` cites this file at
`:78-81`, `:98`, `:879`, `:1061`, `:1100-1120`, `:1124`, `:1125` and `:1132` --
every one of them below 762 -- so a marker line inserted above the edit would
falsify all eight.  A one-token change that costs no lines is worth less than
eight citations that still land where they say.  This entry is the record, and
this file is the place a rebaser is sent.

THE THIRD EDIT IS THREE OF THEM, THEY ARE ADDITIONS RATHER THAN CHANGES, AND
THEY ALL CARRY MARKERS.  The two edits above are one-token changes that could
not carry a marker without falsifying a citation.  These are the opposite
case: each ADDS text, each addition carries its own `// PIE:` marker, and step
1 therefore removes the comment and leaves the code -- which for an addition
is not enough.  Step 3 has to take these out, and here is what it takes out:

    attention/default_decode_params.cuh   `IdType* q_indptr;` appended to
                                          `BatchDecodeParams`, after
                                          `partition_kv`, plus its
                                          `q_indptr(nullptr)` in both of that
                                          struct's constructor init lists.
                                          `BatchDecodeParamsMLA` and the two
                                          Single structs are untouched.
    attention/decode.cuh                  the two `q + batch_idx * q_stride_n`
                                          in the batch decode device function
                                          become
                                          `q + params.q_indptr[batch_idx] * ...`;
                                          and, once `chunk_start`/`chunk_end`
                                          are known, a block whose tile starts
                                          at or past `kv_len` stores the empty
                                          state's partial and returns. A body
                                          replayed from a CUDA graph keeps the
                                          tile split of the lengths it was
                                          captured with; a shorter request has
                                          tiles past its end, and upstream's
                                          `chunk_end - chunk_start` wraps to
                                          ~4e9 there.
    attention/cascade.cuh                 a trailing
                                          `const uint32_t* __restrict__ win`
                                          on the two Persistent kernels, the
                                          merge and its attention-sum
                                          twin, and the two lines each derives
                                          from it: the row count is bounded by
                                          `win[0]` and every `v_merged` /
                                          `s_merged` / `v_sum` write is shifted
                                          by `win[1]`.  A null `win` is the
                                          upstream kernel exactly.

WHY THEY EXIST, IN ONE SENTENCE EACH.  A decode request's query row was
`batch_idx` itself, which is true only while the launch's request numbers and
the plane's rows are one numbering; PIE's engine can hand these kernels a
plane a window sits inside, and prefill has read a `q_indptr` for that reason
all along.  And the cascade merge wrote its output dense from row zero, which
is the same assumption on the output side -- `win` is the `(count, start)`
seat every other kernel in this tree already takes, spelled here for the one
launch of the fa2 family that touches the fire's own planes.

THE COLUMNS MOVED WITH THEM.  The three rows above were re-derived from the
files; `added` was raised by each file's net line delta, which for these edits
includes two continuation lines that are a REWRAP of an upstream line rather
than an addition.  The column has no surviving definition to be exact against
-- the test that re-derived it is gone (see THE TEST IS GONE above) -- and
overcounting by two is the direction that does not claim upstream text is
ours.

HOST C++ IS GONE, AND THAT IS WHY THIS TABLE LOST FOUR ROWS.  The move
carried 28 files.  Four of them contained no device code at all -- not one
`__global__` body, not one `__device__` function -- and they are deleted:

    attention/scheduler.cuh    1     13    1840     87507
    exception.h                1      9     115      4194
    allocator.h                1      7     111      4038
    attention/heap.h           1      6      72      2108

That is 2,138 lines and 97,847 bytes of pure host C++ that NVRTC was being
handed and could never have compiled.  Every one of them was guarded in its
ENTIRETY, so NVRTC saw four empty files; the guards were doing the work a
deletion does, at the cost of shipping the text.  Each has a Rust owner that
predates this deletion:

    attention/scheduler.cuh    src/plan/{decode,prefill,mla,sm90}.rs
    exception.h                src/plan/error.rs
    allocator.h                src/plan/alloc.rs
    attention/heap.h           src/plan/heap.rs

The measurement that made this safe, and it was the one worth running:
`scheduler.cuh` is `#include`d by NOTHING in this tree, and `allocator.h`
and `heap.h` are `#include`d only by `scheduler.cuh`.  Walking the four
compilation roots -- `attn/fa2.cuh`, `attn/attention_mla_fa2.cuh`,
`attn/attention_xqa_mha.cuh`, `cascade/merge_states.cuh` -- reached 40 of
the 43 internalised files, and the three it missed were exactly those.  They
were unreachable weight, not a dependency.

`exception.h` is the fourth and the only one that was reachable, from
`page.cuh` and from `utils.cuh`.  Deleting it therefore deleted two
`#include` lines, which is step 3 above and was, at the time, the only edit
in this tree that a strip could not undo.  It is safe for a reason that is
easy to state and worth stating: because `exception.h` was guarded in its
entirety, `FLASHINFER_CHECK` and `FLASHINFER_ERROR` were ALREADY undefined in
every NVRTC translation unit here.  24 expansions survived it, inside
`#define` bodies NVRTC never expands or inside host launchers NVRTC never
reaches.  Removing the launchers took the second kind with them: 13
expansions remain and every one is a `#define` continuation line.  The
deletion changed what NVRTC sees by nothing at all, and still does.

`scheduler.cuh`'s three `__global__` lines were forward declarations with no
bodies, present only so `cudaOccupancyMaxActiveBlocksPerMultiprocessor` had a
function pointer to ask about.  Under the JIT that query is Rust's, so there
was never a kernel in that file to keep.

AND A FIFTH FILE WENT, FOR THE OPPOSITE REASON.  The four above were deleted
because they are host C++ that NVRTC could never compile.  `fp16.h` is the
other case: 179 lines of pure device text, every one of it `constexpr` and
device-callable, deleted because **no compile has ever reached it.**  It is
Marat Dukhan's software fp32<->fp16 conversion, redistributed by FlashInfer
under MIT, and `attention/prefill.cuh` opened it behind a guard:

    #ifdef FP16_QK_REDUCTION_SUPPORTED
    #include "../fp16.h"
    #endif

Nothing in this workspace defines `FP16_QK_REDUCTION_SUPPORTED`, so the
directive was never taken -- and defining it would not have helped.
`prefill.cuh:4208-4210` selects a `half` QK accumulator only when
`USE_FP16_QK_REDUCTION && is_same_v<DTypeQ, half>`, and PIE's `DTypeQ` is
`bf16` (`attn/fa2.cuh:227-230` states this and says why).  The feature is off
structurally, not by configuration.

It took two shims with it.  `shim/bit` (`std::bit_cast`) and
`shim/boost/math/ccmath/fabs.hpp` (one `constexpr` absolute value)
existed for `fp16.h` and for nothing else, and `shim/README.md` is the
register for their retirement.  Both are marked RETIRED in the list under
"Not upstream, written by PIE" below, which leaves four of its six live.

The measurement, and it is the one that matters here: with `fp16.h` and both
shims deleted, `tests/every_instantiation_compiles.rs` still passes -- every
`__global__` a unit names lowers under NVRTC 13.0, 37s of real compilation.
A file that can be deleted without a compiler noticing was not a dependency.
`prefill.cuh` carries a `// PIE: REMOVED` marker at the include site and its
row above counts it.  The `#ifdef` blocks at `:247` and `:1312` are upstream's
own and are untouched: their `#else` arms are what compiles, and each already
`static_assert`s that a `__half` accumulator needs the flag, so the diagnostic
a future asker gets names the feature rather than a missing file.

AND THEN THE REST OF IT WENT.  Deleting four whole files left the rule
holding in effect and not in letter: `csrc/` still held host C++, it was
merely hidden.  34 guarded regions and 13 unguarded launchers later, it does
not.  Measured against the commit before the removal, per file, lines then
bytes:

    flashinfer/pos_enc.cuh                             614    29148
    flashinfer/attention/prefill.cuh                   579    32444
    flashinfer/page.cuh                                288    13527
    flashinfer/attention/decode.cuh                    270    12161
    flashinfer/attention/default_prefill_params.cuh    266     9578
    xqa/mha.cuh                                        240     9285
    flashinfer/attention/cascade.cuh                   227    10940
    flashinfer/utils.cuh                                93     2730
    flashinfer/attention/mla.cuh                        37     1342
    flashinfer/layout.cuh                               27      818
    xqa/utils.h                                         20       79
    xqa/hostUtils.h                                     12      495
    flashinfer/math.cuh                                 10      550
    flashinfer/fastdiv.cuh                               6      307
    flashinfer/fp16.h                                    5      238
                                                     -----  -------
    15 files                                          2694   123642

That is net: each site keeps a `// PIE: REMOVED` marker naming what stood
there, 40 of them, and those markers are in the count.  What went was 34
guarded regions -- `#ifndef __CUDACC_RTC__` in FlashInfer, upstream's own
`#ifndef GENERATE_CUBIN` in XQA -- carrying 1,503 lines of upstream host
text, and 13 `cudaError_t` launchers carrying 1,055 more.

CITATIONS INTO THESE FILES MOVED, AND THE FIX IS NOT IN THIS DIRECTORY.
Rust in `crates/kernels-cuda/src/` and `crates/driver-cuda/src/` cites
these files by line -- `utils.cuh:116-133`, `decode.cuh:823`, `mla.cuh:1079`
and 238 more.  Removing 2,694 lines moved 202 of those 241 citations.

Whether they were citations into OUR copy or into upstream's is the whole
question, because upstream's numbering never had our guards in it and was
never going to match anyway.  It was settled by discriminator, not by
assumption: of the 64 citations that name a C++ identifier alongside the line
number, 52 find that identifier within +-3 lines of OUR text at the commit
before the removal, 1 finds it in upstream's, and 11 find it in neither.  They
are citations into this tree.  The convention `mha.cu` = upstream and
`mha.cuh` = ours has no counterpart on the FlashInfer side, where both are
spelled `.cuh`, and that is why this had to be measured.

Eleven of the twelve `DISPATCH_*` macros and `FLASHINFER_CUDA_CALL` still
stand, in `utils.cuh` and `mla.cuh`, holding all 13 remaining `std::` tokens
in this tree -- every one a `\`-continued `#define` line, none of them ever
expanded.  They are dead text by the same argument that removed everything
above, and they were KEPT anyway: they are the single most-cited construct in
the two crates, `src/` is not this task's to edit, and deleting them would
move a further 18 citations while the owner is holding a tripwire in
`src/x/abi.rs`.  Dead text that nothing compiles is a smaller debt than a
hundred wrong line numbers, and this paragraph is the reason, not an excuse.
`DEFINE_HAS_MEMBER` is the twelfth and it is not dead and not host: its
`std::false_type`, `std::void_t` and `std::declval` are `<type_traits>`, it
has eleven real expansions, and `shim/` provides all three.

The translation, per file: at and after each upstream-side line below, add
the offset to get the line in the file as it now stands.

    flashinfer/attention/cascade.cuh                788->-227
    flashinfer/attention/decode.cuh                 33->-6  836->-193  1188->-270
    flashinfer/attention/default_prefill_params.cuh 27->-3  132->-66  310->-175  462->-266
    flashinfer/attention/mla.cuh                    30->-5  1158->-37
    flashinfer/attention/prefill.cuh                2571->-194  4364->-579
    flashinfer/fastdiv.cuh                          38->-6
    flashinfer/fp16.h                               22->-5
    flashinfer/layout.cuh                           29->-6  61->-16  142->-27
    flashinfer/math.cuh                             35->-10
    flashinfer/page.cuh                             28->-6  128->-26  166->-45  465->-115
                                                    675->-205  754->-225  795->-247  917->-288
    flashinfer/pos_enc.cuh                          26->-4  33->-7  100->-20  1650->-614
    flashinfer/utils.cuh                            33->-11  41->-12  46->-13  51->-14  467->-93
    xqa/mha.cuh                                     2948->-99  3093->-240
    xqa/utils.h                                     29->-5  55->-20

`xqa/mha.cuh` is the one file whose first offset is deliberate.  Four live
citations sit at `:1598`, `:2783`, `:2804` and `:2814` in `src/x/xqa.rs` and
`src/x/abi.rs`, so the removal at `:56` was padded to be exactly as long as
what it replaced -- 8 lines for 8 -- and nothing above `:2948` moved.  That
was possible there because one marker replaced one small block.  It is not
possible for 2,694 lines, and pretending otherwise is how a file ends up with
padding nobody can explain.

Citations that name a DELETED entity rather than a moved one are a separate
and smaller list, and they cannot be repaired by an offset:
`families/cascade.rs:10` cites `prefill.cuh:4351` for
`BatchPrefillWithPagedKVCacheDispatched` and `:11` cites `decode.cuh:823` for
`BatchDecodeWithPagedKVCacheDispatched`.  Both launchers are gone.  The prose
around them is still true -- it describes what the Rust replaced -- so those
become upstream citations and should say so.

THAT PARAGRAPH IS RIGHT ABOUT ITS TWO AND WRONG ABOUT "SMALLER", AND THE
TABLE ABOVE IT NEEDS A DOMAIN.  Both were found by consuming this ledger to
do the repair, at `a91cadcec`, and both errors point the same way: toward
applying an offset where an offset is destructive.

The table is a STEP FUNCTION, so it is total.  It returns a number for every
input, including a line that no longer exists and a citation that was never
written against the old numbering.  Re-deriving it independently -- `git diff
-U0` from `d737aad29` to here, an exact per-line map -- agrees with it on all
12,384 surviving lines, 0 disagreements.  That is the strongest check
available and it CANNOT see either error, because it ranges only over the
lines the table is valid for.  Two derivations of one quantity catch a
mis-measured quantity; they say nothing about a quantity measured over the
wrong set.  Read the table as: valid at and after each breakpoint, FOR LINES
THAT SURVIVED.  It is silent, not correct, elsewhere.

Errors, measured:

1. The citation population is MIXED, and mixed within single files.  Of the
   27 citations this table would renumber, 21 ALREADY cite the post-strip
   numbering, 1 was wrong before the strip, and 5 are genuinely pre-strip.
   Applying the table wholesale breaks 21 to fix 5.  `fa2/params.rs` cites
   `page.cuh` 45-51 for the seven-field run `page_size..stride_h`, which is
   where those fields are TODAY and not where they were; the same file's
   `fastdiv.cuh:47`/`:48` are pre-strip.  Six `utils.cuh` citations land
   inside their `DISPATCH_*` macro today and inside the wrong macro before.
   The discriminator that settles a citation is its enclosing entity's
   position in BOTH versions, and the answer is per citation, not per file.

   The discriminator recorded above -- 52 of 64 identifiers resolving in our
   pre-strip text -- is reproducible and is not what was wrong.  What was
   wrong is that it was never run against the post-strip text.  A test given
   one hypothesis cannot rank two, and this one errs toward "needs repair".

2. The deleted-entity class has EIGHT members and 104 citations, not two.  The
   two named above are real and are the two most-cited.  The rest:
   `MergeStates` 22, `SingleDecodeWithKVCacheDispatched` 10,
   `VariableLengthMergeStates` 7, `SinglePrefillWithKVCacheDispatched` 3,
   `BatchMLAPagedAttention` 3, and the `__host__` constructors in
   `fastdiv.cuh` and `page.cuh`.  66 of the 104 cite a line PAST THE END of
   the file they name -- `cascade.cuh` is 564 lines and is cited at 637-736,
   `prefill.cuh` is 3788 and is cited at 4176-4356.  That measurement needs
   neither this table nor this document, which is why it is the one to trust
   and the one to re-run.  For every one of the 104 the table yields a
   plausible wrong answer: `decode.cuh:823` becomes `:817`, unrelated MLA
   device code.  They need their entity's NAME, which every `// PIE: REMOVED`
   marker below already carries.

   CORRECTION, same pass, one commit later.  The first version of this
   paragraph said 98 citations and 91 past-EOF.  Both were wrong, and wrongly
   in the way this document exists to catch: 141/91 came from a scan at
   `0b0a4f829` and 98 from a later partial re-run at `a91cadcec`, and I
   published the numerator of one beside the denominator of the other.
   `a91cadcec` deleted 36,616 comment lines from `kernels-cuda/src`, and
   citations live in comments, so the population genuinely shrank between the
   two runs -- neither number was fabricated, and that is exactly why the pair
   read as consistent.  A count is only a count AT a commit.  Re-derived whole
   at `55c3e144c`: 211 citations resolve to a header here, 104 un-repairable,
   66 of those past-EOF and 38 deleted-but-in-range.

   Add 4.  The extractor matched `file.cuh:NNN` and silently ignored a bare
   `` `:NNN` `` continuation -- the form `` `cascade.cuh:649` and `:702` ``
   uses.  425 such continuations exist in the two src trees; 4 share a line
   with an already-broken citation.  So 108 dead numbers across 104 citations.
   The gap was found by reading one line of `families/cascade.rs` while
   repairing it, not by the extractor, which is the ordinary way a grammar gap
   is found.

3. SEVEN of the 104 are not this strip's doing, and finding them is the reason
   to widen the population before believing any of the above.  All seven are in
   `driver-cuda/src/fire/mla_naive.rs` and all cite
   `attn/attention_mla_naive.cuh`, which is OUTSIDE flashinfer/ and xqa/ and
   has not been touched since `d737aad29`.  They rot for an older reason:
   `d69beb982` unmixed that header, moving four host functions
   (`launch_mla_naive_paged_raw`, `launch_mla_mma_paged_raw`,
   `mla_mma_supported`, `mma_detail::smem_bytes`) out of a 736-line file that
   is now 635.  Every citation in that block resolves exactly at `d69beb982^`
   -- `:45` is `constexpr int kMlaNaiveBlock = 256;`, `:682` is
   `inline std::size_t smem_bytes() {`, `:717-723` is the `std::once_flag`
   block -- and four MORE of them, in range and so never flagged as
   suspicious, resolve to nothing today.  Two carry a further off-by-one:
   `:725`/`:726` are cited for `dim3 grid` and the launch, which are `:724`
   and `:725`.  Left unrepaired; that file's citations are a different
   commit's debt, and naming it is worth more than half-fixing it.

   That header's own "# What left, and where it went" is stale by one more
   step: it says the four moved to `kernels-cuda/csrc/src/attn/attention_mla.cu`,
   and the archive crate was deleted at `85c6c674b`.

   And ZERO of the 104 are xqa.  The `xqa/mha.cuh` padding described above --
   8 lines for 8, so nothing above `:2948` moved -- is the only intervention in
   this strip that provably cost its citations nothing.  It is the control for
   everything else on this page.

4. The dead `DISPATCH_*` retention note above says 11 macros and 10
   `std::ostringstream` lines.  Re-derived independently at `55c3e144c`:
   13 definitions under `csrc/`, ALL with zero expansions, and 10
   `ostringstream` code lines (12 hits, 2 of them prose).  The 10 sit in 10 of
   the 13; the note counted the macros carrying the marker and reported that
   as the dead set.  The zero is confirmed -- it is the population that was
   short.  They stay: the condition for removing them was that some pass
   repair the files citing them, and the six `utils.cuh` citations turned out
   to need no repair, so nothing here touched those files.

   ITS OTHER NUMBER IS A NUMBER FROM THE PARAGRAPH ABOVE IT.  The note says
   the macros hold "all 13 remaining `std::` tokens in this tree".  There are
   **12**: `std::cerr` and `std::endl` in `FLASHINFER_CUDA_CALL`
   (`utils.cuh:62`, `:63`) and ten `std::ostringstream`, nine in `utils.cuh`
   and one at `mla.cuh:1111`.  13 is the count of surviving
   `FLASHINFER_CHECK`/`FLASHINFER_ERROR` expansions, which the paragraph two
   above the note states correctly and which the note then reused for a
   different quantity.  Both are `\`-continued `#define` lines and neither is
   ever expanded, so nothing about the argument changes -- only the number.
   Note also that "in this tree" cannot be read literally for `std::`: the
   closure holds 104 `std::` tokens.  85 of the other 92 are device text --
   `is_same_v` 37, `conditional_t` 16, `bit_cast` 14, `is_same` 8,
   `numeric_limits` 4, the two `_type`s 4, `void_t` and `declval` one each --
   answered by `shim/`'s `type_traits`, `bit`, `limits` and
   `cuda/std/limits`, and the remaining 7 are prose INSIDE `// PIE:` markers
   naming host text that is gone (`std::string`, `std::exception`,
   `std::make_tuple`, two `std::ostringstream`).  The 12 are the host `std::`
   still standing in code, which is what the sentence is about.

5. `families/cascade.rs` DOES NOT EXIST, and neither does the directory it
   sat in.  It is named twice above -- once as the home of the two most-cited
   un-repairable citations, once as where the extractor's grammar gap was
   found.  `4f9ccba92` deleted `src/families/` outright, with `unit.rs` and
   `device.rs`, and the file is now `src/x/cascade.rs`; the two citations are
   at `:14` and `:19` of it.  Half of what the paragraph above recommended has
   since
   happened on its own: `:14` cites `cascade.cuh:637-668` for `MergeStates`
   and says "upstream's numbering" in the same parenthesis, which is exactly
   the repair asked for.  `:19` cites `prefill.cuh:4350-4352` for
   `VariableLengthMergeStates` and does not say it.  `src/families/attn.rs`,
   named further down as the passer of `-DGENERATE_CUBIN=1`, went the same
   way; it is `src/x/xqa.rs`'s `XQA_COMMON_OPTIONS` now.

   The census itself has to be re-stamped rather than adjusted, for the reason
   the CORRECTION under 2 gives at length: a count is only a count AT a
   commit, and between `55c3e144c` and here the sweep (`1a08b179a`) deleted
   32,974 lines of tests and examples and the descent took the modules just
   named.  Re-derived whole at `751a0df11`,
   over `kernels-cuda/src` and `driver-cuda/src`: **141 citations name a
   file in these two trees, and 54 of them cite a line past the end of the
   file they name.**  `cascade.cuh` is 564 lines and is cited at 637-736;
   `prefill.cuh` is 3,788 and is cited at 4,191-4,356.  They are concentrated
   in four files -- `src/fa2.rs` 15, `src/x/cascade.rs` 15,
   `driver-cuda/src/fire/flashinfer_fa2_dispatch.rs` 11, `src/x/fa2.rs` 7 --
   with the last 6 spread over `driver-cuda`'s `device/device.rs`,
   `fire/merge_states.rs`, `fire/mod.rs` and `tower/qwen3_vl/attn.rs`.

   Past-EOF is a floor and not the whole class: it is the half a walk can
   settle without knowing the entity's name, and the in-range-but-deleted half
   needs the names every `// PIE: REMOVED` marker below already carries.  The
   66/38 split recorded under 2 was over a population that has since halved,
   so it is left as the measurement it was and not carried forward.  What has
   not changed is the conclusion: an offset cannot repair any of these, the
   marker names the entity, and the repair is per citation.


`xqa/hostUtils.h` is a fifth whole-file deletion in the shape of the first
four: one `inline cudaLaunchConfig_t makeLaunchConfig(...)`, zero device text,
and a single includer -- `mha.cuh`, inside the host block that went with it.
It was found by the deletion, not before it; nothing had ever reached it.

WHAT WAS NOT REMOVED, and why, because a rule applied to 34 of 46 candidates
needs the other 12 named.  `xqa/mha_stdheaders.cuh` and two sites in
`attention/mla.cuh` and `xqa/utils.cuh` hold twelve TWO-SIDED regions: an
`#if` whose host arm is upstream's `std::` and whose `#else` arm is our
device `mha::` replacement.  Deleting the host arm there would leave our
device text standing where step 1 expects upstream's, so a strip would
quietly promote OUR code into a file claiming to be upstream's.  The host arm
is what makes the `#else` legible as a substitution.

TWO MORE STAY AND THEY ARE NOT TWO-SIDED, which the sentence above claimed
until the classification was re-run over all fourteen instead of over the ten
that fit it.  `xqa/mha_stdheaders.cuh:20` and `:976` are one-sided
`#ifndef GENERATE_CUBIN` blocks; `:20` carries four host `#include`s
(`<algorithm>`, `<array>`, `<limits>`, `<utility>` -- the other three in it
are shim-provided), `:976` carries none.  They stay for a different reason
than the twelve, and it is a better one: that file is byte-identical to
upstream, carries ZERO `// PIE:` markers, and the guard is UPSTREAM'S OWN
switch, not ours.  `src/x/xqa.rs` passes `-DGENERATE_CUBIN=1`, so
NVRTC never reads them.  Removing them would turn the one XQA file that needs
no entry in any ledger into a modified file needing markers and a row, to
delete text no compiler reaches.  The owner's own earlier survey recorded
`mha_stdheaders.cuh` as having 0 real host lines, and that was right about
the thing that matters.

AND THE SECOND OF THE TWO IS THE OPPOSITE POLARITY, so "NVRTC never reads
them" is true of one and false of the other.  `:20` is `#ifndef
GENERATE_CUBIN` and NVRTC skips it.  `:976` is `#if GENERATE_CUBIN` -- the
arm NVRTC TAKES, seven `using uint8_t = mha::uint8_t;`-shaped aliases that
promote the file's own device typedefs to global scope, which is how the
`GENERATE_CUBIN` build gets fixed-width names without `<cstdint>`.  It carries
no host text, so it stays for the same conclusion by a different route: there
is nothing there to remove.  The 12/2 split is confirmed -- 14 regions, 12
with an `#else`, 36 host lines in those twelve arms -- and only the
description of `:976` was wrong.

`:20`'s inventory was also short.  It carries SEVEN host `#include`s, not
four: `<algorithm>`, `<array>`, `<cstddef>`, `<cstdint>`, `<limits>`,
`<type_traits>`, `<utility>`.  Four of the seven have a `shim/` answer
today (`cstddef`, `cstdint`, `limits`, `type_traits`) and three do not
(`algorithm`, `array`, `utility`).  The original sentence put `<limits>` on
the unanswered side, which was true before `c6c3ceab1` added a shim for it and
has not been since.  Nothing follows from either split while the block is
skipped; it is recorded because this is the file that says what is in the
tree, and "four" was a number about a directory that had changed.

Those fourteen are NOT all the host text left in `csrc/` -- an earlier draft
of this paragraph said they were, and the `DISPATCH_*` macros below are the
counterexample.  What is left is: fourteen guarded regions here, and twelve
host `std::` tokens -- ten `std::ostringstream`, plus `std::cerr` and
`std::endl` in an eleventh body that holds no `ostringstream` -- inside eleven
`#define` bodies that nothing expands.  Both sets are named, counted, and
argued for.  The "eleven" was right and this sentence undercounted the tokens
by naming only the `ostringstream`s; item 4 above has the derivation.

THE DISCRIMINATOR CHANGED, and the old one is kept here because it was
right about the facts and wrong about what to do with them.  It read: "a
file is deleted when it has NO device text, and guarded when it has some
... those stay guarded, because cutting a fragment out of a file this tree
must keep byte-recoverable buys nothing NVRTC can see: it already sees
nothing there."  Every clause of that is true.  It answers the wrong
question.  The rule is not "can NVRTC see it" -- a guard already settles
that -- it is "is `csrc/` device text and nothing else", and a guarded host
launcher is still host C++ sitting in `csrc/`.  So the discriminator is now:

    host C++ is REMOVED, wherever it is, whatever hides it.

Three things fell out of running that over the whole set instead of over
the three fragments named above, and all three were surprises:

  * The three named fragments were the SMALL case.  `pos_enc.cuh` alone hid
    591 lines -- fourteen `cudaError_t` launchers -- behind one guard.
  * Not all of it was guarded.  Thirteen host launchers in `cascade.cuh`,
    `decode.cuh`, `mla.cuh` and `prefill.cuh`, 1,055 lines, sat OUTSIDE any
    guard and survived only by being uninstantiated templates.  That is a
    much weaker shield than a guard: it holds while nothing names them and
    while the frontend defers lookup of non-dependent names in a template
    body.  A guard is a decision; non-instantiation is a coincidence.
  * `QKVLayoutToString` and `PosEncodingModeToString` had ZERO callers, not
    one apiece.  Both were briefed as having one.  The single hit each was
    the sentence in this file describing them.

Three files have no device text AND no host text -- `attention/mask.cuh` is
an enum, `attention/mla_params.cuh` a parameter struct, `fp16.h` a set of
`static inline` conversions -- so they are neither case and were never
candidates.  What survives in `csrc/` now is device text, comments, and
`#define` bodies that nothing expands: `FLASHINFER_CUDA_CALL` and every
`DISPATCH_*` macro went from many expansions to zero when the launchers
went, and are kept because a `#define` nothing expands compiles to nothing.

ONE ROW WAS ALREADY WRONG WHEN THE MOVE FOUND IT, and the correction is
below rather than hidden in it.  This paragraph first claimed "all four
numeric columns are the numbers they were" -- true of the move, and not a
statement about whether they were RIGHT.  Re-measuring all 28 files against
the table turned up exactly one row that disagreed with its file:

    attention/mla.cuh    was   1     7    1143     54176
                         is    3    27    1163     55245

`f622dcf8d` put three `#ifndef __CUDACC_RTC__` blocks in that file where
there had been one -- the two stray `.template` keywords at `:806` and `:847`
that NVRTC rejects, each keeping upstream's spelling in the guard and ours in
an `#else` -- and it amended step 1 above to discard the `#else`, which was
the right call, but it did not re-run the columns.  So the row read 1 guard
where the file has 3 and 1143 lines where the file has 1163, and the totals
under it were short by the same 20 lines and 1,069 bytes.  The drift is +4
`added` lines wider than a naive count suggests, because the two `#else`
branches are ours and step 1 discards them, so they are added lines too.
`tests/upstream_manifest.rs` would have caught this the moment it ran; it is
recorded here because the fix and the reason belong in the same place.  That
test has since been deleted, which makes this the sharpest illustration on the
page of what its deletion costs: a row that disagreed with its file stood for
months, a test found it in one run, and there is no such run now.

The reason the transform survived is worth writing down, because the version
of this note that predicted otherwise was WRONG.  It said: "a file that moves
must change its `#include` lines, because upstream reaches its siblings by
RELATIVE path (`"../cp_async.cuh"`, `"../utils.cuh"`: 58 such directives, and
not one `<flashinfer/...>` among them).  So the move edits upstream's bytes."
That is true of a FILE and false of a SUBTREE.  A relative directive names a
path relative to the includer, so it keeps resolving as long as the includer
and its target move together -- and this tree moved intact, one `git mv` of
the whole directory.  58 relative directives, zero rewritten.  The ten
`#include` lines that did change are ours, in four of our files:
`kernels/attn/{fa2,attention_mla_fa2,attention_score_capture}.cuh` and
`kernels/cascade/merge_states.cuh`, which reach IN from outside and had to
learn the new prefix.  They also changed bracket style, `<flashinfer/...>` to
`"attn/flashinfer/..."`, which is not cosmetic: `source.rs`'s reachability
walk follows quoted spellings and skips angled ones, so quoting these ten put
the internalised closure inside the graph the tests check.

So step 2 is still the identity and no file below spells `<pie/...>`.  The map
it would have used was empty and lived in `tests/upstream_manifest.rs`, which
is deleted; step 2 stays in the transform because the prediction it was built
for is DEFERRED, not refuted -- separating the device stdlib from the
attention algorithms still rewrites upstream's lines, and the whole risk of
that cut is still that internalising quietly becomes a fork.  With the map's
file gone, whoever makes that cut writes the map down here, where the rest of
the transform already is.  A version bump stays re-fetch plus re-apply.

`grep -rn "// PIE:" kernels/flashinfer/ --include='*.cuh' --include='*.h'`
finds all 56 markers -- 40 in the attention closure and 16 in `comm/`, which
is where the jump came from.  The `--include` is not decoration and this
sentence went without it for some time: THIS FILE sits in the directory being
searched and quotes the marker form throughout, so the bare grep answers a
number larger than 56 and larger again after every edit to this page.  The XQA sentence in
`NOTICE` needs no filter -- that directory holds sources and `LICENSE` and
nothing else.

THE COMM DIRECTORY, INTERNALISED 2026-08-16, AND IT BROKE A RULE THIS
DOCUMENT HAD NEVER HAD TO BREAK.  Three files arrived:

    comm/vllm_custom_all_reduce.cuh    the plain P2P all-reduce
    comm/trtllm_allreduce_fusion.cuh   the fused residual + RMSNorm landing
    fp4_layout.cuh                     one enum, `QuantizationSFLayout`

They are the closure of `kernels_cuda::comm`'s two symbols, which are the
tensor-parallel collectives, and they are what made
`kernels_cuda::comm::CAN_LAUNCH` flip from `false` to `true`.  The other two
files their `#include`s reach -- `utils.cuh` and `vec_dtypes.cuh` -- were
already here and are unchanged; `comm/` reaches them by `../utils.cuh` and
`../vec_dtypes.cuh` exactly as `attention/` does, which is why the move added
one spelling to the carried set (`../fp4_layout.cuh`) and no others.

WHERE THE BYTES CAME FROM, AND IT IS NOT THE VERSION THIS FILE SAYS.  Every
paragraph above says v0.6.15.  These three were copied from

    /root/.venv/vllm/lib/python3.12/site-packages/flashinfer/data/include/
        flashinfer/

whose `_build_meta.py` reads `__version__ = "0.6.14"`, git
`19f1a41e6b21f0c422d775e377b6fdf9a1fc9d23`.  That is a DIFFERENT upstream
revision from the one the twenty-three rows above were taken at, and pretending
otherwise is exactly the drift this document exists to refuse.

What makes the mix safe to record rather than to fix immediately is that the
overlap was MEASURED, with step 1 of the recipe above rather than by eye.
`utils.cuh` and `vec_dtypes.cuh` are the only two files both closures share.
Stripping every `// PIE:` marker block from this tree's copies and diffing the
result against the WHEEL's:

    vec_dtypes.cuh   upstream 2195 lines -> 2195 kept, 0 deleted,
                     0 insertions, 0 replacements  (byte-identical)
    utils.cuh        upstream  551 lines ->  472 kept, 79 deleted,
                     0 insertions, 0 replacements

Zero insertions and zero replacements in both: every line this tree keeps of
those two files is present, unedited and in order, in the 0.6.14 copy as well
as the 0.6.15 one.  The two revisions do not disagree anywhere this tree can
see.  So a rebase has TWO tarballs to fetch until somebody re-fetches the whole
tree at one version, and this paragraph is both the reason and the evidence
that it is not urgent.

TWO REMOVALS TAKE DEVICE TEXT, WHICH NOTHING ABOVE DOES.  Every removal
recorded in this file until now was host C++, under the discriminator "host
C++ is REMOVED, wherever it is, whatever hides it".  Both comm headers got
that treatment too -- `class CustomAllreduce` and its `cuda_error`,
`CHECK_CUDA_SUCCESS`, `IPC_KEY` and `rangeStartAddrAttr` on one side;
`getSMVersion`, `getSMRegisters`, `get_sm_count`, `launch_oneshot_lamport`,
`launch_twoshot_sync`, the two `get_registers_per_thread_*`, `use_oneshot`,
`allreduce_fusion_kernel_launcher` and `allreduce_fusion_op` on the other, 296
lines in one marker.  That is the ordinary case and needs no new rule.

Two removals are NOT that case, and each carries its full argument at the site
as well as the summary here:

  1. THE INTEGER<->HALF-PRECISION `cuda_cast` FAMILY, six specialisations in
     `trtllm_allreduce_fusion.cuh`'s `namespace maths`:
     `cuda_cast<int8_t, half>`, `cuda_cast<int16_t, half2>`,
     `cuda_cast<half2, int16_t>`, `cuda_cast<__nv_bfloat16>(int32_t)`,
     `cuda_cast<__nv_bfloat16>(int8_t)`, `cuda_cast<int8_t>(__nv_bfloat16)`,
     plus `bf1622int16` and the two casts that forward to them.

     Every one needs an IMPLICIT conversion between a built-in arithmetic type
     and a half-precision one -- `make_half2(int8[0], int8[1])`,
     `return static_cast<float>(val);` out of a function returning
     `__nv_bfloat16`, `static_cast<short>(val.x)` on a `__nv_bfloat16`.
     NVIDIA's classes have a converting constructor per arithmetic type;
     `kernels/prelude/device.cuh`'s `pie::f16` and `pie::bf16` declare
     every one of theirs `explicit` on purpose, and THE EDIT THAT IS NOT A
     REMOVAL above is a second, independent record of that explicitness being
     load-bearing.

     Nothing calls them.  Measured: the only references to `maths::` outside
     `namespace maths` are four lines inside `utils::cvt_warp_fp16_to_fp4`,
     and all four name `cuda_abs` and `cuda_max`.  The removed six reference
     only each other.

  2. BOTH `fp32_vec_to_e2m1` OVERLOADS, `float (&)[8]` and `float2 (&)[4]`.
     Their sub-sm_100 bodies call `__nv_cvt_float2_to_fp4x2`, which
     `shim/cuda_fp4.h` deliberately does not declare -- that file's
     banner says a missing fp4 conversion is *"the correct moment to decide
     what that conversion should do, on hardware that can be measured"*, and
     the one GPU here is an sm_89 L40S that cannot run the instruction the
     emulation emulates.  Neither overload has a caller in the file.

Both are step-3 work on a rebase and neither compiles when restored until the
thing it collides with moves first -- the prelude's `explicit` for the first,
a measured `__nv_cvt_float2_to_fp4x2` for the second.  That is the correct
failure and it is why the markers say so at the site.

TWO `// PIE: ADDED` MARKERS, WHICH ARE THE THIRD KIND.  Not a removal and not
a guard: an `#include` line that upstream does not have.

    vllm_custom_all_reduce.cuh   + #include <type_traits>
    trtllm_allreduce_fusion.cuh  - #include <tuple>   + #include <array>

Both are the same fact.  Upstream reaches a name TRANSITIVELY through a header
whose only other content is host: `std::is_same` (used by `upcast`/`downcast`)
comes in behind `<array>`/`<iostream>`/`<limits>`/`<map>`/`<sstream>`/
`<unordered_map>`/`<vector>`, and `std::array` (used in
`allreduce_fusion_kernel_twoshot_sync`'s PARAMETER LIST) comes in behind
`<tuple>`, because libstdc++'s `<tuple>` includes `<array>` for the
`tuple_size` specialisations.  `shim/` answers a directive by the literal
string in it and carries no `tuple`, so the transitive route does not exist
here.  Naming the header that actually carries the name is one line and is
recoverable; leaving the host include in to get the name would have kept host
C++ in `csrc/`, which is the rule the whole strip is under.

AND STEP 1 WAS RUN ON THESE TWO RATHER THAN ASSERTED.  Every other row in this
document rests on the transform having been applied by hand and believed.  For
the two comm files it was executed: strip every contiguous comment block
containing a `// PIE:` line, then `difflib.SequenceMatcher` the result against
the wheel's copy, line for line.  The result:

    vllm_custom_all_reduce.cuh    upstream 518 lines -> 265 kept, 247 deleted
                                  0 insertions
                                  1 replacement: the seven host `#include`s
                                    for `<type_traits>`
    trtllm_allreduce_fusion.cuh   upstream 1772 lines -> 1289 kept, 483 deleted
                                  0 insertions
                                  1 replacement: `<tuple>` for `<array>`

**Zero replacements outside those two, and zero insertions.**  Every line this
tree keeps is upstream's, unedited, in upstream's order; every line it does not
keep is inside a `// PIE: REMOVED` marker's range; and the only two lines that
are OURS are the two the paragraph above names.  That is the property the whole
document asserts, run rather than claimed, on the two newest files.

WHAT THIS VENDORING COST OUTSIDE THIS DIRECTORY, because a reader chasing a
`// PIE:` marker will not find these by grepping here.  All of it is ours and
none of it is upstream's:

    shim/array                     new: `std::array<T, N>`
    shim/cuda/std/optional         new: `cuda::std::optional`, `nullopt`
    shim/cuda_fp16.h               + `__hadd`, `__half2half2`, `__habs`,
                                          `__habs2`, `__half_raw(unsigned short)`
    shim/cuda_bf16.h               + `__hadd`, `__bfloat162bfloat162`,
                                          `__bfloat16_as_ushort`, `__habs`,
                                          `__habs2`, `__hmax`, `__hmax2`,
                                          `__nv_bfloat16_raw(unsigned short)`
    shim/cooperative_groups.h      + `cluster_group`, `this_cluster`,
                                          `grid_group::{cluster_rank,num_clusters}`
    kernels/prelude/device.cuh          `bf16`/`f16`'s bit constructors became
                                          `constexpr`; both gained
                                          `operator=(float)`

`shim/README.md` is the register for the first five and carries each
addition's asking site.  The last is the one that touches a type every
FA2 and XQA instantiation is built on, and it is worth naming twice: `explicit`
was NOT relaxed on any constructor, so `bf16 x = 1.0f;` is still refused; what
was added is `constexpr` (which changes nothing about what compiles, only
where) and an assignment operator NVIDIA's own `__nv_bfloat16` declares.  Both
have their asking site in the file's own doc comments.

Columns: guards, added lines, file lines, file bytes.

`bytes` is the fourth column and the newest.  `lines` alone is blind to an
in-place edit that keeps the line count -- change a constant, reflow an
expression, and all three of the other columns are unmoved.  A byte count is
not a checksum and does not catch a length-preserving edit either, but it
catches everything that changes a file's LENGTH, which is nearly every real
drift, and unlike a checksum it is computed identically by `wc -c` and by
`str::len` so the two sides cannot disagree about the algorithm.  The residue
-- a length-preserving edit outside every guard -- is what the network-side
check covers; `.wiki/driver/new-horizon.md` §23.8 records it being run.

AND THAT PARAGRAPH IS EXACTLY WHERE THE REGENERATION WENT WRONG.  The strip
of host C++ recomputed all four columns with a throwaway script, and the
script measured `bytes` as the length of the DECODED text rather than of the
file.  Rust's `str::len` is a byte count; Python's `len(str)` is a character
count; they agree on every ASCII file and this tree has two that are not.
Upstream's own comments in `attention/prefill.cuh` carry twelve non-ASCII
bytes -- em dashes and a `->` arrow, in `The SF smem layout is a plain flat
byte array` and three lines like it -- so its row was written 12 low, and the
totals row with it.  `xqa/mha_components.cuh` has two more and is not in this
table.  The row said 190667; the file is 190679.

It is recorded here rather than quietly fixed because the sentence above
promised the two sides *cannot* disagree about the algorithm, and that is
true only of `wc -c` and `str::len`.  A third implementation was introduced
to regenerate the table and it was not checked against either.  The columns
below were re-derived once more, reading bytes, and every row now equals its
file -- but the way this was caught is worth more than the fix: the per-file
removal ledger and a whole-tree before/after count were computed separately
and disagreed by 14 bytes, 12 in one tree and 2 in the other.  Neither number
was suspicious on its own.  Two derivations of one quantity are how the
defect surfaced, and the reason `f622dcf8d` stood for months is that nobody
had a second one.

file                                       guards   added    lines      bytes
-----------------------------------------------------------------------------
attention/cascade.cuh                            1      20      578      24483
attention/decode.cuh                             3      18      926      41279
attention/default_decode_params.cuh              0       8      287       9138
attention/default_prefill_params.cuh             4      12      208       6253
attention/mask.cuh                               0       0       30        922
attention/mla.cuh                                4      28     1127      53926
attention/mla_params.cuh                         0       0       84       2251
attention/prefill.cuh                            3      20     3795     191573
attention/state.cuh                              0       0       83       2554
attention/variant_helper.cuh                     0       0      104       5118
attention/variants.cuh                           0       0       97       3223
comm/trtllm_allreduce_fusion.cuh                 0     156     1443      54879
comm/vllm_custom_all_reduce.cuh                  0      40      303      11734
cp_async.cuh                                     0       0      265      10227
fastdiv.cuh                                      1       3       61       2078
fp4_layout.cuh                                   0       0       37       1482
frag_layout_swizzle.cuh                          0       0       75       3481
layout.cuh                                       3       9      117       4956
math.cuh                                         1       3      274       8344
mma.cuh                                          0       0      705      26618
page.cuh                                         9      31      632      28238
permuted_smem.cuh                                1       7      197       7261
pos_enc.cuh                                      4      13     1039      45656
profiler.cuh                                     0       0      149       6708
utils.cuh                                        6      26      497      26136
vec_dtypes.cuh                                   0       0     2194      74561
-----------------------------------------------------------------------------
26 files                                        40     394    15307     653079
12 files are untouched. `LICENSE` is upstream's, copied and not counted.

Not upstream, written by PIE, and NO LONGER IN THIS DIRECTORY:

    bit                          std::bit_cast, for fp16.h        [RETIRED]
    boost/math/ccmath/fabs.hpp   one constexpr fabs, for fp16.h   [RETIRED]
    cstdint                      the eight fixed-width typedefs
    cuda.h                       CUDA_VERSION, derived from __CUDACC_VER_*
    cuda_runtime.h               ushort/uchar, and <cstdint>
    type_traits                  integral_constant and six things built on it

These six sat here for a provenance reason -- they exist BECAUSE FlashInfer
asks for them, so they were filed under FlashInfer.  Two are marked RETIRED:
`fp16.h` was their only asker and it is deleted, so four remain.  The rows
stay because this list records what was MOVED, and both were.  That is the wrong axis.
What they DO is impersonate a header this compiler does not ship, which is
also what `cuda_fp16.h`, `cuda_bf16.h`, `cuda_fp8.h`, `cuda_fp4.h`,
`cooperative_groups.h` and `cuda/{cmath,pipeline,std/limits}` do, and those
eight sat beside our own device text for the mirror-image reason: our own
text asks for them too.  All fourteen are now in `shim/`, filed by role.  That
directory holds twenty today, and its own README is the register of what
was added and retired since; none of the seven additions is upstream's and
none is counted in the table above.  One of the eight has since been RETIRED
-- `cuda/pipeline`, whose only asker in this tree was the vestigial
`#include` that `permuted_smem.cuh`'s row now counts -- so the sentence above
is the set as it moved, not as it stands.

The move cannot be seen from inside a compile.  NVRTC has no `-I`: it matches
the literal string in the directive against the names in `includeNames[]`, and
those names are unchanged, because for a shim THE NAME IS THE CONTRACT -- a
file called `cuda_fp16.h` is answering for NVIDIA's `cuda_fp16.h` and cannot
be renamed without ceasing to do its job.  nvcc needed one line: a second
`-iquote` at each of the four sites that passed the first.  There are no such
sites left -- nothing in the workspace compiles this device text with nvcc any
more -- so that sentence is history and `shim/README.md` records the same
change at length.

One row in this table changed, and not because of the move: see the
correction near the top.  Four rows LEFT it, and that was not the move
either: see HOST C++ above.  Every file still listed moved, together, from
`csrc/vendor/flashinfer/` to `csrc/src/attn/flashinfer/`, byte for byte.

WHAT THE TRANSFORM RECOVERS IS NOW SMALLER, and saying so is the point of
keeping this document.  Eleven files still come back by steps 1 and 2 alone
-- the ones with a `0` in the `guards` column, which is what a `0` there now
means.  Twelve need step 3, and step 3 needs the upstream tarball, because
what a `// PIE: REMOVED` marker records is what a file no longer contains.
The property this document asserts is therefore weaker than it was, and it is
still checkable: every line that REMAINS in the twenty-three files below is
upstream's, and every line that does not is named, counted and attributed.
The columns used to be re-derived from the tree by the same code the test ran,
rather than adjusted by hand -- which is the failure this file exists to have
caught once already, at `attention/mla.cuh`.  The test is gone and so is that
code.  For this revision they were re-derived with `wc -l` and `wc -c` and
every row equals its file; the next reader has to do the same thing, and
nothing will tell them if they do not.
The shim group was done first for a reason that turned out to generalise: it
was the one role whose files are neither upstream's nor included by relative
path from an upstream file, so it moved without editing a single upstream
byte.  Moving the closure as one subtree bought the same property for a
harder case, and the 58 relative directives are why.
