Skip to content

nanobind Based Interface - #2427

Open
philip-paul-mueller wants to merge 235 commits into
spcl:mainfrom
philip-paul-mueller:nanobind-compiled-sdfg
Open

nanobind Based Interface#2427
philip-paul-mueller wants to merge 235 commits into
spcl:mainfrom
philip-paul-mueller:nanobind-compiled-sdfg

Conversation

@philip-paul-mueller

@philip-paul-mueller philip-paul-mueller commented Jul 7, 2026

Copy link
Copy Markdown
Collaborator

Adds a Python <-> C interface that uses nanobind instead of ctypes.
The NanobindCompiledSDFG wrapper is not intended to replace the original ctypes based CompiledSDFG class but to act as an alternative and is thus provided as an opt-in.
The advantages of it is that calling it is much faster than before as (most) argument processing is moved into C++.

However, since nanobind works a bit differently and the interface was designed for performance, not everything that the ctypes interface supports is supported by the nanobind interface as well.

Most Important Changes

Here is a short summary of the most important changes a user should be aware before switching.

  • Before CompiledSDFG made sure that for every instance there was a separate .so file and performed a copy if needed.
    This is no longer possible as the name of the file now has a significance for Python.
    Instead now multiple NanobindCompiledSDFG instances can use the same .so file.
  • The above change implies that the .so file can not have a state.
    All state information is contained in the handle, which is (still) managed by the compiled SDFG object.
  • Before it was possible to unload a .so file, which was done when the CompiledSDFG was garbage collected or finalize() was called.
    This is no longer possible as Python does not allow to unload modules.
    If an SDFG is compiled while its identity is already loaded (see below) it will be silently renamed.
  • The generated .so files are now Python C extensions that are "normal" modules that can be imported.
    They are imported under the dace.generated package, as dace.generated.<magic>.<name>, where <magic> is a hash of the resolved build folder.
    A module's identity is thus the pair (build folder, SDFG name): SDFGs with the same name that build into different folders coexist in one process.
    Loading the same pair again returns the already imported module; the loader verifies that it refers to the same file and fails loudly otherwise.
    The generated C++ types additionally carry a content hash in their namespace, so two different programs sharing an SDFG name can not be dispatched into each other (nanobind's type registry is process wide and keyed by type name).
  • Recompiling an UNCHANGED SDFG whose identity is already loaded reuses the loaded module: the module bakes the source SDFG's pre-codegen content hash, compile() compares against it and, on a match, returns a fresh handle - no rename, no rebuild.
    A changed SDFG still renames and rebuilds (a module cannot be reloaded), and instrumentation changes count as changes (they alter the generated code even though the semantic SDFG hash ignores them).
    Hash instability can only cause a redundant rebuild, never a wrong reuse.
    SDFGs carrying unresolved external nested SDFGs participate as well: the reuse key is computed on a resolved deepcopy (via the new SDFG.load_external_nsdfgs() helper), so the caller's object is never mutated by the lazy loading.
    Gated by compiler.nanobind_reuse_loaded (default: enabled).
  • compile() now returns a handle whose sdfg attribute is an isolated deepcopy (pinned to the build folder) on EVERY return path.
    Previously the cached-binary (use_cache) path and the no-regenerate path passed the original object, so later mutations of the caller's SDFG leaked into the handle.
    Note that this deliberately also applies to the ctypes interface (the paths are shared).
  • The renaming on a collision follows two regimes.
    If the build folder is derived from the configuration (cache set to name, hash or unique) the renamed program also derives its own build folder from the new name, thus it never overwrites the artifacts of the original folder.
    If the build folder was set explicitly (which essentially behaves as cache set to single) the folder is fixed and the renamed program builds in place inside it; keeping such a folder consistent is the responsibility of the user.
    Instead of renaming, an error can be requested through compiler.nanobind_name_collision.
  • NanobindCompiledSDFG no longer stores the (processed) arguments of the previous call.
    This means that calls such as get_workspace_sizes() and set_workspace() now needs the symbols.
    However, they accept the same call arguments as __call__() and initialize().
  • Compilation is significantly longer and the resulted files are larger.
    It is thus recommended to set compiler.build_folder_mode to production.
  • Scalar pyobject arguments are supported: the object passes through as an opaque PyObject* (typical use: forwarded into a callback) and arrives there as the very same object.
    The wrapper keeps the object alive for the duration of the call; the program must not retain the pointer beyond it.
    Arrays of pyobjects are supported too: a numpy dtype=object array binds through
    __array_interface__ (DLPack refuses object arrays outright), and the caller's array keeps
    the contained objects alive for the call - the same lifetime contract as ctypes, which
    also hands out pointers into the caller's buffer.
    pyobject return values are supported and decay to the single contained object, matching
    the ctypes convention (_return_arrays[i].item()).
  • The interface is able to perform symbol inferences in C.
    This means if the symbol argument N is also used as shape/strides of an array and is not provided, then it is taken from there.
    If it is present then no consistency check is performed.
    Note that symbols that are listed in arg_names can not be inferred but can act as source for symbol inferences.
    As an example, assume you have an array with symbolic shape of (a + b,) and you have added b to arg_names, then b can be inferred.
    However not all arrays can be used as source, for example arrays of structs are not considered only fundamental types can be used.
    Since return values are allocated inside the binding after inference ran, inferred symbols may also size return arrays.

Structured Call Signatures (user_args)

As an opt-in, an SDFG can declare a structured call signature through the new SDFG.user_args property.
Its entries are argument names or (nested) tuples of argument names; the property is serialized and thus enters the cache hash.
If it is set, the bindings additionally generate a user_call entry point, exposed as NanobindCompiledSDFG.user_bind_call().
This call path takes the declared positional layout directly (tuple entries arrive as Python tuples), performs no keyword processing at all and is the fastest way to call a compiled SDFG with full argument processing.
Every argument of the SDFG must either be listed in user_args or be inferable; this is verified at code-generation time and violations refuse to generate, naming the argument.
Return values are refused on this path and only primitive scalars, scalar pyobjects and plain arrays may be listed.
An empty string entry is an ignored placeholder slot: the position exists in the caller's convention, accepts any value (including None) and is never read; nested placeholder slots still count toward the tuple length check.
The GPU error-record check still applies (inside the compiled binding), unless disabled through the gpu_error_check property.

Both arg_names and user_args are implemented by the new ArgumentSignatureProperty type.
arg_names is still only allowed to have a flat list of strings, while user_args may have the nested structure.
While the json format is different than the old one, it is possible to load files that were generated before.

Things Still done in Python

There are some things that are not ported to C and are still done in Python.

  • Hooks are still managed in Python.
  • The callbacks are also fully processed in Python and are passed as arguments.

GPU errors are checked in the compiled binding, using the same mechanism as ctypes:
after each program call the binding reads (and clears) the error the generated code recorded for this SDFG (__dace_gpu_last_error) and raises it.
The process-global CUDA last-error slot is never consulted; it is per-host-thread and shared with every other GPU user in the process, so it can carry third-party state.
The check is only compiled in when a GPU target emitted its init/exit pair.
Translating a nonzero __dace_exit code in finalize() still goes through the Python-side GPU runtime.

Return values are allocated by the compiled binding itself (through NumPy/CuPy via the Python API, so ownership and dtype semantics are unchanged), AFTER the compiled symbol inference - so inferred symbols may size a return array.
The binding also builds and returns the result (bare array or tuple); a call with return values no longer needs any Python-side processing.
Whether a caller-provided __return* buffer is accepted is baked in at code generation from compiler.nanobind_allow_return_override (changing it requires a recompile); an accepted buffer is validated against the symbol-derived return shape (struct-element returns excepted).
Return arrays must have offset 0 (violations refuse at code generation), and when do_not_execute suppresses the program run the call returns None.

Good to Know

To achieve best performance the following rules should be followed:

  • Only pass arrays of primitive types.
  • Do not use callbacks (as their processing is done in Python and slow).
  • NanobindCompiledSDFG.gpu_error_check (also a constructor argument) can disable the per-call GPU error check, but it is now a plain in-library read of the SDFG's own error record, so disabling it buys next to nothing.
  • If your call arguments can be decomposed in a nested structure of tuples, i.e. struct of array, use SDFG.user_args and call through user_bind_call().
  • It is no longer possible to copy and rename the generated file.
  • The original ctypes based CompiledSDFG class was renamed to CtypesCompiledSDFG and put into a different file.
    compiled_sdfg.py now contains the definition of a protocol but also exports CompiledSDFG as a deprecated alias.
  • The fast_call() interface is no longer supported, a call can only be made using __call__().
  • An SDFG with external (workspace) memory refuses to run before set_workspace() was called, with a clear error naming the storage type.
    Under ctypes the same call silently dereferences a null workspace pointer (usually a segfault).
    The requirement is per state: finalize() drops the association, so a re-initialized handle must be given its workspace again.
  • Generated code compiles with the same flags as under ctypes:
    nanobind_add_module()'s defaults are disabled (NOMINSIZE, PROTECT_STACK), since they would apply to the whole module and thus override compiler.build_type's optimization level (-Os) and drop the stack protector on the kernels.

Not Supported

  • Return values are arrays, or a pyobject (which decays to the single contained object, as under ctypes).
    Python scalar returns are unsupported on BOTH interfaces - SDFG.validate rejects them, so this is not a compatibility gap.
  • Scalars of the low-precision types bfloat16, float8_e4m3fn and float8_e5m2 are rejected at code generation (they would need nanobind value type-casters).
    Low-precision RETURN arrays are supported: the in-binding allocation imports ml_dtypes (without it NumPy cannot resolve the dtype name) and the pointer comes from the array-interface dict.
    The storage is not special-cased - a GPU return hands the dtype to CuPy exactly as the ctypes allocator does, and CuPy succeeds or raises on its own (CuPy gained experimental ml_dtypes.bfloat16 support in v15.0.0a1; float8 is not covered).
    Arrays of these types ARE supported, but through a weaker mechanism than other arrays: numpy cannot export ml_dtypes-backed arrays via DLPack or the buffer protocol, so the raw pointer is extracted from __array_interface__ (__cuda_array_interface__ on GPU) - the same protocol the ctypes marshaller uses.
    The only check is the itemsize from the interface's typestr; there is no dtype identity and no contiguity check, and such arrays cannot serve as symbol-inference sources nor appear in user_args.
  • SDFG.safe_call() is refused (it hides the compiled object, whose collision rename would make post-call queries on the original SDFG unsound, e.g. get_latest_report()).
    Use compile() and CompiledSDFG.safe_call() instead, and query instrumentation reports via the compiled object's sdfg.
  • The same rename unsoundness affects data-instrumentation report lookups:
    a same-name recompile (e.g. re-instrumenting a program that is already loaded) renames into its own build folder, so an implicit "latest report" lookup on the new compiled object finds nothing.
    Pass the report explicitly (dace.instrument_data(..., restore_from=dreport)); SDFG.call_with_instrumented_data always receives it explicitly and works unchanged.
    Programs sharing one name also share report folders under ctypes already - give programs distinct names when their reports must stay apart.
  • Passing None for an array requires the descriptor to be marked optional, e.g. an Optional[T] type hint.
    ctypes silently passes a null pointer for any array argument.
  • Lossy scalar casts, e.g. float -> int parameters are now rejected.
  • NumPy scalars are accepted unless compiler.nanobind_strict_scalar_cast is specified.
  • Passing float16 scalars as arguments is not possible, but passing float16 arrays is supported.
  • Python lists and __array_interface__-style objects are not coerced to arrays; pass numpy arrays.
  • Views pass zero-copy via DLPack; compiler.allow_view_arguments is a ctypes-marshalling concept and is not consulted.
  • No caching of previous call arguments (_lastargs):
    the workspace getters/setters take the symbol values they need per call.
  • An already-imported module cannot be reloaded after a same-path recompile, CPython cannot unload extension modules, ctypes ReloadableDLL covers that workflow.
    Recompilation under a taken (build folder, name) identity renames instead.
  • Thread-safety model is documented on the class and in the generated code
    (same initialization race as ctypes; no shared per-call state, unlike _lastargs).

Configuration added

The following configurations were added:

  • compiler.interface (ctypes | nanobind, default ctypes):
    Decides which backend is used.
  • compiler.nanobind_strict_scalar_cast (default off):
    Decides if (safe) widening casts are allowed (only applies to Python scalars)
  • compiler.nanobind_allow_return_override (default off):
    Decides if it is allowed to explicitly pass the implicit __return arguments to __call__().
    Consulted at code generation only - the decision is baked into the module, changing the option requires a recompile.
  • compiler.nanobind_name_collision (rename | error, default rename):
    Decides what happens when an SDFG is compiled while its (build folder, name) identity is already loaded: silently rename or refuse.
  • compiler.nanobind_reuse_loaded (default on):
    Reuse an already-loaded, content-identical module on recompile instead of rename-and-recompile (see "Most Important Changes").

Modified Tests

Here is an explanation for all skip that were added.

Python lists not Coerced (skip_list_arg_on_nanobind)

ctypes' marshaller runs np.asarray on every argument, so a Python list passed for an array parameter is silently converted to an array (and any writes to that array are not propagated).
The nanobind binding is nb::ndarray<T>.noconvert(), which accepts only objects exposing the DLPack or buffer protocol, a bare list has neither and is rejected at dispatch.

__array_interface__-Style Objects not Coerced (skip_arraylike_args_on_nanobind)

ctypes reads NumPy's __array_interface__ protocol to obtain a data pointer, so an object exposing it (wrapping real memory) is passed by reference.
nanobind's nb::ndarray speaks DLPack / buffer protocol, not __array_interface__, so such an object is not recognised at dispatch and leads to an error.
(The low-precision types are the deliberate exception: their arrays bypass nb::ndarray and read __array_interface__ directly, see "Most Important Changes".)
Implementing it in C would mean to reimplement the array processing in nanobind.
A Python fallback would be possible, but be slow.

Module Reload After Same-Path Recompile (skip_lib_reuse, skip_recompile_folder_mode, skip_recompile_reload)

These assert that recompiling an SDFG in place yields a fresh, reloaded library within the same process.
ctypes wraps the shared object in ReloadableDLL (dlclose + dlopen).
A nanobind module is a CPython extension module, which cannot be unloaded or reloaded once imported, a recompile is therefore renamed (into its own build folder for configuration-derived folders, in place for explicitly set ones), and a same-path reimport returns the already-loaded module.
The reload semantics cannot exist under nanobind.

ctypes Pointer-Array Input Form (skip_ctypes_pointer_array_on_nanobind)

A ContainerArray (array of pointers) can be supplied on ctypes as a ctypes pointer array ((POINTER(c_double) * m)(...)).
nanobind takes the same pointers as a NumPy uint64 array instead.
It is the ctypes-specific input form that is skipped, the nanobind form has its own container-array test.

View Rejection not Applicable (skip_view_rejection_on_nanobind)

compiler.allow_view_arguments is a ctypes-marshalling switch that, when off, rejects non-contiguous / view arrays.
nanobind passes any DLPack-compatible view zero-copy and never consults that config, so the "views are rejected" assertion does not apply.

Return-override is opt-in (skip_return_override_on_nanobind)

ctypes lets a caller pass their own buffer for a __return value.
nanobind forbids this by default, this can be enabled by setting compiler.nanobind_allow_return_override (baked in at code generation; changing it requires a recompile).
The default-behaviour assertion differs, the opt-in path has its own nanobind test.

Lossy Scalar Cast Rejected vs. Warned (test_bad_cast_csdfg)

ctypes' marshaller emits a Casting UserWarning and truncates a lossy scalar, e.g. 0.1 -> int parameter.
nanobind rejects a lossy cast outright (overflow-checked casters), so there is no warning to assert.
This is an implementation detail of nanobind.

SDFG.safe_call() Refused (skip_sdfg_safe_call_on_nanobind)

SDFG.safe_call() compiles internally and hides the compiled object; a compiled nanobind module cannot be reloaded, so a recompile under an already-loaded identity is silently renamed into its own build folder.
Post-call queries on the original SDFG (e.g. get_latest_report()) would then silently look in the wrong folder, hence the method raises on the nanobind interface.
The four tests exercising it are skipped; the *_precompiled variants cover the supported route (compile() + CompiledSDFG.safe_call(), reports queried via the compiled object's sdfg).

Weakened

  • Missing-argument error type: a missing program argument raises KeyError
    on ctypes and TypeError from the nanobind dispatcher; the test accepts both.
  • Workspace symbols per call: the external-memory API takes the symbol values
    it needs per call under nanobind (there is no _lastargs cache),
    so the test passes symbols to the workspace getters under nanobind and none under ctypes.

Claude and others added 30 commits July 6, 2026 09:48
Implements the first slice of the nanobind port: with
`compiler.interface=nanobind`, the program library is built as a nanobind
extension module (single artifact, exporting both the C ABI and
PyInit_<name>), loaded via importlib under `dace.generated.<name>`, and
called through a generated C++ handle that marshals arguments and
releases the GIL around `__program_<name>`. The ctypes path is untouched
and remains the default; an INTERFACE marker file (analogous to
FOLDER_MODE) records which interface produced a build folder.

Covers: CPU arrays and scalars, lazy init, one-module-many-handles
builder, no-implicit-conversion array casts (wrong dtype raises).
Not covered yet: GPU arrays, workspace/external memory, callbacks,
name-collision handling, return values, load-only path details.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…type fix

- SDFG.is_loaded() dispatches on compiler.interface: on the nanobind path
  it is a sys.modules lookup on dace.generated.<name>, which makes the
  existing rename-and-recompile loop in SDFG.compile() handle module-name
  collisions (silent sys.modules increment, as today's behavior).
- NanobindCompiledSDFG allocates __return* arrays (fresh each call,
  symbolic shape evaluation in Python) and returns them, with an early
  exit that skips all shape logic when the SDFG has no return values.
- The generated handle struct is now named uniquely per SDFG
  (DaceHandle_<name>): nanobind shares internals across NB_STATIC modules
  and keys bindings by std::type_index, so identically-named handle types
  from different generated modules collided - a handle minted by one
  module could dispatch into another module's call method.

Existing suite sample under DACE_compiler_interface=nanobind: 30 passed
(codegen: alias/arraywrite/atomic_xchg/dependency_edge; python_frontend:
assignment_statements/augassign_wcr). ctypes path unchanged and green.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…_gpu_code

Addresses review comments:
- Bound methods now use typed positional/keyword parameters, so argument
  matching and casting happen in nanobind's dispatcher instead of
  hand-written kwargs lookups. Bound-parameter order is the user-facing
  positional order (sdfg.arg_names first), while __program is called in C
  signature (arglist) order. A trailing nb::kwargs parameter absorbs extra
  keyword arguments, as the old interface allowed. The Python shell
  forwards *args/**kwargs directly when there are no return values; the
  name-mapping loop remains only on the return-value path (shape
  evaluation needs the symbol values).
  Note: no nb::call_guard - it would copy Python parameter objects with
  the GIL released; the GIL is released explicitly inside the body around
  the pure-C init/program calls instead.
- __dace_exit status is now checked: the state counts as deallocated even
  on failure, and finalize() raises on a nonzero exit code (old-interface
  behavior); the destructor never throws.
- has_gpu_code property on handle (baked in at codegen time, same
  detection as the ctypes CompiledSDFG) and on the shell.
- nanobind_bindings import moved to compiler.py top level.

Tests: 6/6 nanobind interface tests green (2 new: positional+extra-kwargs
absorption, has_gpu_code); parity sample 23 passed; ctypes baseline green.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Accessing state_pointer on an uninitialized or finalized handle now
throws instead of returning 0; a raw null pointer handed to callers of
get_state_struct-style APIs would only fail later and far away.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…odule

Addresses review comments:
- load_nanobind_module / load_nanobind_compiled_sdfg (and the
  GENERATED_NAMESPACE constant) move to compiler.py, next to their ctypes
  equivalents (get_program_handle / load_precompiled_sdfg), so both
  interfaces are used transparently through the same entry points.
- The interface tag file requested in review is the existing INTERFACE
  marker: generate_program_folder() consults Config and writes it;
  loading inspects only the marker (missing marker = old-style ctypes
  folder). get_program_interface() now documents this rule explicitly.
- nanobind_support.py -> nanobind_compiled_sdfg.py; only the wrapper
  class remains, documented like CompiledSDFG (same tasks, delegated to
  the generated module; single fast path instead of the three-step
  advanced interface).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
pyobject arguments/returns (including the PR#2206 bug-compatible decay
of pyobject return arrays in _initialize_return_values) are part 2 of
the port. Until then the bindings generator refused nothing: pyobject's
ctype is not a C++ type, so such SDFGs died with an opaque C++ compile
error. The generator now raises NotImplementedError with a pointer to
the ctypes interface, and _allocate_return_arrays documents why the
bug-compat block is intentionally absent and where the decision lives.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…error option

Name collisions on the compile path now follow the proposal:

- A collision-renamed program (name_0, name_1, ...) is compiled into its
  own build folder, derived from the new name, instead of in-place next
  to the original artifact - the in-place recompile left the folder with
  artifacts that no longer match its program.sdfgz.
- New config option compiler.nanobind_name_collision (rename | error,
  default rename): "error" refuses to compile under a name that is
  already loaded, for pipelines that consider a silent rename a bug.

Both are gated on compiler.interface=nanobind; the ctypes path keeps its
historical in-folder rename untouched.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
get_exported_function

- get_workspace_sizes/set_workspace: the external-memory entry points
  take the init symbols as arguments, so the generated handle stores the
  init symbol values at initialization (the ctypes wrapper's _lastargs
  equivalent). The handle speaks raw StorageType values (baked in at
  codegen time); the shell converts to the Python enum, which is not
  exposed to C++. DaCe's existing external_memory_test passes unmodified
  under DACE_compiler_interface=nanobind.
- state_fields(): the frame generator's statestruct declarations are
  attached to the Frame CodeObject and the pointer-field names baked
  into the module - replaces the ctypes path's regex parsing of
  generated sources.
- get_state_struct(): returns the raw state pointer (with disclaimer),
  per the proposal; combine with state_fields() if needed.
- get_exported_function(): ctypes.CDLL on the imported module file (same
  library handle), with the wrapper attached as __compiled_sdfg__
  keep-alive; returns None for absent symbols.
- initialize() accepts positional arguments like the ctypes wrapper.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ping

Review question: the mapping looks unnecessary in initialize() (unlike
__call__, no return-shape evaluation happens here). It is needed for a
different reason: the generated C++ initialize's positional parameters
are the init symbols only (e.g. initialize(int N, **kwargs)), while
callers pass positionals in user-facing arg_names order, as in
initialize(a, N=20) from DaCe's external_memory_test. A pass-through
would match `a` positionally against `N` - or silently accept an
integer first argument as N. The workspace test now pins the positional
call form.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Philip Mueller and others added 28 commits August 10, 2026 08:42
Temporary, part of a four-arm matrix; reverts to a plain pytest invocation once
the matrix has been walked.

The matrix
----------
  arm 1  ctypes,   nanobind test file ignored   -> 3 consecutive passes  [done]
  arm 2  nanobind, nanobind test file ignored   -> this commit
  arm 3  nanobind, nanobind test file enabled
  arm 4  ctypes,   nanobind test file enabled   (the original configuration)

Arm 1 established that the job is stable when no nanobind code runs at all. If
arms 2 and 3 also pass and arm 4 fails, the fault lies in mixing the two
interfaces within a single job rather than in either interface on its own.

This arm is genuinely nanobind-only
-----------------------------------
Five tests pin compiler.interface to ctypes explicitly and would therefore keep
running ctypes even here, but none of them is reachable in this job: they live in
tests/codegen/compiled_sdfg_protocol_test.py and tests/codegen/external_memory_test.py,
neither of which has any gpu-marked test, and in the non-gpu
test_get_state_struct_refused_in_production_folder_mode. The one gpu-marked test
in tests/parse_state_struct_test.py does not pin an interface. So the GPU test
set really does run entirely on nanobind here.

Two corrections to the plan this arm belongs to
-----------------------------------------------
1. tests/library/reduce_test.py::test_multidim_gpu is multi-DIMENSIONAL, not
   multi-GPU. It is parametrized over itertools.product(_impls, _case_params) =
   4 implementations x 11 reduction shapes = 44 tests, and the failing id
   [CUDA (device)-test_case17] resolves to the 'CUDA (device)' implementation
   reducing a [1000000] float64 array along axis 0 to a scalar. Excluding "the
   multi gpu tests" would therefore not exclude it. If the intent is to drop the
   heavy reductions, the discriminator is size: the same list contains a
   [512, 555, 257] float64 case that allocates ~584 MB, which is a plausible
   problem when 32 workers share 4 devices.

2. Worth noting for the arm-4 hypothesis: this branch already modifies
   dace/libraries/standard/nodes/reduce.py and adds
   tests/library/reduce_cub_workspace_test.py, i.e. it changes CUB reduction
   workspace handling -- and the failing test is a CUDA-device reduction. That is
   a lead independent of the ctypes/nanobind interaction theory and should not be
   ruled out by the matrix result alone.
… syncdebug

Root cause evidence
-------------------
The intermittent GPU CI failures all report "invalid device ordinal" from the
same generated line. Regenerating the exact failing translation unit locally
(codegen needs no GPU) identifies it:

    test_sdfg_cuda.cu:67
    DACE_GPU_CHECK(cub::DeviceReduce::Reduce(nullptr, __cub_ssize_test_sdfg_0_2, ...))

That is the CUB temp-storage SIZE QUERY inside __dace_init_cuda. All three tests
seen failing this way - tests/wcr_cudatest.py, block_allreduce_cudatest.py and
reduce_test.py::test_multidim_gpu[CUDA (device)] - are CUB reductions. The
failures appeared under BOTH the ctypes and the nanobind interface, on different
tests, so the interface was never the variable; CUB is the common factor.

Why the query is the victim rather than the culprit
---------------------------------------------------
A size query passes d_temp_storage = nullptr and launches nothing. It does,
however, consult the CUDA runtime, and CUB checks the runtime's error state
internally. That state is per-host-thread and shared with every other GPU user in
the process, so a value another party left pending is returned as though the
query had failed. The size query is frequently the first DACE_GPU_CHECK-wrapped
call in a generated module, which makes it the usual place a foreign error
surfaces.

This is the same reasoning that already motivates __dace_gpu_last_error reading
the SDFG's own record instead of the runtime slot; the init path simply had no
equivalent protection.

The change
----------
__dace_init_cuda now drains the runtime's pending error before doing any work of
its own, and reports what it discarded, naming it as not ours. A stale
non-sticky error is then harmless, and the CUB query below is judged on its own
result.

Deliberately a warning rather than an exception: an error pending on entry is by
definition not this SDFG's, and throwing would fail an innocent SDFG - the very
symptom this is meant to remove. If the pending error is sticky, the next real
call fails anyway and is reported through the normal path.

Also exports DACE_compiler_cuda_syncdebug=1 in the GPU job, which the failure
message itself recommends: it surfaces errors at their origin instead of at the
next synchronization point, which is what distinguishes "CUB failed" from "CUB
inherited someone else's failure".

Verified on the host: the drain is emitted inside __dace_init_cuda at line 54 and
the CUB size query at line 82 of the regenerated failing file, so the ordering is
correct. gpuError_t, gpuGetErrorString and {backend}GetLastError are defined for
both the CUDA and HIP backends.
The previous commit changed two things at once - it drained the runtime's pending
GPU error at the top of __dace_init_cuda AND enabled synchronous debugging - and
the job then passed twice. That does not attribute the fix: syncdebug serializes
GPU operations, so it shifts timing enough to suppress this class of failure on
its own.

Removing it while holding everything else constant isolates the variable. If the
job keeps passing, the drain is what fixed it. If the failures return, syncdebug
was masking them and the drain is not sufficient.

It has to come out regardless: it is a debugging aid, and leaving it on would
permanently hide the asynchronous-error behaviour that produced these failures in
the first place.

Note for future debugging: the drain's warning cannot be observed in a passing
run. It is a printf from generated code, so pytest captures it and only shows
captured stdout when a test FAILS. Confirming that the drain fired on a green run
would need -rP or --capture=no.

Remaining experiment scaffolding, to be removed once the fix is confirmed:
DACE_compiler_interface=nanobind and --ignore of the nanobind interface test file.
Removes the --ignore added while bisecting the intermittent GPU failures. That
exclusion has served its purpose: the failures turned out to be a CUB
temp-storage size query in __dace_init_cuda inheriting another party's pending
GPU error, which appeared under both interfaces and had nothing to do with these
tests.

Two green runs now stand with the drain fix and without
DACE_compiler_cuda_syncdebug, so the codegen change is carrying the result on its
own rather than being masked by synchronous execution.

This restores five GPU tests: gpu_arrays, gpu_return_values, gpu_has_gpu_code,
gpu_workspace and gpu_callback. gpu_callback has never executed on hardware -- it
was added after the run that segfaulted in the since-removed container-array
test, and every run since has either crashed before reaching it or excluded the
file. Treat a failure there as new information about that test rather than as a
regression of the fix.

One piece of scaffolding remains: DACE_compiler_interface=nanobind. Removing it
returns the job to ctypes, which is the upstream default and the real
configuration, and should be the next and last step.
Removes the last piece of scaffolding, DACE_compiler_interface=nanobind. The job
is back to its normal configuration: 32 concurrent workers, the ctypes interface
(the upstream default), the nanobind interface tests running, no debugging flags.
ci/cscs_gpu.yml now differs from main by a single line, the 'not mpi' marker,
which predates this work.

What the bisection concluded
----------------------------
The intermittent failures were never about the interface. They were a CUB
temp-storage size query in __dace_init_cuda reporting a GPU error it had
inherited rather than caused: the CUDA runtime's last-error slot is per-thread
and shared with every other GPU user in the process, and the size query is
typically the first DACE_GPU_CHECK-wrapped call in a generated module. The same
"invalid device ordinal" surfaced under both ctypes and nanobind, on a different
test each run, which is what the four-arm matrix was measuring without knowing
it.

__dace_init_cuda now drains that slot before doing any work of its own and
reports what it discarded. Serialization and DACE_compiler_cuda_syncdebug were
diagnostic steps along the way and have both been reverted; the drain is the only
behavioural change that remains, and it held across green runs after syncdebug
was removed.

Not proven, and worth stating plainly: the drain warning was never observed
firing, because pytest hides captured stdout for passing tests. The evidence is
failures before and green runs after, not a direct observation of the mechanism
engaging.

Still open, neither blocking
---------------------------
- tests/persistent_fusion_cudatest.py::test_persistent_fusion is a confirmed race
  (it failed and passed on the same commit and configuration). It is now
  instrumented to classify a mismatch, but not repaired.
- dace/codegen/targets/cuda.py hardcodes the device ordinal in
  cudaDeviceGetDefaultMemPool(&mempool, 0) and checks neither that call nor the
  cudaMemPoolSetAttribute after it. An unchecked call that can leave an error
  pending is exactly the hazard this commit's drain exists to absorb.
… returns

Three ctypes-compatibility gaps closed. All three are variations on one theme:
nb::ndarray cannot ingest the dtype, so the pointer comes from the
array-interface dict instead - the same protocol the ctypes marshaller reads.

pyobject ARRAY arguments
------------------------
A numpy dtype=object array is a flat run of PyObject* slots. DLPack refuses
object arrays outright ("DLPack only supports signed/unsigned integers, float and
complex dtypes"), exactly as it refuses the ml_dtypes-backed low-precision types,
so the argument binds as nb::object and the pointer is read from
__array_interface__. Object arrays report typestr '|O' with NO size suffix, so
the guard checks the kind letter rather than an itemsize.

Lifetime is the caller's array: it owns the references to the contained objects
and the nb::object parameter keeps it alive across the call, so the slots stay
valid while the program dereferences them. ctypes does the same - it hands out
pointers into the caller's buffer.

pyobject RETURNS
----------------
Allocated in-binding as a numpy dtype=object array and decayed to the single
contained object on the way out. That decay is the ctypes convention: it returns
`_return_arrays[i].item()` whenever the return is a pyobject, for a proper Scalar
or for an Array wrapping one. The E2E test asserts nanobind and ctypes against
each other in the same test, since matching that convention is the point.

Note this was previously documented as a shared limitation, which it was not:
ctypes supports pyobject returns via _retarray_is_pyobject. What IS shared is
that scalar returns of other types are rejected by SDFG.validate on both
interfaces - so those were never a compatibility gap.

Low-precision RETURN arrays
---------------------------
The in-binding allocation now imports ml_dtypes first: without it NumPy cannot
resolve the dtype NAME (np.dtype("bfloat16") raises TypeError), and CuPy
re-exports numpy.dtype so it needs the registration too. Extraction skips the
nb::cast<nb::ndarray>, which can never work for these dtypes.

Storage is deliberately NOT special-cased. A GPU return hands the dtype to CuPy
exactly as the ctypes allocator does, and CuPy succeeds or raises on its own;
refusing it in codegen would make nanobind stricter than ctypes for no reason.
Practical consequence: a GPU low-precision return is expected to fail on the
current CI image - CuPy gained experimental ml_dtypes.bfloat16 support only in
v15.0.0a1 (an alpha, which `uv pip install cupy` does not select) and float8 is
not covered at all. It fails the same way under ctypes.

Tests
-----
test_nanobind_interface_pyobject_rejected asserted the two refusals removed here
and is replaced by five tests: binding-shape assertions plus an E2E for each
feature. All are CPU-testable; none needs a GPU.

The skip_pyobject_return_on_nanobind markers in callback_autodetect_test.py and
callee_autodetect_test.py are removed - those tests were confirmed to pass under
nanobind before the markers were deleted, rather than assumed to.

Verified: 179 passed under ctypes and 174 passed / 7 skipped under
nanobind+production across nanobind_interface_test, callback_autodetect,
callee_autodetect and retval_test. The remaining nanobind skips are the
by-design divergences (no list / __array_interface__-style coercion, opt-in
return override).

HANDOFF.md and pr_description.md are updated to match (they live outside this
repository).
Follow-up to 5115115, which added pyobject ARRAY arguments without extending
the user_call eligibility check.

The fast path has no setup scope - _user_call_binding ends its argument walk with
`assert not setup  # the scope restriction excludes everything that needs setup`,
and the eligibility rules are what uphold that. pyobject arrays take their
pointer from __array_interface__ via a setup statement, exactly as the
low-precision arrays do, so they have to be excluded for the same reason. They
were not, and the omission surfaced as a bare AssertionError instead of a usable
message.

pyobject SCALARS remain eligible: they forward `.ptr()` inline and need no setup.

The other two features from that commit do not reach user_call at all - it
refuses SDFGs with return values outright, so neither pyobject returns nor
low-precision returns can apply there.

Covered by two additions to test_nanobind_interface_user_args_validation: the
array form now raises the documented ValueError, and the scalar form still
generates.
MPI was well covered but only ever on the ctypes interface: heterogeneous-ci.yml
and the MPI leg of gpu-ci.yml both run without DACE_compiler_interface set, so
they take the default. Nothing had ever run an MPI program through the nanobind
CompiledSDFG. This adds one step that repeats the existing "Test MPI with pytest"
leg with the interface pinned to nanobind.

Why there was no reason to expect breakage
------------------------------------------
The MPI data descriptors are not arguments. A ProcessGrid lives in
sdfg.process_grids (SubArray and RedistrArray likewise in their own registries)
and is emitted into the state struct by code generation; it never appears in
sdfg.arglist(). Confirmed directly:

    sdfg.add_pgrid(shape=[2, 1])
    -> process_grids: ['__pgrid']  (ProcessGrid)
    -> arglist:       {'A': 'Array'}

Neither dace/codegen/ctypes_compiled_sdfg.py nor dace/codegen/nanobind_bindings.py
mentions those types at all, so there is no marshalling difference between the
interfaces to go wrong. The nanobind generator's catch-all refusal for unknown
descriptor kinds is therefore unreachable for them.

Verified before adding the step
-------------------------------
Rather than adding a CI leg and waiting to find out, the suite was run locally on
two ranks with mpi4py and pytest-mpi installed:

    DACE_compiler_interface=nanobind DACE_cache=unique \
      mpirun -n 2 --oversubscribe python -m pytest tests/ --with-mpi \
      -m "mpi and not gpu" --ignore=tests/tutorials

    52 passed, 30 skipped, exit 0

(tests/tutorials is excluded only because it fails to COLLECT in this environment
for unrelated reasons - it contains no MPI tests. The CI leg does not need the
exclusion, since the existing ctypes MPI step collects that file fine there.)

So this commit records a property that already held rather than fixing a defect.
The point is that it is now checked: "expected to work" and "known to work" are
different things, and the gap was invisible precisely because every MPI job
silently used the default interface.
…l setup

Follow-up to the earlier init-only drain, addressing review feedback: the drain
belongs on every invocation, not once per state, and the memory-pool calls should
be checked like everything else.

Per-call drain
--------------
The CUDA runtime's last-error slot is per-host-thread and shared with every other
GPU user in the process (CuPy, another SDFG, any library). A value pending in it
on entry is not ours, and the first DACE_GPU_CHECK-wrapped call inside reports it
as its own failure. __dace_init_cuda already drained it, but initialization runs
once per state while contamination can arrive between any two calls.

The drain moves into an exported `__dace_gpu_drain_error`, defined in the
generated .cu and called from two places: the top of __dace_init_cuda (as
before), and the top of every `__program_<name>` invocation. The frame code
declares it and emits the call only when the cuda target is in use, exactly as it
already does for __dace_init_<target> - the .cpp cannot include the CUDA headers,
so the call has to cross into the .cu.

Memory-pool calls now checked
-----------------------------
`cudaDeviceGetDefaultMemPool` and `cudaMemPoolSetAttribute` were unchecked. An
unchecked call that can fail is exactly what leaves an error pending for the next
checked call to inherit - the hazard the drain absorbs - so they are the "usual
suspects" worth wrapping. The other unchecked backend calls were left alone
deliberately: `GetDeviceCount` has its own explicit error handling, the
`DeviceSynchronize` in __dace_exit_cuda has its result returned, and
`LaunchKernel` is followed by DACE_KERNEL_LAUNCH_CHECK.

Two things had to be fixed for that wrapping to be safe rather than harmful
------------------------------------------------------------------------------
DACE_GPU_CHECK records into `__state->gpu_context`, and the pool setup was
emitted BEFORE the context is constructed - so wrapping those calls as written
would have turned a mempool failure into a null/indeterminate dereference. The
pool setup now runs after the context exists, so a failure there is recorded
normally.

Separately, `gpu_context` was a raw pointer with no initializer in a state struct
allocated via `new T` (default-initialized, so the member was indeterminate), and
the runtime warm-up allocation in __dace_init_cuda is checked before the context
is assigned. It now has a `= nullptr` initializer and DACE_GPU_CHECK guards the
recording on a non-null context, printing either way. That is a pre-existing
latent bug on the init error path, independent of this change.

Verification
------------
Code generation only - there is no CUDA toolchain on the development machine, so
the emitted GPU code is checked by regenerating it, not by compiling or running
it. Confirmed on a real GPU SDFG (tests/wcr_cudatest.py): the drain is defined in
the .cu, called from __dace_init_cuda, declared before use in the .cpp and called
at the top of __program_<name>; a pooled SDFG emits both mempool calls wrapped and
places the context construction before them; and a CPU-only SDFG emits no drain
at all.

The CPU path is exercised for real: 130 tests pass across the codegen, external
memory, build cache, compile and state-struct suites. The framecode change
touches every generated program, so that coverage is the point.
All three were mine, and all three were invisible to the checks I ran, which
asserted that expected strings were PRESENT in generated code. None of them can
detect output that is present but malformed, or a path not exercised by the
probe SDFG. Adds tests that do.

1. A format placeholder inside a generated comment
--------------------------------------------------
__dace_init_cuda is built with str.format, and I wrote `{initcode}` inside a C++
comment. It was substituted there like anywhere else: the whole init body landed
in the middle of the comment and the rest of the sentence became a bare
statement. nvcc:

    test_sdfg_cuda.cu(55): error: identifier "often" is undefined
    test_sdfg_cuda.cu(55): error: expected a ";"

Six GPU tests failed to compile. The comment no longer names the placeholder.

The regression test checks the TEMPLATE SOURCE, not generated output: reproducing
this through generated code needs an SDFG whose {initcode} is non-empty (a CUB
reduction), so a test built on a simpler SDFG passes while the bug is present.
My first attempt did exactly that and passed with the bug reintroduced. The
source-level check catches the whole class for every SDFG shape. {backend} is
allowed - it expands to a single identifier and has always been used this way.

2. target_name assumed on every used target
-------------------------------------------
The per-call drain is emitted only when the cuda target is in use, which I wrote
as `any(target.target_name == 'cuda' ...)`. A user-registered code generator need
not define target_name - the codegen tutorial's MyCustomLoop does not - and this
walks every used target, so it raised

    AttributeError: 'MyCustomLoop' object has no attribute 'target_name'

and failed tests/tutorials/tutorials_test.py on the CPU job. The pre-existing
loop below only reads target_name inside a has_initializer/has_finalizer branch,
which is why it never hit this. Now uses getattr with a default.

3. A struct-field initializer broke get_state_struct
----------------------------------------------------
To make gpu_context safe to read before assignment I declared it
`dace::cuda::Context *gpu_context = nullptr;`. But CompiledSDFG.get_state_struct
recovers the layout by parsing these declarations, and its field regex is
anchored at the end of the declaration, so `= nullptr` stopped the parse at that
field and silently dropped every field after it:

    AttributeError: 'State' object has no attribute '__0_persistent_transient'

The declaration goes back to having no initializer. The pointer is zeroed by
value-initializing the whole state instead - `new T()` rather than `new T` in the
frame code - which is what DACE_GPU_CHECK's null guard relies on, and which fixes
the same hazard for every other pointer member rather than just this one.

Verification
------------
The tutorial failure is reproduced and fixed locally: the test fails with the
target_name bug reintroduced and passes with it restored. The template-comment
guard likewise fails with {initcode} put back and names the offending file, line
and placeholder. 264 tests pass across tests/codegen, tests/tutorials,
tests/parse_state_struct_test.py and tests/compile_sdfg_test.py.

The GPU compile itself still cannot be verified here - there is no CUDA toolchain
on this machine - which is precisely how defect 1 reached CI.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants