docs(research): record the CUDA Tile additive-reduction gap

[lusoris]

Summary

Timeboxed spike closing out the CUDA Tile adoption question, run against the
kernel family where Tile looked most plausible — float_vif_compute.

The result is sharper than "not worth it": the Tile API in CUDA 13.3 has no
additive reduction primitive at all. reduce_max, reduce_min,
reduce_bitand, reduce_bitor and reduce_bitxor are builtins; reduce_add
does not exist. Every hot kernel in this fork — float_vif_compute, the
integer_adm accumulators, float_moment, the SpEED covariance pass — is built
from additive reductions.

Measured on an RTX 4090 (sm_89), same 32×32 tile, only the operator differs:

kernel	reduction	REG	SHARED
tile_max_native	reduce_max (builtin)	24	0 B
tile_sum_via_scan	additive, two partial_sum scans	47	4096 B

~2× the registers plus a 4 KB shared-memory staging buffer, to produce one
scalar the current SIMT path gets from a register-only warp shuffle.

Two further findings, both build-system rather than capability:

• A __global__ function cannot call __tile__ code (calling a __tile__ function ... from a __global__ function is not allowed). Tile kernels need
  the separate __tile_global__ entry point, so a Tile port is a second
  parallel entry point plus host-side dispatch, not an incremental kernel edit.
• A plain -cubin build with --enable-tile emits no SASS for tile kernels
  — only .note.nv.tkinfo and GLOBAL:0. Codegen lands only under
  --tilefatbin, in a separate fatbin section: a second build artifact for
  core/src/meson.build and the CUDA loader to manage.

Recommendation unchanged (do not adopt), now resting on a mechanism rather than
a judgement call. Revisit if a future CUDA release adds reduce_add.

Reproducer / smoke-test command

# reduce_add is absent from the API:
grep -oE 'reduce_[a-z_]+' /opt/cuda/include/crt/cuda_tile.h | sort -u
# -> reduce_bitand reduce_bitor reduce_bitxor reduce_max reduce_min

# Build the two kernels from the digest and compare resource usage:
nvcc --enable-tile --tilefatbin -std=c++20 -arch=sm_89 tile_spike.cu -o tile_spike.tilefatbin
cuobjdump -res-usage tile_spike.tilefatbin

The digest carries the full tile_spike.cu source.

Deep-dive deliverables (ADR-0108)

• Research digest — docs/research/2040-cuda-tile-additive-reduction-gap.md.
• Decision matrix — no alternatives: this PR records a measurement, it does not decide anything. The adoption decision it informs is unchanged.
• AGENTS.md invariant note — no rebase-sensitive invariants; no code changes.
• Reproducer / smoke-test command — above.
• Changelog fragment — changelog.d/added/2040-cuda-tile-spike.md, CHANGELOG.md regenerated.
• Rebase note — no rebase impact: docs-only, no upstream counterpart.

Docs (rule 10)

Docs-only PR; no user-discoverable surface changes.

State (rule 13)

no state delta: a research measurement, no bug opened or closed.

🤖 Generated with Claude Code