fix: widen the CUDA stream parameter to intptr_t and forward it for the distance ops by balbasty · Pull Request #35 · fastfields/fastfields-cuda-lib

balbasty · 2026-08-01T20:38:05Z

Part of #4 (tracked in fastfields-lib: fastfields/fastfields-lib#4). Repo 2 of 5; depends on fastfields/fastfields-cuda-impl#25.

What changed

(a) Width. A cudaStream_t is a pointer and therefore 64-bit, but every public stream parameter in this dispatch layer was int. A real stream handle coming down from Python is silently truncated to its lower 32 bits and reinterpreted as a different, invalid stream. All 128 public stream parameters are now intptr_t, matching the launcher signatures in fastfields-cuda-impl:

distance 12 · posdef 16 · pushpull 16 · pushpull_backward 8 · reg_field 31 · reg_field_rls 4 · reg_flow 31 · reg_flow_rls 4 · resize 2 · restrict 2 · splinc 2

The _reg_stream() helpers in the four reg_* translation units are widened to match; their reinterpret_cast<cudaStream_t>(static_cast<std::intptr_t>(...)) body is now a genuine no-op conversion rather than a widening of an already-truncated value.

(b) Forwarding.dt_euclidean, dt_l1 and dt_mesh marked the parameter /* stream <unused> */, because the impl launchers took no stream at all. Now that they do, the value is threaded through the internal _dt_* helpers and the dtype dispatch macros so the caller's stream reaches the kernel launch.

Deliberately left unforwarded

Two cases still read /* stream <unused> */, and I did not change them:

dt_spline_table / dt_spline_brent / dt_spline_gaussnewton — there is no CUDA spline-distance host launcher yet (impl/distance_spline.h has the kernels but no CUHOST entry point), so there is nothing to forward to.
posdef / resize / restrict / splinc — their impl launchers do accept a stream, so this layer never passing one is a real remaining gap: those ops always run on the default stream regardless of what the caller asks for. It is out of scope for chore(ci): pass CI_SUBMODULE_TOKEN to checkout for private submodules #4, which scopes the forwarding fix to the distance launchers, so I have left it and am flagging it here instead. Happy to open a follow-up issue.

Submodule pin

impl bumped 81079a8 → cc2155c (fastfields/fastfields-cuda-impl#25).

How this was verified

Compiled, really run:nvcc 12.0. make build/distance.o and make build/posdef.o both succeed against the bumped impl. distance.o exercises the full widened + forwarded path end to end (public entry point → _dt_* helper → impl launcher → kernel launch). resize.o and restrict.o also compiled successfully in a later pass.
Not run:splinc.o fails to compile — but this is pre-existing on main, confirmed by compiling the pristine (pre-change) splinc.cpp/splinc.h and getting the identical error (calling a __device__ function("...sqrt...") from a __host__ function). splinc is not in this repo's MODULES list, so it isn't built by CI either. Unrelated to this change.
Not verified — no GPU available. Runtime stream behaviour is unexercised; this is compile+link verification only, as the issue anticipated.

PR series

Workstream: claude-jitfields-to-fastfields

🤖 Generated with Claude Code

…he distance ops A `cudaStream_t` is a pointer and therefore 64-bit, but every public `stream` parameter in this dispatch layer was declared `int`. A real stream handle passed in from Python is silently truncated to its lower 32 bits and then reinterpreted as a different, invalid stream. Widen all 128 public `stream` parameters to `intptr_t`, matching the launcher signatures in fastfields-cuda-impl. `dt_euclidean`, `dt_l1` and `dt_mesh` also stopped forwarding the value: they marked the parameter `/* stream <unused> */` because the underlying launchers took no stream at all. Now that those launchers accept one, thread it through the internal `_dt_*` helpers and the dtype dispatch so the caller's stream reaches the kernel launch. The three `dt_spline_*` entry points keep `/* stream <unused> */`: the CUDA spline distance launcher does not exist yet, so there is nothing to forward to. `posdef`/`resize`/`restrict`/`splinc` likewise still mark it unused here -- their impl launchers do accept a stream, so wiring this layer up to them is a real remaining gap, but it is outside the scope of issue #4, which scopes the forwarding fix to the distance launchers. Bumps the `impl` submodule pin to cc2155c (the matching cuda-impl change). Verified by compiling `distance.o` and `posdef.o` with nvcc 12.0 against the bumped impl. No GPU is available, so this is compile+link verification only. Refs fastfields/fastfields-lib#4 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016AjQcY78NgbagPSbPJRr6Z

balbasty · 2026-08-02T06:48:41Z

Review summary (triage, jitfields-to-fastfields): merged as squash.

128 stream params widened int→intptr_t across distance/posdef/pushpull/reg_field/reg_flow (+ RLS variants), impl submodule pin bumped to the merged cuda-impl#25 fix, distance dispatch (distance.cpp) now forwards the real stream instead of ignoring it.
CI green: build completed/success (real nvcc compile of distance.o/posdef.o/resize.o/restrict.o); splinc.o failure confirmed by the author as pre-existing on main, not caused by this change, and not in CI's MODULES list.
Compile-only gate (no GPU in CI) — appropriate given this repo's own bar; CPU behavior is provably untouched since the change is signature-only plus CUDA-side forwarding.
Independent of the CPU regulariser regression found in the sibling fastfields-cpu-lib#61 PR in this same series — that one is held pending investigation, this one isn't affected by it.

Workstream: claude-jitfields-to-fastfields

Generated by Claude Code

…dary (#22) * fix: widen the CUDA stream parameter to intptr_t at the nanobind boundary Closes#4. A `cudaStream_t` is a pointer and therefore 64-bit, but every `stream` parameter accepted at this boundary was declared plain `int` in the nanobind lambda signatures. A real stream handle from a caller (e.g. a torch/cupy stream pointer forwarded as an integer) whose upper 32 bits are non-zero is silently truncated before it ever reaches the C++ library -- the last and most externally visible link in the truncation chain fixed by this issue across all five repos. Widens all 45 `stream` parameters in `src/ext.cpp` to `intptr_t`. Bumps the `_fastfields_lib` submodule pin to fastfields-lib's 16f3eb7 (fastfields/fastfields-lib#50), which carries the matching widening through the hub and both device-dispatch layers (fastfields/fastfields-cpu-lib#61, fastfields/fastfields-cuda-lib#35) down to the CUDA launchers themselves (fastfields/fastfields-cuda-impl#25). Verified end to end on this machine (CPU path; no GPU available): - `pip install --no-build-isolation .` succeeds: builds `_core.so` via nanobind against the bumped `_fastfields_lib`, links `-lfastfields`, and installs the wheel. - `import fastfields.dlpack._core` succeeds and exposes the expected entry points (dt_euclidean, dt_l1, dt_mesh, field_matvec, ...). - Regression check for the actual bug: calling `dt_euclidean(x, 1.0, 1 << 40)` -- a stream value with bits set above the 32-bit range -- completes without truncation or OverflowError, which a plain `int` parameter could not guarantee. - `pytest tests/`: 36 passed, 0 failed. Refs fastfields/fastfields-lib#4 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016AjQcY78NgbagPSbPJRr6Z * fix: bump _fastfields_lib pin past the RLS broadcast-stride fix fastfields-cpu-lib#62 (field_matvec_rls's RLS/shared-weight broadcast path reading the wrong stride, corrupting output for wc=1, C>1) has now been fixed and threaded all the way up through cpu-impl#48, cuda-impl#27, cpu-lib#64, cuda-lib#38, and fastfields-lib#55. This repo's _fastfields_lib pin (16f3eb7) predated fastfields-lib#55, so the fix was not reachable through this PR's pin -- the extension built here would still have linked against the broken field_matvec_rls. Bumps _fastfields_lib -> fd9e6c5 (fastfields-lib main HEAD, which now carries cpu -> 4d46a02 and cuda -> e96440c, both past the fix). Verified for real: CXX=clang++ pip install --no-build-isolation . cd /tmp && python -m pytest <repo>/tests -q Install succeeds end to end (builds libfastfields.so/-cpu.so via the bumped _fastfields_lib's Makefile, links the nanobind extension, packages and installs the wheel). pytest: 36 passed, 0 failed. Also re-ran the stream-widening regression check this PR introduced (stream = 2**40 via dt_euclidean) against the freshly built extension: still passes. part of fastfields-cpu-lib#62 Co-authored-by: Claude <noreply@anthropic.com> --------- Co-authored-by: Claude <noreply@anthropic.com>

balbasty added the claude-jitfields-to-fastfields label Aug 1, 2026 — with Claude

balbasty merged commit 2111c8f into mainAug 2, 2026
1 check passed

balbasty deleted the claude/4-cuda-stream-intptr_t branch August 2, 2026 06:48

Navigation Menu

Uh oh!

fix: widen the CUDA stream parameter to intptr_t and forward it for the distance ops - #35

fix: widen the CUDA stream parameter to intptr_t and forward it for the distance ops#35
balbasty merged 1 commit into
mainfrom
claude/4-cuda-stream-intptr_t

balbasty commented Aug 1, 2026•
edited
Loading

Uh oh!

Uh oh!

balbasty commented Aug 2, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

balbasty commented Aug 1, 2026• edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What changed

Deliberately left unforwarded

Submodule pin

How this was verified

PR series

Uh oh!

Uh oh!

balbasty commented Aug 2, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

balbasty commented Aug 1, 2026•
edited
Loading