KernelInterface: sub-group shuffles, votes and a constant width

[vchuravy]

Implements KernelInterface's sub-group contract from JuliaGPU/KernelAbstractions.jl#831 in
the CUDA back-end (CUDACore/src/CUDAKernels.jl). Stacked on #3314 (ka-0.10).

• Shuffles KI.shfl, shfl_down, shfl_up, shfl_xor and their width variants, each
  a single shfl.sync with the full warp mask. They and KI.supports_shuffle only match the
  primitive types CUDA.jl's warp shuffles handle (Bool, Base.BitInteger,
  Base.IEEEFloat); KernelInterface shuffles other isbits types (Complex, structs,
  tuples) as packed UInt32 words. shfl.sync only uses the low 5 bits of an offset, so offsets of
  32 and more become 0, which gives the thread's own value as the contract requires. Offsets
  are compared before they are narrowed, so wide ones (2^32 + 1, 2^64 + 1, a regression
  test) do too.
• Votes KI.sub_group_any, sub_group_all, sub_group_ballot (bit i - 1 for lane i).
• Capabilities: KI.supports_linear_subgroups and KI.supports_independent_subgroups
  are true: warps are formed from consecutive linear thread indices, and are independent.
• Constant width: the warp size is the constant 32 instead of a read of %WARP_SZ, so
  get_max_sub_group_size is a constant of the generated code; the other queries use
  unsigned 32-bit shifts.
• redux.sync: on sm_80+, KI.sub_group_reduce of + min max & | ⊻ on Int32/UInt32
  is a single redux.sync over the full warp mask (not activemask(), which threads that
  aren't converged could read differently). The operators are commutative, so the lane order the contract
  specifies doesn't matter. Everything else uses KernelInterface's ordered fallback.
• Partitionable sub-groups: KI.supports_partitionable_subgroups is true. The functions with a width
  pass the member mask of the caller's segment to the .sync intrinsics instead of the full
  mask: the width shuffles, the votes (masking other segments out of the ballot and shifting
  it to the segment), the redux.sync reductions, and KI.sub_group_barrier(width), which
  becomes bar.warp.sync on the segment. A width of 32 still compiles to the full mask. This
  lets KernelAbstractions put several narrow tiles in a warp: a DG-like kernel with tiles of
  16 runs 2.4x faster with 16 tiles per block than with a block per tile.

The [TEMP] commit gets KernelAbstractions and KernelInterface from vc/ki-subgroup-ops;
drop it once KernelAbstractions.jl#831 is merged.

Tested locally (RTX 5080): KernelInterface testsuite 6091/6091, KernelAbstractions testsuite 3103 pass,
17 broken.

[codecov]

Codecov Report

❌ Patch coverage is 90.47619% with 4 lines in your changes missing coverage. Please review.
✅ Project coverage is 38.34%. Comparing base (34e43e0) to head (41b9f97).

Files with missing lines	Patch %	Lines
CUDACore/src/CUDAKernels.jl	90.47%	4 Missing ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           ka-0.10    #3331      +/-   ##
===========================================
+ Coverage    37.01%   38.34%   +1.33%     
===========================================
  Files          191      191              
  Lines        19296    18698     -598     
===========================================
+ Hits          7142     7170      +28     
+ Misses       12154    11528     -626     

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.