You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
[cuda.compute] Fix device memory leaks caused by caching - #12197
No actionable comments were generated in the recent review.
Summary
Summary by CodeRabbit
Bug Fixes
Cached operations no longer keep device arrays alive unnecessarily after use.
Cache keys better distinguish operators and account for captured values, reducing accidental collisions.
Transform iterators distinguish operations with different behavior when generating their identities.
Clearing the compilation cache removes cached entries, which can be rebuilt as needed.
Pointer and iterator state binding no longer retains the input object; related APIs no longer accept an owner argument.
Proxy-backed iterators use a placeholder address during compilation, with actual device addresses bound before launch.
Performance
Concurrent requests to compile the same operator can share compilation work, while different compilations can proceed concurrently.
Documentation
Keep your own references to arrays until asynchronous work submitted to the CUDA stream completes.
The changes remove owner references from cached pointer state, revise operator cache keys, and add explicit compilation memoization. They also use placeholder addresses for proxy-backed builds. Tests cover array collection, cache keys, proxy iterator state, and cache behavior. Documentation states that callers must retain arrays until asynchronous work completes.
Changes: Cached array retention
Layer / File(s)
Summary
Update pointer state and proxy-backed builds python/cuda_cccl/cuda/compute/_bindings.pyi, python/cuda_cccl/cuda/compute/_bindings_impl.pyx, python/cuda_cccl/cuda/compute/_cccl_interop.py, python/cuda_cccl/cuda/compute/_caching.py, python/cuda_cccl/cuda/compute/_proxy.py, python/cuda_cccl/tests/compute/test_multi_cc_serialization.py, python/cuda_cccl/tests/compute/test_caching.py
Pointer and iterator binding APIs no longer accept owner objects. Factory builds substitute proxies for contiguous device-array arguments. Proxy iterator state uses a placeholder address. Tests cover proxy-backed merge-sort state and array collection.
Function keys handle recursive references and captured values, and retain objects that use identity keys. Transform iterator keys use the operator cache key. Tests check key equality, distinctness, recursion, and captured-array behavior.
compile_op uses an explicit memo and coalesces concurrent same-key misses. Tests check cache clearing and compilation behavior. Documentation states that callers must retain arrays until asynchronous work completes.
important: NVIDIA/cccl#11388: Iterator binding no longer retains the input object. The changes do not make wrappers stateless or thread-shareable.
important: NVIDIA/cccl#12021: Proxy-backed builds use a placeholder address for iterator state, including the merge-sort test path.
Change: Bug fix · Severity of issue fixed: High
Merge Risk: Moderate · up to 88dea
Pointer bindings no longer retain caller arrays, so callers must keep arrays alive until asynchronous work finishes. The documentation does not yet cover arrays captured by operators. The compiled-operator memo is also unbounded, which can grow memory in long sweeps until caches are cleared. These are bounded concerns that should be resolved or explicitly accepted before merge.
Security architecture risk: Low · up to 88dea
The inspected execution paths support the intended separation between cached compilation and caller-owned GPU memory. Callers must retain arrays until asynchronous work completes. No introduced security defect was established, but dependent-path coverage is incomplete.
Retained concerns
No architecture-level concerns identified.
Security review details
Security Blast Radius
inferred — The demonstrated security-relevant surface is allocation lifetime and kernel correctness inside the calling application's CUDA execution path. The inspected path consumes caller-provided operators and array addresses; it does not establish a new remote, cross-tenant or privileged-service entrypoint.
Trust Boundaries and Controls
observed — The non-null proxy address is build-time metadata, not an allocation. Proxy data access raises an error, including through dictionary get, while the inspected execution path retrieves real pointers before launch. This guards against executing a build-only placeholder as an ordinary array.
inferred — A cached helper kernel outliving its original callable is not by itself a demonstrated use-after-free. In the inspected unary-transform path, each launch receives an operator whose closure can retain the helper and captured array. A different pointer changes the baked-array key; identical pointer metadata preserves address-based compiled semantics. This conclusion depends on callers preserving allocation lifetime through stream completion.
- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts
Reviewing files that changed from the base of the PR and between dc60e83 and 88dea82.
Files selected for processing: 2
python/cuda_cccl/cuda/compute/_caching.py
python/cuda_cccl/tests/compute/test_caching.py
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
CUDA Array Interface specification object lifetime producer consumer asynchronous operations
💡 Result:
<source_evidence>
<source kind="open_page">
<title>CUDA Array Interface (Version 3) — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html</location>
<excerpt>citeturn1view0 [wordlim: 200] Crawled: today; Content type: text/html; Source: open({"ref_id":"turn0search0","lineno":null}); Total lines: 361
L0: Back to top
L1:
L2: [Input: Search the docs ...] `Ctrl`+`K`
L3:
L4: Table of Contents
L5: * 【0†User guide】
L6: * 【1†Installation】
L7: * 【2†Writing CUDA Kernels】
L8: * 【3†Memory management】
L9: * 【4†Writing Device Functions】
L10: * 【5†Global Variables and Captured Values】
L11: * 【6†Supported Python features in CUDA Python】
L12: * 【7†CUDA Fast Math】
L13: * 【8†Supported Atomic Operations】
L14: * 【9†Cooperative Groups】
L15: * 【10†Random Number Generation】
L16: * 【11†Device management】
L17: * 【12†Examples】
L18: * 【13†Debugging Numba CUDA Programs with Visual Studio Code and CUDA GDB】
L19: * 【14†Debugging CUDA Python with the the CUDA Simulator】
L20: * 【15†GPU Reduction】
L21: * 【16†CUDA Ufuncs and Generalized Ufuncs】
L22: * 【17†Sharing CUDA Memory】
L23: * 【18†CUDA Array Interface (Version 3)】
L24: * 【19†External Memory Management (EMM) Plugin interface】
L25: * 【20†CUDA Bindings】
L26: * 【21†Calling foreign functions from Python kernels】
L27: * 【22†Compiling Python functions for use with other languages】
L28: * 【23†CUDA device call conventions】
L29: * 【24†On-disk Kernel Caching】
L30: * 【25†CUDA Minor Version Compatibility】
L31: * 【26†CUDA Frequently Asked Questions】
L32: * 【27†Reference documentation】
L33: * 【28†CUDA Host API】
L34: * 【29†CUDA Kernel API】
L35: * 【30†Cache Hints for Memory Operations】
L36: * 【31†CUDA-Specific Types】
L37: * 【32†Memory Management】
L38: * 【33†Libdevice functions】
L39: * 【34†Environment Variables】
L40: # CUDA Array Interface (Version 3)【35†#】
L41:
L42: The CUDA Array Interface (or CAI) is created for interoperability between different implementations of CUDA array-like objects in various projects. The idea is borrowed from the 【36†NumPy array interface†docs.scipy.org】.
L43:
L44: Note
L45:
L46: Currently, we only define the Python-side interface. In the future, we may add a C-side interface for efficient exchange of the information in compiled code.
L47: ## Python Interface Specification【37†#】
L48:
L49: Note
L50:
L51: Experimental feature. Specification may change.
L52:
L53: The `__cuda_array_interface__` attribute returns a dictionary (`dict`) that must contain the following entries:
L54:
L55: * shape: `(integer, ...)`
L56:
L57: A tuple of `int` (or `long`) representing the size of each dimension.
L58:
L59: * typestr: `str`
L60:
L61: The type string. This has the same definition as `typestr` in the 【36†NumPy array interface†docs.scipy.org】.
L62:
L63: * data: `(integer, boolean)`
L64: The data is a 2-tuple. The first element is the data pointer as a Python `int` (or `long`). The data must be device-accessible. For zero-size arrays, use `0` here. The second element is the read-only flag as a Python `bool`.
L65: Because the user of the interface may or may not be in the same context, the most common case is to use `cuPointerGetAttribute` with `CU_POINTER_ATTRIBUTE_DEVICE_POINTER` in the CUDA driver API (or the equivalent CUDA Runtime API) to retrieve a device pointer that is usable in the currently active context.
L66:
L67: * version: `integer`
L68:
L69: An integer for the version of the interface being exported. The current version is 3.
L70:
L71: The following are optional entries:
L72:
L73: * strides: `None` or `(integer, ...)`
L74: If strides is not given, or it is `None`, the array is in C-contiguous layout. Otherwise, a tuple of `int` (or `long`) is explicitly given for representing the number of bytes to skip to access the next element at each dimension.
L75:
L76: * descr
L77:
L78: This is for describing more complicated types. This follows the same specification as in the 【36†NumPy array interface†docs.scipy.org】.
L79:
L80: * mask: `None` or objec…[truncated]</excerpt>
</source>
<source kind="search">
<title>CUDA Array Interface (Version 3) — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html</location>
<excerpt>citeturn0search0 [wordlim: 200] Crawled: 3 months ago; The CAI specification does not prescribe the exact mechanism by which these options are set, or related options that Producers or Consumers might provide to allow the user further control over synchronization behavior. ... Obtaining the value of the `__cuda_array_interface__` property of any object has no effect on the lifetime of the object from which it was created. ... Numba provides two mechanisms for creating device arrays from objects exporting the CUDA Array Interface.
</excerpt>
</source>
<source kind="search">
<title>Numba: docs/source/cuda/cuda_array_interface.rst | Fossies</title>
<location>https://fossies.org/linux/numba/docs/source/cuda/cuda_array_interface.rst</location>
<excerpt>citeturn0search1 [wordlim: 200] Crawled: 3 weeks ago; This follows the same specification as in the NumPy array interface. ... Numba takes no steps to maintain the lifetime of an object from which the interface is exported - it is the user's responsibility to ensure that the underlying object is kept alive for the duration that the exported interface might be used. ... Numba provides two mechanisms for creating device arrays from objects exporting the CUDA Array Interface.
</excerpt>
</source>
<source kind="search">
<title>Numba Documentation</title>
<location>https://numba.readthedocs.io/_/downloads/en/0.53.1/pdf/</location>
<excerpt>citeturn0search12 [wordlim: 200] Published: 5.6 years ago; Obtaining the value of the `__cuda_array_interface__` property of any object has no effect on the lifetime of the object from which it was created. ... Numba takes no steps to maintain the lifetime of an object from which the interface is exported - it is the user’s responsibility to ensure that the underlying object is kept alive for the duration that the exported interface might be used.
</excerpt>
</source>
<source kind="search">
<title>User guide — Numba CUDA</title>
<location>https://nvidia.github.io/numba-cuda/user/index.html</location>
<excerpt>citeturn0search2 [wordlim: 200] Crawled: 3 months ago; * CUDA Array Interface (Version 3) * Python Interface Specification ... * Lifetime management in Numba
</excerpt>
</source>
<omitted_sources count="28">Additional source records were omitted from this bounded result.</omitted_sources>
</source_evidence>
Citations:
- 1: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 2: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 3: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
- 4: https://nvidia.github.io/numba-cuda/user/cuda_array_interface.html
🌐 Web query:
official CCCL DeviceReduce stream asynchronous API documentation
💡 Result:
<source_evidence>
<source kind="search">
<title>Streams — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/stream.html</location>
<excerpt>citeturn0search0 [wordlim: 200] Crawled: 3 weeks ago; It is passed as an argument to all asynchronous operations like kernel launch, memory copy and allocations. ... Availability: CCCL 2.2.0 / CUDA 12.3 ... // Pass to a stream-ordered API
# Streams#
Stream is conceptually a queue of operations for a specific device. It is passed as an argument to all asynchronous operations like kernel launch, memory copy and allocations.
## `cuda::stream_ref`#
`cuda::stream_ref` is a non-owning wrapper around a `cudaStream_t`. It prevents unsafe implicit constructions from `nullptr` or integer literals and provides convenient helpers for:
* `sync()`: wait for the recorded work to complete
* `is_done()`: non-blocking completion query
* comparison operators against other `cuda::stream_ref` or `cudaStream_t`
Availability: CCCL 2.2.0 / CUDA 12.3
Example:
#include <cuda/stream>
cudaStream_t stream;
cudaStreamCreate(&stream);
cuda::stream_ref ref{stream};
ref.sync(); // synchronizes the stream via cudaStreamSynchronize
assert(ref.is_done()); // verifies that the stream has finished all operations via cudaStreamQuery
// compare against other stream_ref or cudaStream_t
assert(ref == stream);
assert(ref != cuda::invalid_stream);
cudaStreamDestroy(stream);
## `cuda::stream`#
`cuda::stream` is an owning wrapper around a `cudaStream_t` that manages the lifetime of the underlying CUDA stream. It derives from `cuda::stream_ref`, provides all of its functionality, and can be used anywhere a `cuda::stream_ref` is expected. It can be constructed for a specific `cuda::device_ref`, moved (but not copied), and converted from or to a `cudaStream_t` via `from_native_handle`/`release()`. Constructing a new `cuda::stream` always creates a non-blocking stream; see non-blocking stream creation for CUDA Runtime interop details.
Availability: CCCL 3.1.0 / CUDA 13.1
#include <cuda/stream>
#include <cuda/devices>
int main() {
{
// Create a stream on a specific device
cuda::stream s{cuda::devices[0]};
// Pass to a stream-ordered API
// Synchronize the stream
s.sync();
} // Stream is automatically destroyed here
}</excerpt>
</source>
<source kind="search">
<title>cub::DeviceReduce — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html</location>
<excerpt>citeturn0search2 [wordlim: 200] Crawled: 4 days ago; > __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::ReducePolicy ... > const auto error = cub::DeviceReduce::Reduce( ... The problem size is read in stream order by the reduction kernels. ... See Two-Phase API (explicit temporary storage management) for usage guidance.
# cub::DeviceReduce#
struct DeviceReduce#
DeviceReduce provides device-wide, parallel operations for computing a reduction across a sequence of data items residing within device-accessible memory.
## Overview#
A reduction (or fold) uses a binary combining operator to compute a single aggregate from a sequence of input elements.
## Usage Considerations#
> * Dynamic parallelism. DeviceReduce methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported.
>
> * Determinism. The default reproducibility guarantee is `run_to_run`. A different guarantee can be requested through the execution environment with `cuda::execution::require`. See Determinism in CUB for the supported guarantees.
## Performance#
The work-complexity of reduction, reduce-by-key, and run-length encode as a function of input size is linear, resulting in performance throughput that plateaus with problem sizes large enough to saturate the GPU.
## Deferred problem sizes#
`Reduce`, `Sum`, `Min`, `Max`, and `TransformReduce` allow specifying the problem size from a value that resides in device memory through a single-value `cuda::args::deferred` argument. The deferred source may be a device pointer, a one-element span, or a random-access fancy iterator whose element is a non-`bool` 32- or 64-bit integer.
The problem size is read in stream order by the reduction kernels. Work that produces the count in the same stream is ordered automatically; a producer in another stream requires an event or an equivalent dependency. The source and its value must remain accessible and unchanged until all reduction kernels complete. The value must be nonnegative and `[d_in, d_in + num_items)` must be accessible. A temporary-storage query does not dereference the deferred source.
Deferred reductions are CUDA Graph capturable. The pointed-to count may change between graph replays without updating or recapturing the graph. Compile-time and runtime bounds are accepted as caller preconditions, but do not currently change temporary storage, grid dimensions, or pass selection. When those conditions are not met, the call falls back to `run_to_run` rather than failing.
Public Static Functions
template<typename InputIteratorT, typename OutputIteratorT, typename ReductionOpT, typename T, typename NumItemsT>
static inline cudaError_t Reduce(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
NumItemsT num_items,
ReductionOpT reduction_op,
T init,
cudaStream_t stream = nullptr
)#
Computes a device-wide reduction using the specified binary `reduction_op` functor and initial value `init`.
Added in version 2.2.0: First appears in CUDA Toolkit 12.3.
* Does not support binary reduction operators that are non-commutative.
* Provides “run-to-run” determinism for pseudo-associative reduction (e.g., addition of floating point types) on the same GPU device. However, results for pseudo-associative reduction may be inconsistent from one device to a another device of a different compute-capability because CUB can employ different tile-sizing for different architectures.
* The range `[d_in, d_in + num_items)` shall not overlap `d_out`.
* Temporary storage for this operation. If `d_temp_storage` is `nullptr`, the required size is written to `temp_storage_bytes` without dereferencing iterators or launching kernels. Otherwise, `d_temp_storage` must point to a device-accessible allocation of at least `temp_storage_bytes` bytes. No special alignment is requi…[truncated]</excerpt>
</source>
<source kind="search">
<title>Lower-level API — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cudax/stf/lower_level_api.html</location>
<excerpt>citeturn0search1 [wordlim: 200] Crawled: 3 weeks ago; # Lower-level API# ... This object provides access to the local description of the data associated with the task and a CUDA stream that can be used to submit work asynchronously.
}
auto lX = ctx.logical_data(X);
auto lY = ctx.logical_data(Y);
double alpha = 3.14;
/* Compute Y = Y + alpha X */
auto t = ctx.task(lX.read(), lY.rw());
t.start();
slice<double> sX = t.get<0>();
slice<double> sY = t.get<1>();
axpy<<<16, 128, 0, t.get_stream()>>>(sX.size(), alpha, sX.data_handle(), sY.data_handle());
t.end();
ctx.sync();
}
The `ctx.task()` call returns a task object. This object provides access to the local description of the data associated with the task and a CUDA stream that can be used to submit work asynchronously. The beginning of the task body and its end are delimited by the `.start()` and `.end()` calls. Failing to call either of these methods or calling them more than once or in the wrong order results in undefined behavior.
Asynchrony is achieved by using the CUDA stream, which provides a mechanism to submit work on the execution place (here, implicitly the current CUDA device). </excerpt>
</source>
<source kind="search">
<title>CUDA Runtime interactions — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/cudart_interactions.html</location>
<excerpt>citeturn0search3 [wordlim: 200] Crawled: last week; // Use default_stream with cccl-runtime APIs. ... In the CUDA Runtime API, the blocking/non-blocking stream creation flag controls synchronization with the CUDA default (NULL) stream.
Instead, it can be wrapped into `cuda::stream_ref` when needed for interop.
Note
When wrapping the NULL stream, the current device must be set explicitly first. CUDA binds the NULL stream to the active device, so the wrapper must be created after selecting the correct device.
## Example: wrapping the default stream#
#include <cuda/stream>
void use_default_stream(int device_id) {
cudaSetDevice(device_id);
cuda::stream_ref default_stream{cudaStreamPerThread};
// Use default_stream with cccl-runtime APIs.
}
The above applies to Driver API interop cases as well, where the current context must be managed by the user rather than the current device setting.
## Non-blocking stream creation#
Constructing a new `cuda::stream` always creates a stream with CUDA Runtime non-blocking behavior. This is the behavior of CCCL Runtime-created streams; wrapping or taking ownership of an existing `cudaStream_t` preserves the behavior of that handle.
In the CUDA Runtime API, the blocking/non-blocking stream creation flag controls synchronization with the CUDA default (NULL) stream. Because CCCL Runtime treats the default stream as an interop case rather than a first-class object, as described in default stream interop, `cuda::stream` does not expose a blocking/non-blocking construction option.
New Runtime code should express ordering between explicit streams directly, for example by making one `cuda::stream_ref` wait on another. Code that needs legacy CUDA Runtime implicit stream semantics should wrap the relevant CUDA Runtime stream handle in `cuda::stream_ref` (or take ownership with `cuda::stream::from_native_handle`); operations submitted through the wrapper use the same native handle and preserve that handle’s CUDA Runtime semantics, including any default-stream synchronization semantics.</excerpt>
</source>
<source kind="search">
<title>CUDASTF — CUDA Core Compute Libraries</title>
<location>https://nvidia.github.io/cccl/unstable/cudax/stf.html</location>
<excerpt>citeturn0search4 [wordlim: 200] Crawled: yesterday; This illustrates how a task in CUDASTF encapsulates a piece of work that is asynchronous with respect to CUDA stream semantics. ... ### Lower-level task API#
* `ctx.submit()` initiates the submission of all asynchronous tasks within the sequence
* `ctx.finalize()` awaits the conclusion of all outstanding asynchronous operations in the context, automatically invoking `ctx.submit()` if not previously called by user code
Usually, creating the task and invoking `ctx.finalize()` is sufficient. However, manually calling `ctx.submit()` can be beneficial in at least two situations. First, it allows for executing additional unrelated work on the CPU (or another GPU) between submission and synchronization. Second, when it’s necessary for two contexts to run concurrently, using the sequence `ctx1.submit(); ctx2.submit(); ctx1.finalize(); ctx2.finalize();` achieves this goal (whereas calling `ctx1.finalize(); ctx2.finalize();` without the `submit()` calls would wait for the completion of the first task before starting the second).
To wait for the completion of all pending operations (tasks, transfers, …), an asynchronous fence mechanism is available :
cudaStream_t stream = ctx.fence();
cudaStreamSynchronize(stream);
Another synchronization mechanism is the `wait` method of the context object. It is typically used in combination with the `reduce()` access mode for dynamic control flow. `auto val = ctx.wait(ld)` is a blocking call that returns the content of the `ld` logical data. The type of the returned value is defined by the `owning_container_of<interface>` trait class where `interface` is the data interface of the logical data. For a logical data with a regular data interface, the `wait` method therefore cannot be called unless that interface overloads this trait class.
This mechanism is illustrated in the dot product example of the Reduce access mode section.
`ctx.wait` can also be called on a token. Since a token carries no content, `ctx.wait(token)` returns `void` and simply blocks the host until the work the token depends on has completed (see the token section for an example).
</excerpt>
</source>
<omitted_sources count="11">Additional source records were omitted from this bounded result.</omitted_sources>
</source_evidence>
Citations:
- 1: https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html?utm_source=openai
- 2: https://nvidia.github.io/cccl/unstable/cub/api/structcub_1_1DeviceReduce.html?utm_source=openai
- 3: https://nvidia.github.io/cccl/unstable/libcudacxx/runtime/stream.html?utm_source=openai
important: Correct the input lifetime contract.
set_cccl_iterator_state binds only a raw pointer, and the CCCL call can leave work pending on the supplied stream. The CUDA Array Interface does not extend the producer's lifetime. Therefore, keeping input_it alive only until the Python call returns can allow its allocation to be released while queued kernels still use it. Require callers to keep the input and its allocation alive until the submitted work completes, or implement per-call stream-ordered retention. Do not retain it in the process-wide cached iterator.
_compile_op_cache is an unbounded dict. Each distinct operator, input type tuple, output type, and target compute capability adds a completed entry. _cache_single_flight removes an entry only when its builder fails, so completed entries are never evicted. The previous _compile_op_impl was decorated with functools.lru_cache(maxsize=256), so this change removes that bound. A long-running sweep that compiles many distinct operators keeps every compiled code blob until compile_op.cache_clear() runs. This can grow process memory until it runs out.
Add bounded eviction for completed entries and keep same-key single-flight. Evict only completed values. An _InFlightBuild entry must stay in the memo until its builder finishes, because waiting callers depend on its event. Apply the eviction in compile_op's memo or in a wrapper around _cache_single_flight. Do not change _cache_single_flight for all callers, because the build-results caches also use it.
The reason will be displayed to describe this comment to others. Learn more.
@oleksandr-pavlyk suggestion: We can defer the bounded cache to a follow-up. clear_all_caches() is a temporary mitigation, but callers must synchronize it with active builds and algorithm execution.
Would you like me to open a GitHub issue for the bounded LRU and single-flight design?
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
important: Include captured arrays in the lifetime rule. The phrase “arrays you pass” does not cover an array captured by an operator, as in test_caching.py. The stateful operator packs that array’s device pointer into runtime state. If the caller releases the array before asynchronous work completes, GPU work can use a freed allocation. Tell callers to retain captured arrays and arrays behind iterators until stream work completes. (raw.githubusercontent.com) As per path instructions, documentation changes must “focus on technical accuracy.”
I feel like retaining a pointer in the cached algorithm without retaining reference to the underlying allocation opens us up for illegal memory access issues. What about normalizing keys by replacing actual owning arrays with ProxyArray?
@oleksandr-pavlyk done. For sorting algorithms, it is a little more complicated because they assume that the values pointer being null (which is true for ProxyArray) means keys only. Had to workaround this with a placeholder address. Alternatively, we can add an absent option to specify this
Reviewing files that changed from the base of the PR and between 9be4474 and 508da50.
Files selected for processing: 2
python/cuda_cccl/cuda/compute/_caching.py
python/cuda_cccl/tests/compute/test_caching.py
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
The reason will be displayed to describe this comment to others. Learn more.
A module-level operator such as def add(x, y): return np.add(x, y) now raises RecursionError while constructing its cache key.
co_names includes the attribute name add; the existing globals lookup resolves that name to the operator itself, and the new function branch recursively hashes it forever.
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
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
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.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Description
closes #12187
Checklist