Uh oh!
There was an error while loading. Please reload this page.
This repository was archived by the owner on Aug 20, 2026. It is now read-only.
fix(bounds): mark index_stays_inbounds CUHOSTDEV - #72
Merged
Conversation
It was a plain host constexpr function called from make_stap (stap.h),
a CUDEV (__device__-only) function -- a hard nvcc error ("calling a
constexpr __host__ function from a __device__ function"), invisible to
the CPU-only build where CUDEV expands to nothing. This took out the
entire CUDA regulariser build (every stencil path reaches make_stap via
stencil::fill_taps).
Closes #71.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for freeto subscribe to this conversation on GitHub.
Already have an account?
Sign in.
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Agent:
claude-fastfields-to-teenyClosesfastfields/fastfields-lib#101.
Why
bound::index_stays_inboundswas a plainconstexprfunction — implicitly__host__under nvcc — called fromstap.h'smake_stap, aCUDEV(__device__-only) function. nvcc hard-errors on this ("calling a constexpr host function from a device function is not allowed"); the CPU-only build never sees it, sinceCUDEVexpands to nothing there. Every stencil path (stencil::fill_taps→make_stap) reaches this, so it took out the entire CUDA regulariser build — found while doing fastfields-cuda-impl#29 (phase 5 of kernels#50), and confirmed pre-existing onteenyat HEAD, unrelated to that work.Fix
One line:
CUHOSTDEV constexpr inline bool index_stays_inbounds(type b).transpose(the other predicate reached near this code, viastencil.h'saxis_bound_adjointstatic member) does not need the same treatment — it's evaluated as a class-template static-member initializer at instantiation time, not called from within a device function body, so nvcc's execution-space check doesn't apply to it; confirmed by this fix alone (no other change) being sufficient.Verification
make -C fastfields-cuda-lib build/reg_field.o build/reg_flow.o(nvcc 12.0,-std=c++17): both compile clean, no--expt-relaxed-constexprneeded — this is the exact command that failed before the fix (per [teeny]bound::index_stays_inboundsis host-only under nvcc — breaks the whole CUDA regulariser build fastfields-lib#101's repro).make -C fastfields-cpu-lib clean && make -j4 CXX=clang++ test: all 11 suites pass, unchanged check counts (test_reg_field525,test_reg_flow1123,test_reg_op186, etc.) — confirms the CPU-only path (CUHOSTDEVexpanding to nothing there) is untouched.Generated by Claude Code