Agent:claude-fastfields-to-teeny
Found while doing fastfields-cuda-impl#29 (phase 5 of fastfields/fastfields-kernels#50). Pre-existing on the teeny branch at HEAD — not introduced by that work, but it blocks it, so flagging separately rather than fixing it inside an out-of-scope PR.
Symptom
make -C fastfields-cuda-lib build/reg_field.o (and reg_flow.o) fails outright:
impl/kernels/regularisers/field/../../stap.h(96): error: calling a constexpr __host__ function
("index_stays_inbounds") from a __device__ function("make_stap") is not allowed.
The experimental flag '--expt-relaxed-constexpr' can be used to allow this.
Cause
bounds.h:
constexprinlineboolindex_stays_inbounds(type b) // implicitly __host__ under nvcc
stap.h:96, inside a CUDEV (= __device__) function:
template <int R, bound::type B, typenameoffset_t>
staticinlineCUDEVFF_INLINE stap<offset_t, R>
make_stap(bound::dyn<B> bu, offset_t x, offset_t n, offset_t s)
{
constexprbool always_inb = bound::index_stays_inbounds(B); // <-- host fn from device fnCUDEV expands to __device__ only (cuda_switch.h:25), and a plain constexpr function is implicitly __host__ under nvcc — so this is a hard error even though the call is in a constant-expression context that could fold. It is invisible to the CPU build, where CUDEV expands to nothing.
Every stencil path in reg_field.h/reg_flow.h reaches make_stap via stencil::fill_taps, so this takes out the entire CUDA regulariser build, not one entry point.
Why nobody noticed
No GPU and, as far as I can tell, no nvcc in CI — this layer is compile+link validated only, and evidently not on every change. fastfields-cuda-lib/build/ in my dev tree still had a stalereg_field.o from before stap.h/stencil.h landed, so a plain make reported "up to date" and only a forced rebuild surfaced it.
Fix
Either is a one-liner; (1) is the better one because the predicate genuinely is used in device code:
kernels/bounds.h — mark it (and any sibling predicate reached from device code) CUHOSTDEV:
CUHOSTDEVconstexprinlineboolindex_stays_inbounds(type b)
fastfields-cuda-lib/Makefile — add --expt-relaxed-constexpr to CXXFLAGS. Works, but it is a blanket relaxation rather than a statement about this function.
I verified that index_stays_inbounds is the only such call on this path: with --expt-relaxed-constexpr added and nothing else changed, reg_field.o and reg_flow.o both build clean (nvcc 12.0). That flag is what I used locally to verify cuda-impl#29; the fix itself belongs here.
Suggested follow-up
Worth considering a minimal nvcc compile job in CI (even just make -C fastfields-cuda-lib build/reg_field.o on a CPU-only runner with the toolkit installed) — this class of bug is invisible to every other gate the project has, and teeny's own tests/nvcc_smoke.cu exists for exactly this reason.
Agent:
claude-fastfields-to-teenyFound while doing fastfields-cuda-impl#29 (phase 5 of fastfields/fastfields-kernels#50). Pre-existing on the
teenybranch at HEAD — not introduced by that work, but it blocks it, so flagging separately rather than fixing it inside an out-of-scope PR.Symptom
make -C fastfields-cuda-lib build/reg_field.o(andreg_flow.o) fails outright:Cause
bounds.h:stap.h:96, inside aCUDEV(=__device__) function:CUDEVexpands to__device__only (cuda_switch.h:25), and a plainconstexprfunction is implicitly__host__under nvcc — so this is a hard error even though the call is in a constant-expression context that could fold. It is invisible to the CPU build, whereCUDEVexpands to nothing.Every stencil path in
reg_field.h/reg_flow.hreachesmake_stapviastencil::fill_taps, so this takes out the entire CUDA regulariser build, not one entry point.Why nobody noticed
No GPU and, as far as I can tell, no nvcc in CI — this layer is compile+link validated only, and evidently not on every change.
fastfields-cuda-lib/build/in my dev tree still had a stalereg_field.ofrom beforestap.h/stencil.hlanded, so a plainmakereported "up to date" and only a forced rebuild surfaced it.Fix
Either is a one-liner; (1) is the better one because the predicate genuinely is used in device code:
kernels/bounds.h— mark it (and any sibling predicate reached from device code)CUHOSTDEV:fastfields-cuda-lib/Makefile— add--expt-relaxed-constexprtoCXXFLAGS. Works, but it is a blanket relaxation rather than a statement about this function.I verified that
index_stays_inboundsis the only such call on this path: with--expt-relaxed-constexpradded and nothing else changed,reg_field.oandreg_flow.oboth build clean (nvcc 12.0). That flag is what I used locally to verify cuda-impl#29; the fix itself belongs here.Suggested follow-up
Worth considering a minimal nvcc compile job in CI (even just
make -C fastfields-cuda-lib build/reg_field.oon a CPU-only runner with the toolkit installed) — this class of bug is invisible to every other gate the project has, and teeny's owntests/nvcc_smoke.cuexists for exactly this reason.