Skip to content

[teeny] Phase B4+B5-cuda: regulariser boundary — reg_field/reg_flow import via from_dlpack<T, gpu_view>, reg_dispatch.h width leaf onto anyrank::reindex #138

Description

@balbasty

Agent:claude-fastfields-to-teeny

Part of the tensor-native-boundaries umbrella (fastfields-lib#57), Phase B. CUDA twin of fastfields-cpu-lib#78, with the CUDA deltas of fastfields-cuda-lib#41 (B0-cuda). One coordinated PR-set with fastfields-cuda-impl#38 (reg_field) + fastfields-cuda-impl#39 (reg_flow); covers this repo's reg_field.cpp, reg_flow.cpp, reg_dispatch.h. Sequencing: LAST of the cuda sets, after cpu-lib#78's gate. Same single-issue justification as cpu-lib#78: reg_dispatch.h (mirrored here, ff::cuda via FF_DEVICE) is shared by both modules, and its Op::run contract changes once for both.

The changes

  1. reg_dispatch.h (cuda mirror) — the dtype×width leaf changes MECHANISM, not policy. Unlike the CPU twin (which deletes the width dimension per D1), CUDA keeps int32 (R5): dispatch_dtype still has two width arms per dtype, but use_32bits (fed by CANUSE32BITS conjunctions) is replaced by the carrier-native test — the wrappers import first, then narrow: all-carriers index_fits<int32_t>()Op::run with reindex<int32_t, FF_REG_*_MAX_RANK>-trimmed carriers, else the reindex<int64_t, FF_REG_*_MAX_RANK> arm (R6 trim in both). Whether the width branch stays in dispatch_dtype or moves into the Op adapters (where the carriers live) is the PR's one structural judgment call — document it; rejection messages and outermost-first order (ndim → bound → dtype) stay verbatim either way.
  2. reg_field.cpp / reg_flow.cpp_field_*/_flow_* wrappers import via from_dlpack<scalar_t, tny::storage::gpu_view> (const-element for read-only operands) and call the carrier entries of cuda-impl#38/reg_field: add field_matvec_rls/diag_rls/relax_rls device dispatch + bump pins #39. copy_if_needed/VOIDPTR/CVOIDPTR/ContiguousStrides uses deleted; CHECK_*, check_selfadjoint_bound, NDIM_MSG, as_weights/vx synthesis, niter/stream plumbing: verbatim.

Gate (CUDA flavor, as cuda-lib#41)

  1. Untouched __global__ kernels PTX-identical, BOTH widths (the int32 instantiations survive here, unlike CPU).
  2. nvcc compile+link; per-launch-family kernel-parameter budget stated (JRLS relax the tight case — see cuda-impl#38).
  3. Full CPU suite green with unchanged counts against the shared teeny pin (clang++ AND g++, true make clean) — including cpu-lib's test_reg_dispatch at its post-tools: add consolidate.sh, the frozen rule set for the six-repo merge #78 baseline.
  4. R7: fastfields' "dlpack.h" before <teeny/dlpack.h> in both .cpp and ahead of teeny includes in reg_dispatch.h.

Scope

The three files above only. NOT other modules, NOT autocast.h itself — note that after this PR-set, autocast.h's remaining CUDA users are only the distance spline/mesh paths (Phases D/E), mirroring the CPU side; its deletion stays tracked at the umbrella.

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions