feat(analysis): say how efficiently a transfer moves its bytes, not only how many

[zhen8838]

Why

The memory family counts how many bytes a transfer moves
(src/tilefoundry/analysis/memory.py:183):

amounts.append(TrafficBytes(moving if moved.read else 0, moving if moved.write else 0))

It says nothing about how efficiently those bytes are moved. There is no notion
of a memory sector or of a per-instruction vector width anywhere in
src/tilefoundry/analysis/ (grep -rn sector returns nothing), and none of the
four families in docs/spec/analysis.md:38-45 reports one.

That gap decided two of the four largest performance steps in an AtomSched
experiment series, a bf16 GEMM at M=8192 K=5120 N=17408 on one H200. In both,
the byte count was already correct and identical between the slow and the fast
kernel. What differed was how many 32-byte sectors those bytes occupied, and in
both cases by exactly a factor of two:

global load   l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum
                unswizzled shared layout   1,069,547,520   32.00 sectors/request   50 % used
                128-byte swizzle             534,773,760   16.00 sectors/request  100 % used
              requests identical at 33,423,360; DRAM bytes within 1.9 %
              7.2127 ms -> 2.8519 ms

global store  l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum
                register tile stored directly   17,825,792
                staged through shared memory     8,912,896    (the output needs 8,912,896)
              Elapsed Cycles 3,143,310 -> 2,947,425

Both were found afterwards, by ncu. Both were predictable before running
anything, from the program alone:

• The load case: a cp.async.cg moves 16 bytes per thread, which is the sm90
  maximum, so per-thread vector width was already full. But the lanes adjacent
  in the unswizzled shared layout were 16 KB apart in global memory, so every
  16-byte chunk sat alone in a 32-byte sector. Width was right; cross-lane
  contiguity was not.
• The store case: the accumulator's register layout gives four lanes eight
  consecutive output columns each, so a warp's store touches eight rows of
  16 bytes and fills half of each sector. That is readable straight off the
  register tile's ShardLayout.

In both cases the only inputs needed are the two layouts, the access relation,
the element size and two target facts. A report that stated this per transfer
would have pointed at the fix in each round instead of leaving it to profiling.

What

A new analysis family, per measured transfer (a Call that moves a tile between
storages), reporting two numbers:

per-lane width      bytes one thread moves per instruction
                    = min(target max vector width, bytes contiguous in both src and dst)
sector utilization  bytes a warp touches / (sectors those bytes occupy x sector size)
                    decided by per-lane width x cross-lane contiguity in global memory

These are two different quantities and both have to be reported: the load case
above had full width and half utilization.

There are two cases, and they differ in what can be claimed:

• Exact. One side is register-distributed. Its ShardLayout states which
  thread holds which element, so the lane-to-address mapping is known and
  utilization is computed exactly. The store case is this one.
• Bound. Both sides are shared or global, broadcast to every thread, and the
  per-thread split is chosen below the IR. The analysis reports the best
  achievable: the maximum per-lane width, and whether any lane group can be
  contiguous in global memory and bank-conflict-free in shared memory at the same
  time. The load case is this one: the unswizzled layout cannot have both, the
  swizzled one can.

Contract

• Pure function of the Call's operand types, their layouts, the access relation,
  and target facts. It does not depend on timing or on any other family's output.
• Two target facts it needs, both belonging on the target rather than in the
  analysis: global memory sector size (32 bytes on sm90) and maximum per-thread
  vector access width (16 bytes on sm90). They should be sourced facts, like the
  existing MemoryHierarchyFacts.
• A transfer in the bound case says it is a bound. It must not print a bound as
  if it were the achieved number.
• The per-warp figure is per warp. The number of warps issuing a transfer
  changes the issue count, not the utilization, and the report should not mix
  the two.

Risk

• The bound case is only a bound for as long as the per-thread split of a
  shared-to-global transfer is decided below the IR. If that split becomes part
  of the IR, or of codegen with a stated rule, the bound case becomes exact; if
  the rule changes silently, the bound silently stops being tight.
• Swizzled shared layouts (ComposedLayout with a Swizzle inner) have to be
  read through: the swizzle moves where an element lands, which is exactly what
  decides bank conflicts, while the access relation is stated on the unswizzled
  outer layout. Getting that split wrong would report a swizzled layout as no
  better than an unswizzled one, which is the one comparison this analysis
  exists to make.