Dry Run Protocol

[achirkin]

The dry run protocol defines a mechanism to simulate the execution of algorithms to get a precise estimate of the memory requirements for a real execution with the same parameters.

#include <raft/util/dry_run_memory_resource.hpp>

raft::resources res;
// auto my_function(const raft::resources& res, my_args...);
auto stats = raft::util::dry_run_execute(res, my_function, my_args...);
// stats.device_global  – peak device memory (bytes)

This PR:

• Introduces new infrastructure: raft::util::dry_run_execute, raft::dry_run_resources, and resource::get_dry_run_flag to let callers estimate peak memory usage of any RAFT algorithm without executing GPU work.
• Makes all public functions across all raft namespaces dry-run compliant: allocations are always visible to the tracker; CUDA work is skipped.
• Adds a small user guide (docs/source/dry_run_protocol.md)

Note for reviewers
The PR contains a lot of small tedious changes to cover all of raft library and the tests components.
Please start reading at docs/source updates to learn more about the topic and the principles guiding these changes.

[huuanhhuyn]

Thanks @achirkin for the heavy-lifting! I have went through this PR halfway.

[huuanhhuyn]

What is the motivation for guarding sync_stream?

Your dry_run_protocol.md suggests the opposite that sync_stream is safe without guard.

[achirkin]

The guide says it's safe to call this function, because it is dry-run compliant (takes the raft::resources as an argument) - hence it guards from any CUDA work inside.
cudaStreamSynchronize is cuda CUDA and it takes time, so we must guard it.

[huuanhhuyn]

If I understand dry-run correctly, when res is in dry-run mode, the stream contains no CUDA work and contains only light-weight fake memory operations. Therefore it would be fast to sync.

If user intentionally adds a CUDA/memory work to the stream without guarding, he wants the sync_stream to work I suppose.

I mean this guarding on sync stream wouldn't bring any benefit but limit its usage a little.

[achirkin]

It is still CUDA work and a call to cuda context, which may briefly lock. In dry run mode, nobody should ever call CUDA, that is a hard constraint.

[huuanhhuyn]

if all CUDA work is guarded / skipped in dry-run. How do we measure allocations whose allocation sizes are produced by a kernel?

If that kernel is dependent on earlier kernels and so on, dry-run must then execute all of them without guarding?

[achirkin]

The case you're describing falls under "allocation size depends on values/data". This is indeed generally a problem and we resort to providing safe upper bounds on allocations. Have a look at the sparse namespace - we have a lot cases like this.
In tests, this is covered by DATA_DRIVEN provenance: we compare that dry run produced stats are not smaller than the real produced stats.

[huuanhhuyn]

Is finding the upper bound always possible? And what if the upper bound is 3x, 10x?

[huuanhhuyn]

Although your implementation avoids peak data (old + new) when the new size is larger than the current size, it has two issues:

• when the new size is the same -> it does NOT reallocate while its name is "reallocate"
• when the new size is smaller -> it DOES allocate while it could more efficiently shrink the current buffer (similar to what resize() already does)

I would suggest to do the current implementation only when the new size is larger. When the new size is less or equal, simply call resize() and clear the buffer.

I am not sure about clearing the buffer. It adds some overhead, but it sounds safer to me.

[achirkin]

Funny enough, both of these are pretty much intentional features. Here's a bit more context for this.

By design, mdarrays and mdspans are not supposed to be resizeable at all. However, at some point we decided to implement sparse matrices using the same container policies as mdarrays. For sparse data, one needs the resize functionality to allow changing sparsity pattern. So the resize function is exclusively used for sparse data and only to allocate a new sparse matrix storage of the required size (initialize_sparsity).

Back to reallocate addressing your concerns. The intention here is to always free up the memory as soon as requested: the released buffer may be huge (e.g. half of GPU memory); the user may want to use the memory as soon as possible, so we want it to be available.

Essentially reallocate is almost the same as running resize + shrink_to_fit on an underlying RMM container. But there's a catch: reallocate guarantees to NOT copy the data. This catch is important for dry run: we fake the memory allocations, so any memory copy would cause OOM / break the program.

[huuanhhuyn]

Thank you for sharing the context with sparse matrix.

I still don't understand why when the new size is less than or equal to the current size, we don't simply free the redundant memory and keep the remaining new size part. This will give back immediately current_size - new_size to the user. Or do you mean we have to free everything at the current address and then allocate the new size some where else? If so, why?

[huuanhhuyn]

The current implementation makes sense to me only when the new size is greater than the current size

[huuanhhuyn]

I suppose we are not able to track non-raft allocations with dry-run. It is probably fine but we should be aware and document it somewhere.

Are we using third-party libraries which do allocations without RAFT somewhere? I am not aware about it.

[achirkin]

Plenty of rapids org libraries use the rmm memory resources directly. They are tracked though.

[huuanhhuyn]

And non-rapidsai libraries like Cutlass, CUB, Thrust, etc.? Are their allocations tracked?

[achirkin]

This is covered in the same file below in the section "What Needs Guarding". In general, I don't think it's sustainable to maintain a list of all possible libraries to tell whether they use rmm or not. But it's normally clear from the API or documentation of those libraries.

Please also refer to the presentation for the relevant overview https://drive.google.com/file/d/1uMkM-Xzhi-2ZAKe_z3zvCfHoefn0bAAx/view?usp=drive_link

[divyegala]

Would you consider guarding the kernel launch behind a dispatch function/macro where the dry run flag is checked, like we talked in your offline brown bag session?

[achirkin]

Hi @divyegala , I definitely want to do that, but maybe in a separate PR in parallel, where we could focus on discussing the pros and cons of different approaches to do that? The main issue I see is that for being useful for dry run, the dispatcher should take raft::resources as an argument, which is not always available in legacy code in detail namespace.

[divyegala]

@achirkin thank you, that makes sense to me. I will start reviewing the PR in its current form. 👍

[coderabbitai]

Actionable comments posted: 1

Note

Due to the large number of review comments, Critical severity comments were prioritized as inline comments.

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (2)

cpp/include/raft/matrix/detail/select_radix.cuh (1)

1280-1337: 🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift

Do not query CUDA occupancy during dry-run.

Lines 1302-1337 dispatch to radix_topk_one_block or radix_topk. Both functions call calc_chunk_size before their if (dry_run) return. calc_chunk_size calls cudaOccupancyMaxActiveBlocksPerMultiprocessor. The multi-block path can also call calc_grid_dim, which makes the same CUDA Runtime call.

A dry-run invocation therefore enters CUDA Runtime before it returns. Guard the occupancy queries. Use a dry-run workspace sizing path that does not call CUDA Runtime APIs.

#!/bin/bash
set -euo pipefail

file="cpp/include/raft/matrix/detail/select_radix.cuh"

# Inspect the dry-run return placement relative to occupancy queries.
sed -n '806,837p;880,920p;1153,1190p;1275,1340p' "$file"

# Confirm all occupancy-query call sites and dry-run guards.
rg -n -C 3 'cudaOccupancyMaxActiveBlocksPerMultiprocessor|calc_chunk_size|calc_grid_dim|if \(dry_run\)' "$file"

As per path instructions, “functions using raw streams or CUDA work must guard execution with resource::get_dry_run_flag(res).”

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/matrix/detail/select_radix.cuh` around lines 1280 - 1337,
Ensure dry-run execution in the radix top-k dispatch does not invoke CUDA
Runtime occupancy APIs. Update calc_chunk_size and calc_grid_dim, and the
radix_topk_one_block/radix_topk paths that call them, to use a deterministic
workspace-sizing path when dry_run is true and only query CUDA occupancy for
real execution; preserve normal dispatch and sizing behavior otherwise.

cpp/include/raft/linalg/power.cuh (1)

79-95: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

Move the dry-run guard below the input validation.

In both power and power_scalar, the guard returns before the RAFT_EXPECTS checks. In dry-run mode, contiguity and size mismatches are therefore not reported, and the caller receives a memory estimate for an input that will fail in the real run.

Neither function allocates memory, so moving the guard after the validation block does not change allocation tracking. It only skips the kernel launch, which is the intent. cpp/include/raft/stats/stddev.cuh and cpp/include/raft/stats/dispersion.cuh in this PR already keep their RAFT_EXPECTS checks ahead of the dry-run dispatch.

🛡️ Proposed fix for both overloads

 void power(raft::resources const& handle, InType in1, InType in2, OutType out)
 {
-  if (resource::get_dry_run_flag(handle)) { return; }
   using in_value_t  = typename InType::value_type;
   using out_value_t = typename OutType::value_type;
 
   RAFT_EXPECTS(raft::is_row_or_column_major(out), "Output must be contiguous");
   RAFT_EXPECTS(raft::is_row_or_column_major(in1), "Input 1 must be contiguous");
   RAFT_EXPECTS(raft::is_row_or_column_major(in2), "Input 2 must be contiguous");
   RAFT_EXPECTS(out.size() == in1.size() && in1.size() == in2.size(),
                "Size mismatch between Output and Inputs");
 
+  if (resource::get_dry_run_flag(handle)) { return; }
   power<in_value_t, out_value_t, typename OutType::index_type>(

 {
-  if (resource::get_dry_run_flag(handle)) { return; }
   using in_value_t  = typename InType::value_type;
   using out_value_t = typename OutType::value_type;
 
   RAFT_EXPECTS(raft::is_row_or_column_major(out), "Output must be contiguous");
   RAFT_EXPECTS(raft::is_row_or_column_major(in), "Input must be contiguous");
   RAFT_EXPECTS(out.size() == in.size(), "Size mismatch between Output and Input");
 
+  if (resource::get_dry_run_flag(handle)) { return; }
   powerScalar<in_value_t, out_value_t, typename OutType::index_type>(

One caveat: power_scalar dereferences *scalar.data_handle() after the guard. That host read is safe only if the scalar is a real host allocation. Confirm that make_host_scalar in dry-run mode still yields a readable pointer, because cpp/include/raft/core/host_mdarray.hpp skips value initialization in dry-run mode.

Also applies to: 118-132

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/linalg/power.cuh` around lines 79 - 95, Move the dry-run
guard below all RAFT_EXPECTS validation in both power and power_scalar, so
contiguity and size checks always execute before skipping the kernel launch.
Preserve the existing validation and dry-run behavior, and verify that
power_scalar’s dereference of scalar.data_handle() remains safe with
make_host_scalar in dry-run mode; adjust only if that pointer is not readable.

🟠 Major comments (20)

cpp/include/raft/sparse/op/detail/reduce.cuh-137-149 (1)

137-149: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Reject scan sizes that do not fit CUB's item count.

Lines 137 and 149 narrow diff.size() to int. If nnz exceeds INT_MAX, the CUB scan receives a truncated or invalid item count. The later read from diff.back() then does not represent a scan over all nonzeros. Validate the bound before the workspace query, or use a supported CUB interface with a wider item-count type.

As per coding guidelines, add input validation when invalid dimensions can cause incorrect behavior.

#!/bin/bash
set -euo pipefail

rg -n -C 4 'max_duplicates|static_cast<int>\(diff.size\(\)\)|nnz_t' \
  cpp/include/raft/sparse/op/detail/reduce.cuh cpp/include/raft/sparse

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/op/detail/reduce.cuh` around lines 137 - 149,
Validate before both CUB scan calls that diff.size() is no greater than the
maximum representable int, and reject the input with the established validation
mechanism when it exceeds that bound. Update the workspace query and scan in the
surrounding reduction flow to use only a validated item count, preserving full
coverage of all nonzeros and preventing narrowing overflow.

cpp/include/raft/sparse/op/detail/reduce.cuh-136-149 (1)

136-149: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Check both CUB scan calls.

Lines 136-137 and 148-149 ignore the CUDA status from cub::DeviceScan::ExclusiveSum. A failed workspace query can leave scan_ws_bytes invalid. A failed scan can continue to the host copy and use invalid scan data. Wrap both calls with RAFT_CUDA_TRY.

Proposed fix

-  cub::DeviceScan::ExclusiveSum(
-    nullptr, scan_ws_bytes, diff.data(), diff.data(), static_cast<int>(diff.size()), stream);
+  RAFT_CUDA_TRY(cub::DeviceScan::ExclusiveSum(
+    nullptr, scan_ws_bytes, diff.data(), diff.data(), static_cast<int>(diff.size()), stream));
...
-  cub::DeviceScan::ExclusiveSum(
-    scan_ws.data(), scan_ws_bytes, diff.data(), diff.data(), static_cast<int>(diff.size()), stream);
+  RAFT_CUDA_TRY(cub::DeviceScan::ExclusiveSum(
+    scan_ws.data(), scan_ws_bytes, diff.data(), diff.data(), static_cast<int>(diff.size()), stream));

As per coding guidelines, “kernel launches, memory operations, and synchronization must use RAFT/CUDA error checking.” As per path instructions, use checked CUDA macros.

#!/bin/bash
set -euo pipefail

rg -n -C 3 'DeviceScan::ExclusiveSum|RAFT_CUDA_TRY' \
  cpp/include/raft/sparse/op/detail/reduce.cuh

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/op/detail/reduce.cuh` around lines 136 - 149, Wrap
both cub::DeviceScan::ExclusiveSum calls in RAFT_CUDA_TRY, including the
workspace-size query and the actual scan, so CUDA failures are propagated before
scan_ws allocation or subsequent data use.

cpp/include/raft/linalg/detail/lstsq.cuh-213-216 (1)

213-216: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Use exception-safe cleanup for gesvdj_params.

The dry-run branch destroys gesvdj_params, but the normal path does not destroy it. A later allocation failure or checked cuSOLVER failure also bypasses this cleanup. Use an RAII owner or a scope guard immediately after cusolverDnCreateGesvdjInfo succeeds.

As per path instructions, review CUDA call sites for “leaks or exception-path cleanup failures.”

#!/bin/bash
set -euo pipefail

rg -n -C 8 \
  'lstsqSvdJacobi|cusolverDnCreateGesvdjInfo|cusolverDnDestroyGesvdjInfo' \
  cpp/include/raft/linalg/detail/lstsq.cuh

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/linalg/detail/lstsq.cuh` around lines 213 - 216, Make
cleanup of gesvdj_params exception-safe in the code surrounding
cusolverDnCreateGesvdjInfo: establish an RAII owner or scope guard immediately
after successful creation that calls cusolverDnDestroyGesvdjInfo exactly once on
every exit path, including dry-run returns, allocation failures, and checked
cuSOLVER errors. Remove the manual dry-run-only destruction in the
lstsqSvdJacobi flow to avoid double cleanup.

cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh-123-126 (1)

123-126: 🗄️ Data Integrity & Integration | 🟠 Major | ⚡ Quick win

Check both CUB scan calls.

cub::DeviceScan::ExclusiveSum returns a CUDA status. If either call fails, the code can continue with invalid workspace or CSR row-offset data. Wrap both calls with RAFT_CUDA_TRY.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh` around lines 123 -
126, Wrap both cub::DeviceScan::ExclusiveSum calls in RAFT_CUDA_TRY, including
the workspace-size query and the actual scan execution. Ensure any CUDA failure
is propagated before using scan_ws or the generated CSR row offsets.

Sources: Coding guidelines, Path instructions

cpp/include/raft/stats/detail/scores.cuh-183-183 (1)

183-183: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Move host-vector allocations after the dry-run return.

std::vector<double> h_sorted_abs_diffs(n) allocates host memory before this check. A large dry-run request can consume or exhaust host memory without contributing to RAFT allocation statistics. Keep the RMM allocations before this check, but construct mean_errors and h_sorted_abs_diffs only after dry-run mode returns.

Proposed fix

-  std::vector<double> mean_errors(2);
-  std::vector<double> h_sorted_abs_diffs(n);
   int thread_cnt = 256;
   int block_cnt  = raft::ceildiv(n, thread_cnt);

   if (dry_run) { return; }

+  std::vector<double> mean_errors(2);
+  std::vector<double> h_sorted_abs_diffs(n);
   RAFT_CUDA_TRY(cudaMemsetAsync(tmp_sums.data(), 0, 2 * sizeof(double), stream));

As per path instructions, dry-run implementations must suppress real allocations while leaving RMM/RAFT allocation attempts unguarded.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/stats/detail/scores.cuh` at line 183, Move the dry_run
return in the surrounding score computation before constructing the host vectors
mean_errors and h_sorted_abs_diffs, while keeping the existing RMM/RAFT
allocation attempts before that return. Ensure dry-run mode performs no
host-vector allocations and exits immediately after those allocation attempts.

Source: Path instructions

cpp/tests/sparse/reduce.cu-65-82 (1)

65-82: 🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift

The COO output out is allocated inside the dry-run scope but declared outside it.

out is constructed at Line 65, outside the wrapper. During the dry-run pass, max_duplicates reaches the dry-run branch in cpp/include/raft/sparse/op/detail/reduce.cuh and calls out.allocate(nnz, m, n, false, stream). That allocation uses the temporary dry-run memory resource installed by execute_with_dry_run_check. When the wrapper exits, out still owns those probe-backed buffers, and they are later freed or reallocated against the restored resource.

This is the same hazard that cpp/tests/core/bitset.cu documents at Lines 201-203, where the whole object lifetime is kept inside the dry-run scope. Declare out inside the callback, or confirm that the dry-run resource outlives every object it allocated.

The inline comment at Lines 66-67 states that the COO output is not tracked. That statement does not match the implementation, which allocates the output during the dry-run pass.

#!/bin/bash
# Check the lifetime semantics of the dry-run resource wrapper and the COO allocate path.
set -euo pipefail

fd -t f 'dry_run_resources.hpp' cpp/include | while IFS= read -r f; do
  echo "=== $f ==="
  cat -n "$f"
done

echo "=== execute_with_dry_run_check definition ==="
rg -n -C 25 'execute_with_dry_run_check' cpp/include cpp/tests --glob '*.{hpp,cuh,h}'

echo "=== COO::allocate ==="
rg -n -C 15 'void allocate' cpp/include/raft/sparse/coo.hpp

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/sparse/reduce.cu` around lines 65 - 82, Move construction of the
COO variable out of the surrounding scope and into the dry-run callback
containing max_duplicates, keeping its entire lifetime within
execute_with_dry_run_check. Remove or update the inaccurate comments claiming
the COO output is not tracked, and ensure the callback’s output remains
available to the subsequent test logic as required.

cpp/tests/core/bitset.cu-388-399 (1)

388-399: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Gate the any and none assertions on the dry-run flag.

Lines 390-391 assert values without checking resource::get_dry_run_flag(h), but every other value-dependent assertion in this callback is gated. In the dry-run pass, my_bitset.reset(h, false) performs no CUDA work, so any(h) and none(h) read probe-backed memory whose contents are undefined. These two assertions can fail during the dry-run pass.

The assertions at Lines 395-399 already cover the same two calls correctly.

💚 Proposed fix

         // Test count() operations
         my_bitset.reset(h, false);
-        ASSERT_EQ(my_bitset.any(h), false);
-        ASSERT_EQ(my_bitset.none(h), true);
+        if (!resource::get_dry_run_flag(h)) {
+          ASSERT_EQ(my_bitset.any(h), false);
+          ASSERT_EQ(my_bitset.none(h), true);
+        }
         raft::linalg::map_offset(h, query_device.view(), raft::cast_op<index_t>{});

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/core/bitset.cu` around lines 388 - 399, Gate the initial
my_bitset.any(h) and my_bitset.none(h) assertions after reset on
resource::get_dry_run_flag(h), matching the existing guarded value-dependent
assertions below. Keep the reset and subsequent setup unchanged, and preserve
the assertions during non-dry-run execution.

cpp/include/raft/solver/detail/lap_functions.cuh-269-271 (1)

269-271: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Calculate dry-run allocation counts in size_t.

Both sites evaluate SP * N before device_uvector receives the value. If the product exceeds the source integer type, overflow produces an invalid allocation size and an incorrect peak-memory estimate.

• cpp/include/raft/solver/detail/lap_functions.cuh#L269-L271: derive a checked size_t matrix_size before allocating predicates_v, addresses_v, and csr_neighbors_v.
• cpp/include/raft/solver/detail/lap_functions.cuh#L432-L436: derive the same checked size_t matrix_size before allocating predicates_v, addresses_v, and elements_v.

#!/bin/bash
set -euo pipefail
rg -n -C 3 'device_uvector<.*>\([^;]*SP \* N' cpp/include/raft/solver/detail/lap_functions.cuh

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/solver/detail/lap_functions.cuh` around lines 269 - 271, In
cpp/include/raft/solver/detail/lap_functions.cuh at lines 269-271 and 432-436,
compute a checked size_t matrix_size from SP and N before constructing the
device_uvector instances; use that variable for predicates_v, addresses_v,
csr_neighbors_v, and elements_v at the respective sites, preserving the existing
allocation behavior while preventing source-type overflow.

cpp/include/raft/stats/detail/contingencyMatrix.cuh-209-219 (1)

209-219: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Prevent overflow in the dry-run workspace estimate.

Line 217 evaluates 4 * nSamples as int. If nSamples > INT_MAX / 4, signed overflow occurs before conversion to size_t. The dry-run estimate can then under-report or corrupt the required workspace size.

Validate nSamples and promote it before every size calculation.

Proposed fix

+  RAFT_EXPECTS(nSamples >= 0, "nSamples must be non-negative");
+  auto sample_count = static_cast<size_t>(nSamples);
   if (dry_run) {
-    auto tmpStagingMemorySize = raft::alignTo<size_t>(nSamples * sizeof(T), 256);
+    auto tmpStagingMemorySize = raft::alignTo<size_t>(sample_count * sizeof(T), 256);
     tmpStagingMemorySize *= 2;
-    size_t cubWorkspaceUpperBound = 4 * nSamples * sizeof(T);
+    size_t cubWorkspaceUpperBound = size_t{4} * sample_count * sizeof(T);

#!/bin/bash
set -euo pipefail
python3 - <<'PY'
import ctypes
max_int = (1 << (ctypes.sizeof(ctypes.c_int) * 8 - 1)) - 1
print(f"INT_MAX={max_int}")
print(f"First overflowing nSamples for 4 * nSamples: {max_int // 4 + 1}")
PY

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/stats/detail/contingencyMatrix.cuh` around lines 209 - 219,
Prevent signed overflow in the dry-run calculations within the
contingency-matrix workspace sizing block by validating nSamples against the
supported range and promoting it to size_t before every multiplication,
including the tmpStagingMemorySize and cubWorkspaceUpperBound expressions.
Preserve the existing alignment and workspace estimate behavior for valid sample
counts, and reject or otherwise handle values that cannot be represented safely.

cpp/include/raft/stats/detail/meanvar.cuh-212-223 (1)

212-223: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Guard the occupancy CUDA call during dry runs.

Line 212 guards the memset and kernel launches, but lines 203-207 still call cudaOccupancyMaxActiveBlocksPerMultiprocessor when dry_run is true. This makes a memory-only probe issue a CUDA runtime query. Move the occupancy calculation and gs.y adjustment into the !dry_run branch.

As per path instructions, dry-run implementations must suppress CUDA work while preserving allocation tracking.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/stats/detail/meanvar.cuh` around lines 212 - 223, Move the
cudaOccupancyMaxActiveBlocksPerMultiprocessor call and related gs.y adjustment
into the existing !dry_run branch around meanvar_kernel_rowmajor and
meanvar_kernel_fill. Ensure dry_run performs no CUDA runtime queries or kernel
work while retaining the existing buffer allocation and tracking behavior.

Source: Path instructions

cpp/include/raft/sparse/solver/detail/lanczos.cuh-158-169 (1)

158-169: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Check all five CUDA kernel launches with RAFT_CUDA_TRY(cudaPeekAtLastError()).

Add the check immediately after kernel_triangular_populate, kernel_triangular_beta_k, kernel_clamp_down, kernel_clamp_down_vector, and kernel_normalize.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/solver/detail/lanczos.cuh` around lines 158 - 169,
Add RAFT_CUDA_TRY(cudaPeekAtLastError()) immediately after each of the five CUDA
launches in lanczos.cuh: kernel_triangular_populate and kernel_triangular_beta_k
in lines 158-169, plus kernel_clamp_down, kernel_clamp_down_vector, and
kernel_normalize in lines 386-415. Keep the checks directly adjacent to their
respective launches.

Sources: Coding guidelines, Path instructions

cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh-309-315 (1)

309-315: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Allocate scan_ws from device_memory. The current constructor uses the default device resource, so custom workspace-resource accounting and isolation do not include the CUB scan workspace.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh` around lines 309 -
315, Construct scan_ws with device_memory, matching the resource used for
sub_nnz, so the CUB scan workspace is accounted for by the configured workspace
resource.

Sources: Coding guidelines, Path instructions

cpp/include/raft/sparse/solver/detail/lanczos.cuh-198-227 (1)

198-227: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Use the handle-bound Thrust policy.

thrust::device runs on the default CUDA stream, while neighboring operations use the handle stream. This can break ordering for non-default handle streams. Use auto thrust_policy = resource::get_thrust_policy(handle) for thrust::sequence and both thrust::sort calls.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/solver/detail/lanczos.cuh` around lines 198 - 227,
Use a handle-bound Thrust execution policy for all operations in this selection
block: define thrust_policy via resource::get_thrust_policy(handle), then pass
it to thrust::sequence and both thrust::sort calls instead of thrust::device.
Preserve the existing index ordering and eigenvalue selection logic.

Sources: Coding guidelines, Path instructions

cpp/include/raft/matrix/detail/select_k-inl.cuh-131-132 (1)

131-132: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Check each CUB sort status before the dry-run return.

SortPairs and SortPairsDescending return CUDA status. The dry-run path returns after an unchecked workspace query. If that query fails, dry-run can report memory usage instead of returning the CUDA error. Wrap both workspace queries and both execution calls with RAFT_CUDA_TRY.

#!/bin/bash
set -euo pipefail

sed -n '90,170p' cpp/include/raft/matrix/detail/select_k-inl.cuh
rg -n -C 2 'RAFT_CUDA_TRY' cpp/include/raft/core/detail/macros.hpp cpp/include/raft

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/matrix/detail/select_k-inl.cuh` around lines 131 - 132,
Update the CUB sort calls in the select-k implementation, including both
SortPairs and SortPairsDescending workspace-size queries and execution calls, to
use RAFT_CUDA_TRY. Ensure each status is checked before the
resource::get_dry_run_flag(handle) early return so failed dry-run workspace
queries propagate the CUDA error.

cpp/include/raft/stats/detail/silhouette_score.cuh-229-230 (1)

229-230: 🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift

Use one ordered stream for dependent operations.

pairwise_distance(handle, ...) and matrix_vector_op(handle, ...) use the stream in handle. The dependent operations use stream. If these streams differ, the reductions can read distanceMatrix or averageDistanceBetweenSampleAndCluster before the producer work completes.

Use resource::get_cuda_stream(handle) for all operations in this function, or add explicit event ordering between the two streams. As per path instructions, maintain asynchronous stream ordering and use RAFT-managed streams for internal work.

Also applies to: 268-276

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/stats/detail/silhouette_score.cuh` around lines 229 - 230,
Update the operations in the silhouette-score function, including
pairwise_distance and the dependent matrix_vector_op/reduction calls, to use the
same RAFT-managed stream obtained via resource::get_cuda_stream(handle).
Preserve asynchronous ordering and ensure every producer and consumer of
distanceMatrix and averageDistanceBetweenSampleAndCluster is submitted to that
stream.

Source: Path instructions

cpp/include/raft/matrix/threshold.cuh-38-38 (1)

38-38: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

Keep the dimension check outside the dry-run guard.

RAFT_EXPECTS(in.size() == out.size(), ...) is skipped when dry-run mode is enabled. Dry-run then accepts an input/output shape that normal execution rejects. This can produce a memory estimate for a call that cannot run.

Move the dry-run guard after the size check, while still skipping detail::setSmallValuesZero.

Proposed fix

-  if (resource::get_dry_run_flag(handle)) { return; }
   RAFT_EXPECTS(in.size() == out.size(), "Input and output matrices must have same size");
+  if (resource::get_dry_run_flag(handle)) { return; }

As per coding guidelines, validate invalid dimensions where they can cause incorrect behavior. As per path instructions, dry-run guards should suppress CUDA work without bypassing required precondition checks.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/matrix/threshold.cuh` at line 38, Move the
resource::get_dry_run_flag(handle) guard in the threshold operation after the
RAFT_EXPECTS input/output size validation, so dimension mismatches are rejected
in both modes while dry-run still skips detail::setSmallValuesZero and other
CUDA work.

Sources: Coding guidelines, Path instructions

cpp/include/raft/mr/dry_run_resource.hpp-181-207 (1)

181-207: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

The probe ignores the alignment of later allocations.

std::call_once creates the probe with the alignment of the first request only. A later request with a larger alignment receives the same pointer, which can be under-aligned for that request. Aligned vector loads and stores on that pointer are then invalid, and the returned pointer breaks the alignment contract of the memory-resource API.

The probe is also 256 bytes for every request size. Any code path that is not dry-run guarded and writes to the returned pointer overruns the probe.

Allocate the probe with the maximum alignment RAFT can request, or reallocate the probe when a request needs a stricter alignment.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/mr/dry_run_resource.hpp` around lines 181 - 207, Update
allocate_sync and allocate so the shared state_->probe always satisfies the
strictest alignment requested by RAFT, rather than preserving the first
allocation’s alignment; reallocate or otherwise replace the probe when a later
request requires greater alignment. Also size the probe to accommodate the
requested bytes, or otherwise ensure returned storage cannot be overrun by an
unguarded request, while preserving allocation/deallocation accounting.

Source: Coding guidelines

cpp/include/raft/core/host_container_policy.hpp-118-127 (1)

118-127: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win

Reset bytesize_ when you free the buffer.

Line 122 frees data_ and Line 123 sets data_ to nullptr, but bytesize_ keeps the old value. If the allocation on Line 125 throws, the container is left with data_ == nullptr and a non-zero bytesize_. Any later operator[] on that object dereferences a null pointer.

The destructor is safe because it tests both fields, so this is an invariant defect and not a double free.

🛠️ Proposed fix

     if (data_ != nullptr) {
       mr_.deallocate_sync(data_, bytesize_);
       data_ = nullptr;
+      bytesize_ = 0;
     }

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/core/host_container_policy.hpp` around lines 118 - 127,
Update host_container::reallocate so bytesize_ is reset to zero immediately
after deallocating data_ and clearing data_. Preserve the existing allocation
and swap flow, ensuring an allocation failure cannot leave a null data_ pointer
paired with a stale nonzero size.

Source: Coding guidelines

cpp/include/raft/random/detail/rng_impl.cuh-16-23 (1)

16-23: 📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win

Add the dry_run_flag.hpp include.

Line 396 calls resource::get_dry_run_flag(res) and Line 397 calls resource::get_cuda_stream(res), but this header does not include <raft/core/resource/dry_run_flag.hpp> or <raft/core/resource/cuda_stream.hpp>. The build then depends on transitive includes and can break when an upstream header changes.

🛠️ Proposed fix

 `#include` <raft/core/detail/macros.hpp>
 `#include` <raft/core/device_mdarray.hpp>
 `#include` <raft/core/math.hpp>
 `#include` <raft/core/operators.cuh>
+#include <raft/core/resource/cuda_stream.hpp>
+#include <raft/core/resource/dry_run_flag.hpp>
 `#include` <raft/linalg/map.cuh>

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/random/detail/rng_impl.cuh` around lines 16 - 23, Add direct
includes for the declarations used by the RNG implementation: include the
dry-run flag and CUDA stream resource headers before the code calling
resource::get_dry_run_flag and resource::get_cuda_stream. Keep the existing
scatter and CUB includes unchanged.

cpp/include/raft/spectral/detail/matrix_wrappers.hpp-393-396 (1)

393-396: 🚀 Performance & Scalability | 🟠 Major | ⚡ Quick win

auto handle = get_handle() copies the raft::resources object in both mv overrides. sparse_matrix_t::get_handle() returns resources const& (Line 323), so auto deduces raft::resources by value and each call constructs and destroys a full handle copy. Both overrides run once per solver iteration, so the copy is repeated on a hot path.

• cpp/include/raft/spectral/detail/matrix_wrappers.hpp#L393-L396: change auto handle to auto const& handle in laplacian_matrix_t::mv.
• cpp/include/raft/spectral/detail/matrix_wrappers.hpp#L459-L462: change auto handle to auto const& handle in modularity_matrix_t::mv.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/spectral/detail/matrix_wrappers.hpp` around lines 393 - 396,
Update the handle declarations in both laplacian_matrix_t::mv
(cpp/include/raft/spectral/detail/matrix_wrappers.hpp:393-396) and
modularity_matrix_t::mv
(cpp/include/raft/spectral/detail/matrix_wrappers.hpp:459-462) to bind the
resources returned by sparse_matrix_t::get_handle() as a const reference instead
of copying them; leave the cublas, stream, and dry-run retrieval unchanged.

🟡 Minor comments (1)

docs/source/developer_guide.md-325-329 (1)

325-329: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

Line 329 renders as part of the last nested bullet.

All attempted allocations in the above resources are tracked... immediately follows a nested list item with no blank line. CommonMark and MyST treat that as a lazy continuation of the list item, so the sentence renders inside the bullet instead of as a standalone paragraph. Insert a blank line.

📝 Proposed fix

   - workspace memory resources managed by `raft::resources`.
+
 All attempted allocations in the above resources are tracked and reported, thus enabling planning of the memory usage with a relatively small overhead of simulated execution.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@docs/source/developer_guide.md` around lines 325 - 329, Insert a blank line
after the final nested bullet under the listed memory resources and before “All
attempted allocations...” so that sentence renders as a standalone paragraph
rather than as continuation text within the bullet.

🧹 Nitpick comments (3)

cpp/tests/sparse/preprocess.cu (1)

166-180: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Align the TF-IDF dry-run expectation with the BM25 expectation.

encode_tfidf and encode_bm25 follow the same allocation pattern in this test, but the BM25 branches use alloc_behavior::DATA_DRIVEN with sizeof(float) * coo_a.nnz, while the TF-IDF branches use alloc_behavior::ARGUMENT_DRIVEN with 1. An expected allocation of 1 byte makes the check almost unconditional, so a regression in TF-IDF allocation estimation would not fail the test.

If encode_tfidf really allocates independently of nnz, add a short comment that states why. Otherwise use the same data-driven bound as the BM25 branches.

🧪 Proposed alignment with the BM25 branches

       if (coo_on) {
         raft::execute_with_dry_run_check(
           handle,
           [&](raft::resources const& h) {
             raft::sparse::matrix::encode_tfidf<float, int>(h, coo_a_matrix, result.view());
           },
-          raft::alloc_behavior::ARGUMENT_DRIVEN,
-          1);
+          raft::alloc_behavior::DATA_DRIVEN,
+          sizeof(float) * coo_a.nnz);
       } else {
         raft::execute_with_dry_run_check(
           handle,
           [&](raft::resources const& h) {
             raft::sparse::matrix::encode_tfidf<float, int>(h, csr_matrix, result.view());
           },
-          raft::alloc_behavior::ARGUMENT_DRIVEN,
-          1);
+          raft::alloc_behavior::DATA_DRIVEN,
+          sizeof(float) * coo_a.nnz);
       }

#!/bin/bash
# Inspect the encode_tfidf / encode_bm25 implementations to compare their allocation patterns.
set -euo pipefail

fd -t f 'preprocessing.cuh' cpp/include | while IFS= read -r f; do
  echo "=== $f ==="
  rg -n -C 12 'encode_tfidf|encode_bm25|get_dry_run_flag|device_uvector|make_device_' "$f"
done

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/sparse/preprocess.cu` around lines 166 - 180, Align both
encode_tfidf dry-run checks in the preprocessing test with the BM25 branches by
using DATA_DRIVEN allocation behavior and a bound based on sizeof(float)
multiplied by coo_a.nnz. If encode_tfidf intentionally allocates independently
of nnz, retain the current expectation and add a concise comment explaining that
behavior.

cpp/tests/util/dry_run_resources.cpp (1)

364-387: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Add coverage for dry_run_resources nested inside dry_run_resources.

dry_run_resources computes active_(!resource::get_dry_run_flag(existing)) and returns an empty memory_stats from get_bytes_peak() when inactive. This test covers dry-run nested in stats, but no test covers dry-run nested in dry-run. That inactive path controls whether the inner object skips init() and whether the destructor restores globals, so a regression there would silently produce zero statistics.

💚 Proposed additional test

TEST(DryRunResources, NestedDryRunIsInactive)
{
  raft::resources res;
  dry_run_resources outer(res);
  {
    dry_run_resources inner(outer);
    EXPECT_TRUE(resource::get_dry_run_flag(inner));
    // Inner is inactive: it reports no statistics of its own.
    EXPECT_EQ(inner.get_bytes_peak().total(), 0UL);
  }
  // The outer handle must still be in dry-run mode after the inner one is destroyed.
  EXPECT_TRUE(resource::get_dry_run_flag(outer));
}

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/util/dry_run_resources.cpp` around lines 364 - 387, Add a test
alongside NestedDryRunInStats that constructs an outer dry_run_resources and an
inner dry_run_resources using the outer resource. Verify the inner remains
marked as dry-run, reports zero from get_bytes_peak().total(), and that the
outer still reports the dry-run flag after the inner is destroyed.

cpp/tests/test_utils.cuh (1)

368-389: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

The diagnostic printf fires for expected DATA_DRIVEN mismatches.

In DATA_DRIVEN mode dry >= actual is the expected outcome, not a failure. The condition at Lines 368-371 triggers on any inequality, so every conforming DATA_DRIVEN test prints a full six-category dump. Gate the diagnostic on the assertion actually failing, or restrict it to the ARGUMENT_DRIVEN and NO_ALLOCATIONS modes.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/test_utils.cuh` around lines 368 - 389, Update the mismatch
diagnostic around the dry/actual workspace comparison so it does not print for
expected DATA_DRIVEN cases where dry exceeds actual. Gate the printf using the
same assertion-failure condition, or limit it to ARGUMENT_DRIVEN and
NO_ALLOCATIONS modes, while preserving diagnostics for genuine mismatches.

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/include/raft/mr/dry_run_resource.hpp`:
- Around line 90-105: Update probe_container’s three-argument construction and
destructor to track whether allocation used allocate_sync or the asynchronous
allocate API. Use that stored allocation mode to call the corresponding
deallocate_sync or deallocate function, rather than selecting solely via
cuda::mr::resource<MR>, while preserving the existing stream choice for
asynchronous deallocation.

---

Outside diff comments:
In `@cpp/include/raft/linalg/power.cuh`:
- Around line 79-95: Move the dry-run guard below all RAFT_EXPECTS validation in
both power and power_scalar, so contiguity and size checks always execute before
skipping the kernel launch. Preserve the existing validation and dry-run
behavior, and verify that power_scalar’s dereference of scalar.data_handle()
remains safe with make_host_scalar in dry-run mode; adjust only if that pointer
is not readable.

In `@cpp/include/raft/matrix/detail/select_radix.cuh`:
- Around line 1280-1337: Ensure dry-run execution in the radix top-k dispatch
does not invoke CUDA Runtime occupancy APIs. Update calc_chunk_size and
calc_grid_dim, and the radix_topk_one_block/radix_topk paths that call them, to
use a deterministic workspace-sizing path when dry_run is true and only query
CUDA occupancy for real execution; preserve normal dispatch and sizing behavior
otherwise.

---

Major comments:
In `@cpp/include/raft/core/host_container_policy.hpp`:
- Around line 118-127: Update host_container::reallocate so bytesize_ is reset
to zero immediately after deallocating data_ and clearing data_. Preserve the
existing allocation and swap flow, ensuring an allocation failure cannot leave a
null data_ pointer paired with a stale nonzero size.

In `@cpp/include/raft/linalg/detail/lstsq.cuh`:
- Around line 213-216: Make cleanup of gesvdj_params exception-safe in the code
surrounding cusolverDnCreateGesvdjInfo: establish an RAII owner or scope guard
immediately after successful creation that calls cusolverDnDestroyGesvdjInfo
exactly once on every exit path, including dry-run returns, allocation failures,
and checked cuSOLVER errors. Remove the manual dry-run-only destruction in the
lstsqSvdJacobi flow to avoid double cleanup.

In `@cpp/include/raft/matrix/detail/select_k-inl.cuh`:
- Around line 131-132: Update the CUB sort calls in the select-k implementation,
including both SortPairs and SortPairsDescending workspace-size queries and
execution calls, to use RAFT_CUDA_TRY. Ensure each status is checked before the
resource::get_dry_run_flag(handle) early return so failed dry-run workspace
queries propagate the CUDA error.

In `@cpp/include/raft/matrix/threshold.cuh`:
- Line 38: Move the resource::get_dry_run_flag(handle) guard in the threshold
operation after the RAFT_EXPECTS input/output size validation, so dimension
mismatches are rejected in both modes while dry-run still skips
detail::setSmallValuesZero and other CUDA work.

In `@cpp/include/raft/mr/dry_run_resource.hpp`:
- Around line 181-207: Update allocate_sync and allocate so the shared
state_->probe always satisfies the strictest alignment requested by RAFT, rather
than preserving the first allocation’s alignment; reallocate or otherwise
replace the probe when a later request requires greater alignment. Also size the
probe to accommodate the requested bytes, or otherwise ensure returned storage
cannot be overrun by an unguarded request, while preserving
allocation/deallocation accounting.

In `@cpp/include/raft/random/detail/rng_impl.cuh`:
- Around line 16-23: Add direct includes for the declarations used by the RNG
implementation: include the dry-run flag and CUDA stream resource headers before
the code calling resource::get_dry_run_flag and resource::get_cuda_stream. Keep
the existing scatter and CUB includes unchanged.

In `@cpp/include/raft/solver/detail/lap_functions.cuh`:
- Around line 269-271: In cpp/include/raft/solver/detail/lap_functions.cuh at
lines 269-271 and 432-436, compute a checked size_t matrix_size from SP and N
before constructing the device_uvector instances; use that variable for
predicates_v, addresses_v, csr_neighbors_v, and elements_v at the respective
sites, preserving the existing allocation behavior while preventing source-type
overflow.

In `@cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh`:
- Around line 309-315: Construct scan_ws with device_memory, matching the
resource used for sub_nnz, so the CUB scan workspace is accounted for by the
configured workspace resource.

In `@cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh`:
- Around line 123-126: Wrap both cub::DeviceScan::ExclusiveSum calls in
RAFT_CUDA_TRY, including the workspace-size query and the actual scan execution.
Ensure any CUDA failure is propagated before using scan_ws or the generated CSR
row offsets.

In `@cpp/include/raft/sparse/op/detail/reduce.cuh`:
- Around line 137-149: Validate before both CUB scan calls that diff.size() is
no greater than the maximum representable int, and reject the input with the
established validation mechanism when it exceeds that bound. Update the
workspace query and scan in the surrounding reduction flow to use only a
validated item count, preserving full coverage of all nonzeros and preventing
narrowing overflow.
- Around line 136-149: Wrap both cub::DeviceScan::ExclusiveSum calls in
RAFT_CUDA_TRY, including the workspace-size query and the actual scan, so CUDA
failures are propagated before scan_ws allocation or subsequent data use.

In `@cpp/include/raft/sparse/solver/detail/lanczos.cuh`:
- Around line 158-169: Add RAFT_CUDA_TRY(cudaPeekAtLastError()) immediately
after each of the five CUDA launches in lanczos.cuh: kernel_triangular_populate
and kernel_triangular_beta_k in lines 158-169, plus kernel_clamp_down,
kernel_clamp_down_vector, and kernel_normalize in lines 386-415. Keep the checks
directly adjacent to their respective launches.
- Around line 198-227: Use a handle-bound Thrust execution policy for all
operations in this selection block: define thrust_policy via
resource::get_thrust_policy(handle), then pass it to thrust::sequence and both
thrust::sort calls instead of thrust::device. Preserve the existing index
ordering and eigenvalue selection logic.

In `@cpp/include/raft/spectral/detail/matrix_wrappers.hpp`:
- Around line 393-396: Update the handle declarations in both
laplacian_matrix_t::mv
(cpp/include/raft/spectral/detail/matrix_wrappers.hpp:393-396) and
modularity_matrix_t::mv
(cpp/include/raft/spectral/detail/matrix_wrappers.hpp:459-462) to bind the
resources returned by sparse_matrix_t::get_handle() as a const reference instead
of copying them; leave the cublas, stream, and dry-run retrieval unchanged.

In `@cpp/include/raft/stats/detail/contingencyMatrix.cuh`:
- Around line 209-219: Prevent signed overflow in the dry-run calculations
within the contingency-matrix workspace sizing block by validating nSamples
against the supported range and promoting it to size_t before every
multiplication, including the tmpStagingMemorySize and cubWorkspaceUpperBound
expressions. Preserve the existing alignment and workspace estimate behavior for
valid sample counts, and reject or otherwise handle values that cannot be
represented safely.

In `@cpp/include/raft/stats/detail/meanvar.cuh`:
- Around line 212-223: Move the cudaOccupancyMaxActiveBlocksPerMultiprocessor
call and related gs.y adjustment into the existing !dry_run branch around
meanvar_kernel_rowmajor and meanvar_kernel_fill. Ensure dry_run performs no CUDA
runtime queries or kernel work while retaining the existing buffer allocation
and tracking behavior.

In `@cpp/include/raft/stats/detail/scores.cuh`:
- Line 183: Move the dry_run return in the surrounding score computation before
constructing the host vectors mean_errors and h_sorted_abs_diffs, while keeping
the existing RMM/RAFT allocation attempts before that return. Ensure dry-run
mode performs no host-vector allocations and exits immediately after those
allocation attempts.

In `@cpp/include/raft/stats/detail/silhouette_score.cuh`:
- Around line 229-230: Update the operations in the silhouette-score function,
including pairwise_distance and the dependent matrix_vector_op/reduction calls,
to use the same RAFT-managed stream obtained via
resource::get_cuda_stream(handle). Preserve asynchronous ordering and ensure
every producer and consumer of distanceMatrix and
averageDistanceBetweenSampleAndCluster is submitted to that stream.

In `@cpp/tests/core/bitset.cu`:
- Around line 388-399: Gate the initial my_bitset.any(h) and my_bitset.none(h)
assertions after reset on resource::get_dry_run_flag(h), matching the existing
guarded value-dependent assertions below. Keep the reset and subsequent setup
unchanged, and preserve the assertions during non-dry-run execution.

In `@cpp/tests/sparse/reduce.cu`:
- Around line 65-82: Move construction of the COO variable out of the
surrounding scope and into the dry-run callback containing max_duplicates,
keeping its entire lifetime within execute_with_dry_run_check. Remove or update
the inaccurate comments claiming the COO output is not tracked, and ensure the
callback’s output remains available to the subsequent test logic as required.

---

Minor comments:
In `@docs/source/developer_guide.md`:
- Around line 325-329: Insert a blank line after the final nested bullet under
the listed memory resources and before “All attempted allocations...” so that
sentence renders as a standalone paragraph rather than as continuation text
within the bullet.

---

Nitpick comments:
In `@cpp/tests/sparse/preprocess.cu`:
- Around line 166-180: Align both encode_tfidf dry-run checks in the
preprocessing test with the BM25 branches by using DATA_DRIVEN allocation
behavior and a bound based on sizeof(float) multiplied by coo_a.nnz. If
encode_tfidf intentionally allocates independently of nnz, retain the current
expectation and add a concise comment explaining that behavior.

In `@cpp/tests/test_utils.cuh`:
- Around line 368-389: Update the mismatch diagnostic around the dry/actual
workspace comparison so it does not print for expected DATA_DRIVEN cases where
dry exceeds actual. Gate the printf using the same assertion-failure condition,
or limit it to ARGUMENT_DRIVEN and NO_ALLOCATIONS modes, while preserving
diagnostics for genuine mismatches.

In `@cpp/tests/util/dry_run_resources.cpp`:
- Around line 364-387: Add a test alongside NestedDryRunInStats that constructs
an outer dry_run_resources and an inner dry_run_resources using the outer
resource. Verify the inner remains marked as dry-run, reports zero from
get_bytes_peak().total(), and that the outer still reports the dry-run flag
after the inner is destroyed.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

[coderabbitai]

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/REVIEW_GUIDELINES.md`:
- Around line 52-54: Revise the dry-run early-return guidance to remove the
broad “or when no other compliant structure works” exception. Permit an early
return only with documented proof that the skipped path and all callees cannot
allocate, backed by a dry-run test; otherwise require unconditional delegation
through public wrappers and guard only meaningful CUDA work, consistent with
docs/source/dry_run_protocol.md and the critical wrapper example.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: c3955160-feeb-42cd-b405-d71bcff7eb40

📥 Commits

Reviewing files that changed from the base of the PR and between d2031cf and c1101d8.

📒 Files selected for processing (3)

• .coderabbit.yaml
• cpp/REVIEW_GUIDELINES.md
• cpp/include/raft/core/dry_run_resources.hpp

🚧 Files skipped from review as they are similar to previous changes (1)

• cpp/include/raft/core/dry_run_resources.hpp

[coderabbitai]

Actionable comments posted: 3

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/include/raft/sparse/linalg/detail/sddmm.hpp`:
- Around line 70-71: Move the resource::get_dry_run_flag(handle) check in the
surrounding SDDMM function to immediately after the workspace query and tmp
allocation, before resource::sync_stream(handle). Preserve those dry-run setup
operations, while ensuring synchronization and subsequent cuSPARSE execution
occur only on the non-dry-run path.

In `@cpp/REVIEW_GUIDELINES.md`:
- Line 216: Update both opening fenced code blocks in cpp/REVIEW_GUIDELINES.md
at lines 216-216 and 228-228 to include the text language identifier, changing
each fence to ```text.

In `@cpp/tests/sparse/solver/randomized_svds.cu`:
- Around line 527-571: Add a dry-run guard in both apply and apply_transpose
after the base operation and local allocations, using
resource::get_dry_run_flag(handle) to return before map, gemv, and gemm
operations. Keep the base_op_ call and allocations before the guard so
allocation requirements remain visible, while preventing CUDA work from
accessing dry-run buffers.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 43afc17c-4abd-45c0-a29f-51b135a6aab1

📥 Commits

Reviewing files that changed from the base of the PR and between d254b0e and c1101d8.

📒 Files selected for processing (231)

• .coderabbit.yaml
• cpp/REVIEW_GUIDELINES.md
• cpp/include/raft/core/bitset.cuh
• cpp/include/raft/core/bitset.hpp
• cpp/include/raft/core/coo_matrix.hpp
• cpp/include/raft/core/copy.cuh
• cpp/include/raft/core/copy.hpp
• cpp/include/raft/core/csr_matrix.hpp
• cpp/include/raft/core/detail/copy.hpp
• cpp/include/raft/core/device_container_policy.hpp
• cpp/include/raft/core/device_mdarray.hpp
• cpp/include/raft/core/dry_run_resources.hpp
• cpp/include/raft/core/host_container_policy.hpp
• cpp/include/raft/core/host_mdarray.hpp
• cpp/include/raft/core/managed_mdarray.hpp
• cpp/include/raft/core/pinned_mdarray.hpp
• cpp/include/raft/core/resource/cuda_stream.hpp
• cpp/include/raft/core/resource/dry_run_flag.hpp
• cpp/include/raft/core/resource/resource_types.hpp
• cpp/include/raft/core/sparse_types.hpp
• cpp/include/raft/label/classlabels.cuh
• cpp/include/raft/label/detail/classlabels.cuh
• cpp/include/raft/linalg/add.cuh
• cpp/include/raft/linalg/coalesced_reduction.cuh
• cpp/include/raft/linalg/detail/axpy.cuh
• cpp/include/raft/linalg/detail/cholesky_r1_update.cuh
• cpp/include/raft/linalg/detail/coalesced_reduction-inl.cuh
• cpp/include/raft/linalg/detail/cublaslt_wrappers.hpp
• cpp/include/raft/linalg/detail/eig.cuh
• cpp/include/raft/linalg/detail/gemv.hpp
• cpp/include/raft/linalg/detail/lstsq.cuh
• cpp/include/raft/linalg/detail/map.cuh
• cpp/include/raft/linalg/detail/matrix_vector_op.cuh
• cpp/include/raft/linalg/detail/norm.cuh
• cpp/include/raft/linalg/detail/pca.cuh
• cpp/include/raft/linalg/detail/qr.cuh
• cpp/include/raft/linalg/detail/reduce.cuh
• cpp/include/raft/linalg/detail/rsvd.cuh
• cpp/include/raft/linalg/detail/svd.cuh
• cpp/include/raft/linalg/detail/transpose.cuh
• cpp/include/raft/linalg/divide.cuh
• cpp/include/raft/linalg/dot.cuh
• cpp/include/raft/linalg/map_reduce.cuh
• cpp/include/raft/linalg/matrix_vector_op.cuh
• cpp/include/raft/linalg/mean_squared_error.cuh
• cpp/include/raft/linalg/multiply.cuh
• cpp/include/raft/linalg/norm.cuh
• cpp/include/raft/linalg/normalize.cuh
• cpp/include/raft/linalg/power.cuh
• cpp/include/raft/linalg/reduce.cuh
• cpp/include/raft/linalg/reduce_cols_by_key.cuh
• cpp/include/raft/linalg/reduce_rows_by_key.cuh
• cpp/include/raft/linalg/sqrt.cuh
• cpp/include/raft/linalg/strided_reduction.cuh
• cpp/include/raft/linalg/subtract.cuh
• cpp/include/raft/linalg/unary_op.cuh
• cpp/include/raft/matrix/argmax.cuh
• cpp/include/raft/matrix/argmin.cuh
• cpp/include/raft/matrix/col_wise_sort.cuh
• cpp/include/raft/matrix/copy.cuh
• cpp/include/raft/matrix/detail/columnWiseSort.cuh
• cpp/include/raft/matrix/detail/gather.cuh
• cpp/include/raft/matrix/detail/gather_inplace.cuh
• cpp/include/raft/matrix/detail/math.cuh
• cpp/include/raft/matrix/detail/matrix.cuh
• cpp/include/raft/matrix/detail/scatter_inplace.cuh
• cpp/include/raft/matrix/detail/select_k-inl.cuh
• cpp/include/raft/matrix/detail/select_radix.cuh
• cpp/include/raft/matrix/detail/select_warpsort.cuh
• cpp/include/raft/matrix/detail/shift.cuh
• cpp/include/raft/matrix/diagonal.cuh
• cpp/include/raft/matrix/gather.cuh
• cpp/include/raft/matrix/init.cuh
• cpp/include/raft/matrix/linewise_op.cuh
• cpp/include/raft/matrix/norm.cuh
• cpp/include/raft/matrix/power.cuh
• cpp/include/raft/matrix/print.cuh
• cpp/include/raft/matrix/reciprocal.cuh
• cpp/include/raft/matrix/reverse.cuh
• cpp/include/raft/matrix/sign_flip.cuh
• cpp/include/raft/matrix/slice.cuh
• cpp/include/raft/matrix/sqrt.cuh
• cpp/include/raft/matrix/threshold.cuh
• cpp/include/raft/matrix/triangular.cuh
• cpp/include/raft/mr/dry_run_resource.hpp
• cpp/include/raft/random/detail/make_blobs.cuh
• cpp/include/raft/random/detail/make_regression.cuh
• cpp/include/raft/random/detail/multi_variable_gaussian.cuh
• cpp/include/raft/random/detail/rmat_rectangular_generator.cuh
• cpp/include/raft/random/detail/rng_impl.cuh
• cpp/include/raft/random/detail/rng_impl_deprecated.cuh
• cpp/include/raft/random/make_blobs.cuh
• cpp/include/raft/random/permute.cuh
• cpp/include/raft/random/rng.cuh
• cpp/include/raft/random/sample_without_replacement.cuh
• cpp/include/raft/solver/detail/lap_functions.cuh
• cpp/include/raft/sparse/convert/detail/adj_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/csr.cuh
• cpp/include/raft/sparse/linalg/detail/laplacian.cuh
• cpp/include/raft/sparse/linalg/detail/sddmm.hpp
• cpp/include/raft/sparse/linalg/detail/spmm.hpp
• cpp/include/raft/sparse/linalg/detail/symmetrize.cuh
• cpp/include/raft/sparse/linalg/detail/utils.cuh
• cpp/include/raft/sparse/linalg/spmm.hpp
• cpp/include/raft/sparse/linalg/transpose.cuh
• cpp/include/raft/sparse/matrix/detail/preprocessing.cuh
• cpp/include/raft/sparse/matrix/preprocessing.cuh
• cpp/include/raft/sparse/op/detail/filter.cuh
• cpp/include/raft/sparse/op/detail/reduce.cuh
• cpp/include/raft/sparse/op/detail/sort.h
• cpp/include/raft/sparse/op/sort.cuh
• cpp/include/raft/sparse/solver/detail/cholesky_qr.cuh
• cpp/include/raft/sparse/solver/detail/lanczos.cuh
• cpp/include/raft/sparse/solver/detail/randomized_svds.cuh
• cpp/include/raft/sparse/solver/detail/svds_sign_correction.cuh
• cpp/include/raft/spectral/detail/matrix_wrappers.hpp
• cpp/include/raft/spectral/detail/modularity_maximization.hpp
• cpp/include/raft/spectral/detail/partition.hpp
• cpp/include/raft/spectral/detail/spectral_util.cuh
• cpp/include/raft/stats/accuracy.cuh
• cpp/include/raft/stats/adjusted_rand_index.cuh
• cpp/include/raft/stats/completeness_score.cuh
• cpp/include/raft/stats/contingency_matrix.cuh
• cpp/include/raft/stats/detail/adjusted_rand_index.cuh
• cpp/include/raft/stats/detail/batched/silhouette_score.cuh
• cpp/include/raft/stats/detail/contingencyMatrix.cuh
• cpp/include/raft/stats/detail/cov.cuh
• cpp/include/raft/stats/detail/dispersion.cuh
• cpp/include/raft/stats/detail/entropy.cuh
• cpp/include/raft/stats/detail/homogeneity_score.cuh
• cpp/include/raft/stats/detail/kl_divergence.cuh
• cpp/include/raft/stats/detail/mean.cuh
• cpp/include/raft/stats/detail/mean_center.cuh
• cpp/include/raft/stats/detail/meanvar.cuh
• cpp/include/raft/stats/detail/mutual_info_score.cuh
• cpp/include/raft/stats/detail/neighborhood_recall.cuh
• cpp/include/raft/stats/detail/rand_index.cuh
• cpp/include/raft/stats/detail/scores.cuh
• cpp/include/raft/stats/detail/silhouette_score.cuh
• cpp/include/raft/stats/detail/stddev.cuh
• cpp/include/raft/stats/detail/sum.cuh
• cpp/include/raft/stats/detail/trustworthiness_score.cuh
• cpp/include/raft/stats/detail/v_measure.cuh
• cpp/include/raft/stats/detail/weighted_mean.cuh
• cpp/include/raft/stats/dispersion.cuh
• cpp/include/raft/stats/entropy.cuh
• cpp/include/raft/stats/histogram.cuh
• cpp/include/raft/stats/homogeneity_score.cuh
• cpp/include/raft/stats/information_criterion.cuh
• cpp/include/raft/stats/kl_divergence.cuh
• cpp/include/raft/stats/mean.cuh
• cpp/include/raft/stats/mean_center.cuh
• cpp/include/raft/stats/meanvar.cuh
• cpp/include/raft/stats/minmax.cuh
• cpp/include/raft/stats/mutual_info_score.cuh
• cpp/include/raft/stats/neighborhood_recall.cuh
• cpp/include/raft/stats/r2_score.cuh
• cpp/include/raft/stats/rand_index.cuh
• cpp/include/raft/stats/regression_metrics.cuh
• cpp/include/raft/stats/stddev.cuh
• cpp/include/raft/stats/sum.cuh
• cpp/include/raft/stats/v_measure.cuh
• cpp/include/raft/stats/weighted_mean.cuh
• cpp/tests/CMakeLists.txt
• cpp/tests/core/bitmap.cu
• cpp/tests/core/bitset.cu
• cpp/tests/core/mdarray.cu
• cpp/tests/core/mdbuffer.cu
• cpp/tests/core/mdspan_copy.cu
• cpp/tests/core/mdspan_utils.cu
• cpp/tests/core/sparse_matrix.cu
• cpp/tests/core/temporary_device_buffer.cu
• cpp/tests/linalg/add.cu
• cpp/tests/linalg/axpy.cu
• cpp/tests/linalg/binary_op.cu
• cpp/tests/linalg/coalesced_reduction.cu
• cpp/tests/linalg/divide.cu
• cpp/tests/linalg/dot.cu
• cpp/tests/linalg/eig.cu
• cpp/tests/linalg/gemm_basic.cpp
• cpp/tests/linalg/gemm_layout.cu
• cpp/tests/linalg/map.cu
• cpp/tests/linalg/matrix_vector_op.cu
• cpp/tests/linalg/mean_squared_error.cu
• cpp/tests/linalg/multiply.cu
• cpp/tests/linalg/norm.cu
• cpp/tests/linalg/normalize.cu
• cpp/tests/linalg/pca.cu
• cpp/tests/linalg/power.cu
• cpp/tests/linalg/reduce.cu
• cpp/tests/linalg/rsvd.cu
• cpp/tests/linalg/strided_reduction.cu
• cpp/tests/linalg/subtract.cu
• cpp/tests/linalg/svd.cu
• cpp/tests/linalg/ternary_op.cu
• cpp/tests/linalg/transpose.cu
• cpp/tests/linalg/unary_op.cu
• cpp/tests/matrix/gather.cu
• cpp/tests/matrix/sample_rows.cu
• cpp/tests/sparse/convert_csr.cu
• cpp/tests/sparse/csr_transpose.cu
• cpp/tests/sparse/diagonal.cu
• cpp/tests/sparse/filter.cu
• cpp/tests/sparse/masked_matmul.cu
• cpp/tests/sparse/norm.cu
• cpp/tests/sparse/preprocess.cu
• cpp/tests/sparse/reduce.cu
• cpp/tests/sparse/sddmm.cu
• cpp/tests/sparse/select_k_csr.cu
• cpp/tests/sparse/solver/lanczos.cu
• cpp/tests/sparse/solver/randomized_svds.cu
• cpp/tests/sparse/spmm.cu
• cpp/tests/sparse/symmetrize.cu
• cpp/tests/stats/cov.cu
• cpp/tests/stats/homogeneity_score.cu
• cpp/tests/stats/mean.cu
• cpp/tests/stats/meanvar.cu
• cpp/tests/stats/minmax.cu
• cpp/tests/stats/stddev.cu
• cpp/tests/stats/sum.cu
• cpp/tests/stats/weighted_mean.cu
• cpp/tests/test_utils.cuh
• cpp/tests/util/dry_run_guards.cu
• cpp/tests/util/dry_run_resources.cpp
• cpp/tests/util/memory_type_dispatcher.cu
• cpp/tests/util/popc.cu
• docs/source/cpp_api/core_resources.rst
• docs/source/developer_guide.md
• docs/source/dry_run_protocol.md

💤 Files with no reviewable changes (1)

• cpp/include/raft/stats/detail/sum.cuh

🚧 Files skipped from review as they are similar to previous changes (226)

• cpp/tests/CMakeLists.txt
• cpp/include/raft/matrix/detail/select_k-inl.cuh
• cpp/include/raft/core/coo_matrix.hpp
• cpp/include/raft/linalg/map_reduce.cuh
• cpp/include/raft/matrix/reverse.cuh
• cpp/include/raft/sparse/convert/detail/adj_to_csr.cuh
• cpp/include/raft/matrix/threshold.cuh
• cpp/include/raft/linalg/divide.cuh
• cpp/include/raft/linalg/reduce_rows_by_key.cuh
• cpp/include/raft/linalg/detail/map.cuh
• cpp/include/raft/linalg/detail/axpy.cuh
• cpp/include/raft/linalg/reduce_cols_by_key.cuh
• cpp/include/raft/matrix/print.cuh
• cpp/include/raft/sparse/linalg/detail/symmetrize.cuh
• cpp/include/raft/matrix/argmax.cuh
• cpp/include/raft/linalg/unary_op.cuh
• cpp/include/raft/stats/detail/homogeneity_score.cuh
• cpp/include/raft/linalg/normalize.cuh
• cpp/include/raft/stats/detail/cov.cuh
• cpp/include/raft/linalg/detail/cublaslt_wrappers.hpp
• cpp/include/raft/stats/detail/rand_index.cuh
• cpp/include/raft/matrix/reciprocal.cuh
• cpp/include/raft/matrix/slice.cuh
• cpp/include/raft/linalg/detail/eig.cuh
• cpp/include/raft/linalg/dot.cuh
• cpp/include/raft/label/classlabels.cuh
• cpp/tests/linalg/add.cu
• cpp/include/raft/matrix/triangular.cuh
• cpp/include/raft/linalg/norm.cuh
• cpp/include/raft/sparse/matrix/preprocessing.cuh
• cpp/include/raft/stats/adjusted_rand_index.cuh
• cpp/include/raft/stats/sum.cuh
• cpp/include/raft/linalg/detail/coalesced_reduction-inl.cuh
• cpp/include/raft/stats/v_measure.cuh
• cpp/include/raft/sparse/convert/detail/csr.cuh
• cpp/tests/stats/minmax.cu
• cpp/include/raft/sparse/op/detail/filter.cuh
• cpp/include/raft/core/copy.cuh
• docs/source/cpp_api/core_resources.rst
• cpp/tests/stats/meanvar.cu
• cpp/tests/core/bitmap.cu
• cpp/include/raft/sparse/solver/detail/svds_sign_correction.cuh
• cpp/include/raft/core/detail/copy.hpp
• cpp/include/raft/random/permute.cuh
• cpp/include/raft/linalg/strided_reduction.cuh
• cpp/include/raft/core/resource/cuda_stream.hpp
• cpp/include/raft/stats/detail/kl_divergence.cuh
• cpp/tests/linalg/mean_squared_error.cu
• cpp/include/raft/matrix/linewise_op.cuh
• cpp/include/raft/core/bitset.cuh
• cpp/include/raft/matrix/argmin.cuh
• cpp/include/raft/linalg/detail/svd.cuh
• cpp/include/raft/core/host_mdarray.hpp
• cpp/include/raft/linalg/detail/cholesky_r1_update.cuh
• cpp/include/raft/linalg/mean_squared_error.cuh
• cpp/include/raft/linalg/multiply.cuh
• cpp/tests/linalg/gemm_basic.cpp
• cpp/include/raft/matrix/init.cuh
• cpp/include/raft/core/sparse_types.hpp
• cpp/include/raft/matrix/detail/matrix.cuh
• cpp/include/raft/matrix/power.cuh
• cpp/include/raft/core/device_mdarray.hpp
• cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh
• cpp/tests/core/mdbuffer.cu
• cpp/tests/core/mdspan_copy.cu
• cpp/include/raft/sparse/solver/detail/randomized_svds.cuh
• cpp/tests/linalg/strided_reduction.cu
• cpp/tests/linalg/multiply.cu
• cpp/include/raft/matrix/detail/math.cuh
• cpp/include/raft/stats/mutual_info_score.cuh
• cpp/include/raft/random/detail/make_regression.cuh
• cpp/include/raft/sparse/linalg/detail/utils.cuh
• cpp/include/raft/stats/detail/batched/silhouette_score.cuh
• cpp/include/raft/core/resource/resource_types.hpp
• cpp/include/raft/stats/accuracy.cuh
• cpp/include/raft/random/sample_without_replacement.cuh
• cpp/include/raft/matrix/detail/scatter_inplace.cuh
• cpp/include/raft/stats/detail/neighborhood_recall.cuh
• cpp/tests/sparse/masked_matmul.cu
• cpp/include/raft/stats/detail/weighted_mean.cuh
• cpp/include/raft/spectral/detail/modularity_maximization.hpp
• cpp/include/raft/stats/detail/dispersion.cuh
• cpp/include/raft/stats/minmax.cuh
• cpp/tests/core/sparse_matrix.cu
• cpp/include/raft/linalg/detail/rsvd.cuh
• cpp/tests/sparse/csr_transpose.cu
• cpp/include/raft/linalg/subtract.cuh
• cpp/tests/sparse/select_k_csr.cu
• cpp/include/raft/stats/meanvar.cuh
• cpp/include/raft/core/managed_mdarray.hpp
• cpp/include/raft/stats/detail/mean.cuh
• cpp/include/raft/sparse/linalg/detail/laplacian.cuh
• cpp/include/raft/linalg/add.cuh
• cpp/include/raft/core/resource/dry_run_flag.hpp
• cpp/include/raft/random/detail/rng_impl_deprecated.cuh
• cpp/include/raft/linalg/sqrt.cuh
• cpp/include/raft/stats/stddev.cuh
• cpp/include/raft/linalg/detail/reduce.cuh
• cpp/include/raft/sparse/linalg/transpose.cuh
• cpp/tests/linalg/power.cu
• cpp/include/raft/linalg/detail/transpose.cuh
• cpp/include/raft/stats/detail/contingencyMatrix.cuh
• cpp/tests/core/mdarray.cu
• cpp/include/raft/random/detail/rng_impl.cuh
• cpp/include/raft/random/detail/rmat_rectangular_generator.cuh
• cpp/tests/linalg/norm.cu
• cpp/tests/linalg/transpose.cu
• cpp/include/raft/matrix/detail/shift.cuh
• cpp/include/raft/stats/information_criterion.cuh
• cpp/include/raft/stats/histogram.cuh
• cpp/include/raft/sparse/op/detail/reduce.cuh
• cpp/include/raft/stats/weighted_mean.cuh
• cpp/include/raft/matrix/detail/gather_inplace.cuh
• cpp/include/raft/core/copy.hpp
• cpp/tests/linalg/binary_op.cu
• cpp/tests/linalg/gemm_layout.cu
• cpp/tests/sparse/sddmm.cu
• cpp/tests/sparse/norm.cu
• cpp/tests/util/popc.cu
• cpp/include/raft/stats/rand_index.cuh
• cpp/include/raft/matrix/norm.cuh
• cpp/tests/linalg/matrix_vector_op.cu
• cpp/include/raft/linalg/power.cuh
• cpp/tests/core/temporary_device_buffer.cu
• cpp/tests/linalg/divide.cu
• cpp/include/raft/label/detail/classlabels.cuh
• cpp/include/raft/linalg/detail/qr.cuh
• cpp/tests/linalg/dot.cu
• cpp/include/raft/matrix/gather.cuh
• cpp/tests/linalg/map.cu
• cpp/include/raft/matrix/detail/gather.cuh
• cpp/include/raft/stats/neighborhood_recall.cuh
• cpp/tests/linalg/ternary_op.cu
• cpp/include/raft/linalg/reduce.cuh
• cpp/include/raft/random/detail/multi_variable_gaussian.cuh
• cpp/tests/stats/sum.cu
• .coderabbit.yaml
• cpp/include/raft/stats/r2_score.cuh
• cpp/include/raft/linalg/coalesced_reduction.cuh
• cpp/include/raft/stats/regression_metrics.cuh
• cpp/include/raft/sparse/matrix/detail/preprocessing.cuh
• cpp/tests/sparse/diagonal.cu
• cpp/tests/stats/homogeneity_score.cu
• cpp/include/raft/matrix/diagonal.cuh
• cpp/include/raft/sparse/linalg/spmm.hpp
• cpp/include/raft/matrix/detail/columnWiseSort.cuh
• cpp/tests/matrix/sample_rows.cu
• cpp/tests/linalg/svd.cu
• cpp/include/raft/stats/homogeneity_score.cuh
• cpp/tests/linalg/unary_op.cu
• cpp/include/raft/matrix/sign_flip.cuh
• cpp/include/raft/linalg/detail/matrix_vector_op.cuh
• docs/source/developer_guide.md
• cpp/tests/util/dry_run_guards.cu
• cpp/include/raft/stats/completeness_score.cuh
• cpp/include/raft/sparse/op/detail/sort.h
• cpp/include/raft/linalg/detail/pca.cuh
• cpp/include/raft/stats/detail/adjusted_rand_index.cuh
• cpp/include/raft/core/dry_run_resources.hpp
• cpp/include/raft/sparse/op/sort.cuh
• cpp/include/raft/stats/detail/trustworthiness_score.cuh
• cpp/include/raft/spectral/detail/partition.hpp
• cpp/include/raft/stats/mean_center.cuh
• cpp/tests/linalg/pca.cu
• cpp/tests/util/memory_type_dispatcher.cu
• cpp/include/raft/random/detail/make_blobs.cuh
• cpp/include/raft/sparse/linalg/detail/spmm.hpp
• cpp/tests/linalg/coalesced_reduction.cu
• docs/source/dry_run_protocol.md
• cpp/include/raft/linalg/detail/gemv.hpp
• cpp/tests/linalg/axpy.cu
• cpp/include/raft/matrix/detail/select_warpsort.cuh
• cpp/tests/sparse/spmm.cu
• cpp/tests/linalg/subtract.cu
• cpp/tests/sparse/preprocess.cu
• cpp/include/raft/solver/detail/lap_functions.cuh
• cpp/tests/linalg/normalize.cu
• cpp/tests/stats/cov.cu
• cpp/include/raft/core/pinned_mdarray.hpp
• cpp/tests/sparse/symmetrize.cu
• cpp/include/raft/matrix/sqrt.cuh
• cpp/tests/linalg/reduce.cu
• cpp/include/raft/stats/detail/mean_center.cuh
• cpp/include/raft/matrix/copy.cuh
• cpp/include/raft/random/rng.cuh
• cpp/include/raft/sparse/solver/detail/lanczos.cuh
• cpp/include/raft/linalg/detail/lstsq.cuh
• cpp/include/raft/stats/contingency_matrix.cuh
• cpp/include/raft/linalg/matrix_vector_op.cuh
• cpp/include/raft/spectral/detail/spectral_util.cuh
• cpp/tests/sparse/solver/lanczos.cu
• cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh
• cpp/include/raft/stats/detail/entropy.cuh
• cpp/include/raft/stats/detail/meanvar.cuh
• cpp/tests/sparse/filter.cu
• cpp/tests/sparse/reduce.cu
• cpp/include/raft/stats/detail/silhouette_score.cuh
• cpp/include/raft/stats/detail/scores.cuh
• cpp/tests/linalg/eig.cu
• cpp/tests/stats/weighted_mean.cu
• cpp/include/raft/stats/detail/v_measure.cuh
• cpp/include/raft/stats/detail/mutual_info_score.cuh
• cpp/tests/test_utils.cuh
• cpp/include/raft/sparse/solver/detail/cholesky_qr.cuh
• cpp/include/raft/stats/kl_divergence.cuh
• cpp/include/raft/matrix/detail/select_radix.cuh
• cpp/include/raft/mr/dry_run_resource.hpp
• cpp/include/raft/linalg/detail/norm.cuh
• cpp/tests/sparse/convert_csr.cu
• cpp/include/raft/core/host_container_policy.hpp
• cpp/include/raft/matrix/col_wise_sort.cuh
• cpp/include/raft/random/make_blobs.cuh
• cpp/tests/matrix/gather.cu
• cpp/tests/linalg/rsvd.cu
• cpp/include/raft/stats/mean.cuh
• cpp/include/raft/stats/entropy.cuh
• cpp/include/raft/core/bitset.hpp
• cpp/include/raft/stats/dispersion.cuh
• cpp/include/raft/core/device_container_policy.hpp
• cpp/tests/stats/stddev.cu
• cpp/tests/stats/mean.cu
• cpp/tests/core/bitset.cu
• cpp/include/raft/stats/detail/stddev.cuh
• cpp/tests/core/mdspan_utils.cu
• cpp/include/raft/spectral/detail/matrix_wrappers.hpp
• cpp/include/raft/core/csr_matrix.hpp

[coderabbitai]

🩺 Stability & Availability | 🔴 Critical | ⚡ Quick win

Guard the mean-centering CUDA work in dry-run mode.

During execute_with_dry_run_check, corr, sums, and ones use dry-run allocations. Lines 531, 537, 541, 558, 564, and 565 then launch work that reads or writes those probe-backed buffers. This can cause invalid device-memory access.

After base_op_.apply or base_op_.apply_transpose and the local allocations, return when resource::get_dry_run_flag(handle) is true. Keep the base_op_ call and allocations before the check so their allocation requirements remain visible.

Proposed guard placement

     base_op_.apply(handle, X, Y);
     auto corr = raft::make_device_vector<ValueType, uint32_t>(handle, bk);
     auto ones = raft::make_device_vector<ValueType, uint32_t>(handle, m_);
+    if (resource::get_dry_run_flag(handle)) { return; }
     raft::linalg::map(handle, ones.view(), raft::const_op{ValueType{1}});

Apply the equivalent change in apply_transpose.

As per path instructions, allocations must remain active and meaningful CUDA work must be guarded in dry-run mode.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/sparse/solver/randomized_svds.cu` around lines 527 - 571, Add a
dry-run guard in both apply and apply_transpose after the base operation and
local allocations, using resource::get_dry_run_flag(handle) to return before
map, gemv, and gemm operations. Keep the base_op_ call and allocations before
the guard so allocation requirements remain visible, while preventing CUDA work
from accessing dry-run buffers.

Source: Path instructions

[coderabbitai]

Actionable comments posted: 1

♻️ Duplicate comments (1)

cpp/include/raft/core/dry_run_resources.hpp (1)

62-70: 🩺 Stability & Availability | 🔴 Critical | ⚡ Quick win

init() still has no rollback if it throws.

The constructor calls init() directly. init() replaces the global host resource at Line 180 and the global device resource at Line 207 before it completes. If any later step throws, the object is not fully constructed, so ~dry_run_resources() never runs. The global host and device resources then point at adaptors that are destroyed during stack unwinding.

Wrap init() in a try/catch that restores old_host_, old_device_, and the dry-run flag, then rethrows.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/core/dry_run_resources.hpp` around lines 62 - 70, Update the
dry_run_resources constructor around init() to provide rollback when
initialization throws: catch the exception, restore old_host_ and old_device_,
reset the dry-run flag to its prior state, then rethrow. Preserve the existing
successful initialization and destructor behavior.

Source: Path instructions

🧹 Nitpick comments (3)

cpp/include/raft/sparse/linalg/transpose.cuh (1)

68-70: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick win

Consider allocating the temporary buffer from the workspace resource.

convert_csc_workspace is a short-lived scratch buffer. It currently uses the default device memory resource, so the dry-run tracker attributes it to device_global instead of device_workspace. Allocating it from resource::get_workspace_resource(handle) gives correct categorization and honors the configured workspace limit.

rmm::device_uvector<char> convert_csc_workspace(
  convert_csc_workspace_size, stream, resource::get_workspace_resource(handle));

The dry-run ordering itself is correct: the buffer-size query and the allocation run before the guard, and only the numeric conversion is skipped.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/sparse/linalg/transpose.cuh` around lines 68 - 70, Update
the convert_csc_workspace allocation to use
resource::get_workspace_resource(handle) as its memory resource, while
preserving the existing allocation-before-dry-run-guard ordering and skipping
only the numeric conversion.

Source: Path instructions

cpp/tests/util/dry_run_guards.cu (2)

153-174: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

DeallocReducesCurrent does not assert the behavior in its name.

The test allocates buf1, frees it, then allocates buf2, but the only assertion is EXPECT_GE(stats.device_global, kSize2). That assertion also passes if deallocation is never accounted for. Add an upper bound so the test fails when the counter does not decrease on deallocation.

🧪 Proposed assertion

   EXPECT_GE(stats.device_global, kSize2);
-  // But could be less than kSize1 + kSize2 (since buf1 is freed before buf2)
-  // This depends on timing/implementation, so we just check the peak is reasonable
+  // buf1 is freed before buf2 is allocated, so the peak must not include both.
+  EXPECT_LT(stats.device_global, kSize1 + kSize2);

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/util/dry_run_guards.cu` around lines 153 - 174, Update the
DeallocReducesCurrent test to assert an upper bound on stats.device_global after
buf1 is freed before buf2 is allocated, ensuring the recorded peak does not
include both allocation sizes. Keep the existing lower-bound check for kSize2
and add a bound that fails when deallocation is not accounted for.

---

42-73: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Consider covering these paths with raft::execute_with_dry_run_check.

AddDoesNotExecute and RngDoesNotExecute toggle set_dry_run_flag on a live resources handle. That checks the guard but not the allocation contract, and it leaves the flag set on res if the call under test throws. The test guidelines require principal public paths to be covered with execute_with_dry_run_check and the matching allocation behavior.

As per coding guidelines: "When adding or materially changing raft::resources algorithms or primitives, cover principal public paths with execute_with_dry_run_check and the appropriate allocation behavior."

Also applies to: 75-100

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/tests/util/dry_run_guards.cu` around lines 42 - 73, Update the dry-run
tests around AddDoesNotExecute and RngDoesNotExecute to use
raft::execute_with_dry_run_check instead of manually toggling set_dry_run_flag
on a live resources handle. Configure the helper to verify the corresponding
allocation behavior while preserving assertions that the operation does not
modify output or execute RNG work, and ensure cleanup remains exception-safe.

Source: Coding guidelines

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/include/raft/core/host_container_policy.hpp`:
- Around line 118-127: Update host_container::reallocate so bytesize_ is reset
to zero immediately when the existing data_ buffer is deallocated and data_ is
cleared, before constructing host_container{count, mr_}; preserve the existing
allocation and swap behavior.

---

Duplicate comments:
In `@cpp/include/raft/core/dry_run_resources.hpp`:
- Around line 62-70: Update the dry_run_resources constructor around init() to
provide rollback when initialization throws: catch the exception, restore
old_host_ and old_device_, reset the dry-run flag to its prior state, then
rethrow. Preserve the existing successful initialization and destructor
behavior.

---

Nitpick comments:
In `@cpp/include/raft/sparse/linalg/transpose.cuh`:
- Around line 68-70: Update the convert_csc_workspace allocation to use
resource::get_workspace_resource(handle) as its memory resource, while
preserving the existing allocation-before-dry-run-guard ordering and skipping
only the numeric conversion.

In `@cpp/tests/util/dry_run_guards.cu`:
- Around line 153-174: Update the DeallocReducesCurrent test to assert an upper
bound on stats.device_global after buf1 is freed before buf2 is allocated,
ensuring the recorded peak does not include both allocation sizes. Keep the
existing lower-bound check for kSize2 and add a bound that fails when
deallocation is not accounted for.
- Around line 42-73: Update the dry-run tests around AddDoesNotExecute and
RngDoesNotExecute to use raft::execute_with_dry_run_check instead of manually
toggling set_dry_run_flag on a live resources handle. Configure the helper to
verify the corresponding allocation behavior while preserving assertions that
the operation does not modify output or execute RNG work, and ensure cleanup
remains exception-safe.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 43afc17c-4abd-45c0-a29f-51b135a6aab1

📥 Commits

Reviewing files that changed from the base of the PR and between d254b0e and c1101d8.

📒 Files selected for processing (231)

• .coderabbit.yaml
• cpp/REVIEW_GUIDELINES.md
• cpp/include/raft/core/bitset.cuh
• cpp/include/raft/core/bitset.hpp
• cpp/include/raft/core/coo_matrix.hpp
• cpp/include/raft/core/copy.cuh
• cpp/include/raft/core/copy.hpp
• cpp/include/raft/core/csr_matrix.hpp
• cpp/include/raft/core/detail/copy.hpp
• cpp/include/raft/core/device_container_policy.hpp
• cpp/include/raft/core/device_mdarray.hpp
• cpp/include/raft/core/dry_run_resources.hpp
• cpp/include/raft/core/host_container_policy.hpp
• cpp/include/raft/core/host_mdarray.hpp
• cpp/include/raft/core/managed_mdarray.hpp
• cpp/include/raft/core/pinned_mdarray.hpp
• cpp/include/raft/core/resource/cuda_stream.hpp
• cpp/include/raft/core/resource/dry_run_flag.hpp
• cpp/include/raft/core/resource/resource_types.hpp
• cpp/include/raft/core/sparse_types.hpp
• cpp/include/raft/label/classlabels.cuh
• cpp/include/raft/label/detail/classlabels.cuh
• cpp/include/raft/linalg/add.cuh
• cpp/include/raft/linalg/coalesced_reduction.cuh
• cpp/include/raft/linalg/detail/axpy.cuh
• cpp/include/raft/linalg/detail/cholesky_r1_update.cuh
• cpp/include/raft/linalg/detail/coalesced_reduction-inl.cuh
• cpp/include/raft/linalg/detail/cublaslt_wrappers.hpp
• cpp/include/raft/linalg/detail/eig.cuh
• cpp/include/raft/linalg/detail/gemv.hpp
• cpp/include/raft/linalg/detail/lstsq.cuh
• cpp/include/raft/linalg/detail/map.cuh
• cpp/include/raft/linalg/detail/matrix_vector_op.cuh
• cpp/include/raft/linalg/detail/norm.cuh
• cpp/include/raft/linalg/detail/pca.cuh
• cpp/include/raft/linalg/detail/qr.cuh
• cpp/include/raft/linalg/detail/reduce.cuh
• cpp/include/raft/linalg/detail/rsvd.cuh
• cpp/include/raft/linalg/detail/svd.cuh
• cpp/include/raft/linalg/detail/transpose.cuh
• cpp/include/raft/linalg/divide.cuh
• cpp/include/raft/linalg/dot.cuh
• cpp/include/raft/linalg/map_reduce.cuh
• cpp/include/raft/linalg/matrix_vector_op.cuh
• cpp/include/raft/linalg/mean_squared_error.cuh
• cpp/include/raft/linalg/multiply.cuh
• cpp/include/raft/linalg/norm.cuh
• cpp/include/raft/linalg/normalize.cuh
• cpp/include/raft/linalg/power.cuh
• cpp/include/raft/linalg/reduce.cuh
• cpp/include/raft/linalg/reduce_cols_by_key.cuh
• cpp/include/raft/linalg/reduce_rows_by_key.cuh
• cpp/include/raft/linalg/sqrt.cuh
• cpp/include/raft/linalg/strided_reduction.cuh
• cpp/include/raft/linalg/subtract.cuh
• cpp/include/raft/linalg/unary_op.cuh
• cpp/include/raft/matrix/argmax.cuh
• cpp/include/raft/matrix/argmin.cuh
• cpp/include/raft/matrix/col_wise_sort.cuh
• cpp/include/raft/matrix/copy.cuh
• cpp/include/raft/matrix/detail/columnWiseSort.cuh
• cpp/include/raft/matrix/detail/gather.cuh
• cpp/include/raft/matrix/detail/gather_inplace.cuh
• cpp/include/raft/matrix/detail/math.cuh
• cpp/include/raft/matrix/detail/matrix.cuh
• cpp/include/raft/matrix/detail/scatter_inplace.cuh
• cpp/include/raft/matrix/detail/select_k-inl.cuh
• cpp/include/raft/matrix/detail/select_radix.cuh
• cpp/include/raft/matrix/detail/select_warpsort.cuh
• cpp/include/raft/matrix/detail/shift.cuh
• cpp/include/raft/matrix/diagonal.cuh
• cpp/include/raft/matrix/gather.cuh
• cpp/include/raft/matrix/init.cuh
• cpp/include/raft/matrix/linewise_op.cuh
• cpp/include/raft/matrix/norm.cuh
• cpp/include/raft/matrix/power.cuh
• cpp/include/raft/matrix/print.cuh
• cpp/include/raft/matrix/reciprocal.cuh
• cpp/include/raft/matrix/reverse.cuh
• cpp/include/raft/matrix/sign_flip.cuh
• cpp/include/raft/matrix/slice.cuh
• cpp/include/raft/matrix/sqrt.cuh
• cpp/include/raft/matrix/threshold.cuh
• cpp/include/raft/matrix/triangular.cuh
• cpp/include/raft/mr/dry_run_resource.hpp
• cpp/include/raft/random/detail/make_blobs.cuh
• cpp/include/raft/random/detail/make_regression.cuh
• cpp/include/raft/random/detail/multi_variable_gaussian.cuh
• cpp/include/raft/random/detail/rmat_rectangular_generator.cuh
• cpp/include/raft/random/detail/rng_impl.cuh
• cpp/include/raft/random/detail/rng_impl_deprecated.cuh
• cpp/include/raft/random/make_blobs.cuh
• cpp/include/raft/random/permute.cuh
• cpp/include/raft/random/rng.cuh
• cpp/include/raft/random/sample_without_replacement.cuh
• cpp/include/raft/solver/detail/lap_functions.cuh
• cpp/include/raft/sparse/convert/detail/adj_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh
• cpp/include/raft/sparse/convert/detail/csr.cuh
• cpp/include/raft/sparse/linalg/detail/laplacian.cuh
• cpp/include/raft/sparse/linalg/detail/sddmm.hpp
• cpp/include/raft/sparse/linalg/detail/spmm.hpp
• cpp/include/raft/sparse/linalg/detail/symmetrize.cuh
• cpp/include/raft/sparse/linalg/detail/utils.cuh
• cpp/include/raft/sparse/linalg/spmm.hpp
• cpp/include/raft/sparse/linalg/transpose.cuh
• cpp/include/raft/sparse/matrix/detail/preprocessing.cuh
• cpp/include/raft/sparse/matrix/preprocessing.cuh
• cpp/include/raft/sparse/op/detail/filter.cuh
• cpp/include/raft/sparse/op/detail/reduce.cuh
• cpp/include/raft/sparse/op/detail/sort.h
• cpp/include/raft/sparse/op/sort.cuh
• cpp/include/raft/sparse/solver/detail/cholesky_qr.cuh
• cpp/include/raft/sparse/solver/detail/lanczos.cuh
• cpp/include/raft/sparse/solver/detail/randomized_svds.cuh
• cpp/include/raft/sparse/solver/detail/svds_sign_correction.cuh
• cpp/include/raft/spectral/detail/matrix_wrappers.hpp
• cpp/include/raft/spectral/detail/modularity_maximization.hpp
• cpp/include/raft/spectral/detail/partition.hpp
• cpp/include/raft/spectral/detail/spectral_util.cuh
• cpp/include/raft/stats/accuracy.cuh
• cpp/include/raft/stats/adjusted_rand_index.cuh
• cpp/include/raft/stats/completeness_score.cuh
• cpp/include/raft/stats/contingency_matrix.cuh
• cpp/include/raft/stats/detail/adjusted_rand_index.cuh
• cpp/include/raft/stats/detail/batched/silhouette_score.cuh
• cpp/include/raft/stats/detail/contingencyMatrix.cuh
• cpp/include/raft/stats/detail/cov.cuh
• cpp/include/raft/stats/detail/dispersion.cuh
• cpp/include/raft/stats/detail/entropy.cuh
• cpp/include/raft/stats/detail/homogeneity_score.cuh
• cpp/include/raft/stats/detail/kl_divergence.cuh
• cpp/include/raft/stats/detail/mean.cuh
• cpp/include/raft/stats/detail/mean_center.cuh
• cpp/include/raft/stats/detail/meanvar.cuh
• cpp/include/raft/stats/detail/mutual_info_score.cuh
• cpp/include/raft/stats/detail/neighborhood_recall.cuh
• cpp/include/raft/stats/detail/rand_index.cuh
• cpp/include/raft/stats/detail/scores.cuh
• cpp/include/raft/stats/detail/silhouette_score.cuh
• cpp/include/raft/stats/detail/stddev.cuh
• cpp/include/raft/stats/detail/sum.cuh
• cpp/include/raft/stats/detail/trustworthiness_score.cuh
• cpp/include/raft/stats/detail/v_measure.cuh
• cpp/include/raft/stats/detail/weighted_mean.cuh
• cpp/include/raft/stats/dispersion.cuh
• cpp/include/raft/stats/entropy.cuh
• cpp/include/raft/stats/histogram.cuh
• cpp/include/raft/stats/homogeneity_score.cuh
• cpp/include/raft/stats/information_criterion.cuh
• cpp/include/raft/stats/kl_divergence.cuh
• cpp/include/raft/stats/mean.cuh
• cpp/include/raft/stats/mean_center.cuh
• cpp/include/raft/stats/meanvar.cuh
• cpp/include/raft/stats/minmax.cuh
• cpp/include/raft/stats/mutual_info_score.cuh
• cpp/include/raft/stats/neighborhood_recall.cuh
• cpp/include/raft/stats/r2_score.cuh
• cpp/include/raft/stats/rand_index.cuh
• cpp/include/raft/stats/regression_metrics.cuh
• cpp/include/raft/stats/stddev.cuh
• cpp/include/raft/stats/sum.cuh
• cpp/include/raft/stats/v_measure.cuh
• cpp/include/raft/stats/weighted_mean.cuh
• cpp/tests/CMakeLists.txt
• cpp/tests/core/bitmap.cu
• cpp/tests/core/bitset.cu
• cpp/tests/core/mdarray.cu
• cpp/tests/core/mdbuffer.cu
• cpp/tests/core/mdspan_copy.cu
• cpp/tests/core/mdspan_utils.cu
• cpp/tests/core/sparse_matrix.cu
• cpp/tests/core/temporary_device_buffer.cu
• cpp/tests/linalg/add.cu
• cpp/tests/linalg/axpy.cu
• cpp/tests/linalg/binary_op.cu
• cpp/tests/linalg/coalesced_reduction.cu
• cpp/tests/linalg/divide.cu
• cpp/tests/linalg/dot.cu
• cpp/tests/linalg/eig.cu
• cpp/tests/linalg/gemm_basic.cpp
• cpp/tests/linalg/gemm_layout.cu
• cpp/tests/linalg/map.cu
• cpp/tests/linalg/matrix_vector_op.cu
• cpp/tests/linalg/mean_squared_error.cu
• cpp/tests/linalg/multiply.cu
• cpp/tests/linalg/norm.cu
• cpp/tests/linalg/normalize.cu
• cpp/tests/linalg/pca.cu
• cpp/tests/linalg/power.cu
• cpp/tests/linalg/reduce.cu
• cpp/tests/linalg/rsvd.cu
• cpp/tests/linalg/strided_reduction.cu
• cpp/tests/linalg/subtract.cu
• cpp/tests/linalg/svd.cu
• cpp/tests/linalg/ternary_op.cu
• cpp/tests/linalg/transpose.cu
• cpp/tests/linalg/unary_op.cu
• cpp/tests/matrix/gather.cu
• cpp/tests/matrix/sample_rows.cu
• cpp/tests/sparse/convert_csr.cu
• cpp/tests/sparse/csr_transpose.cu
• cpp/tests/sparse/diagonal.cu
• cpp/tests/sparse/filter.cu
• cpp/tests/sparse/masked_matmul.cu
• cpp/tests/sparse/norm.cu
• cpp/tests/sparse/preprocess.cu
• cpp/tests/sparse/reduce.cu
• cpp/tests/sparse/sddmm.cu
• cpp/tests/sparse/select_k_csr.cu
• cpp/tests/sparse/solver/lanczos.cu
• cpp/tests/sparse/solver/randomized_svds.cu
• cpp/tests/sparse/spmm.cu
• cpp/tests/sparse/symmetrize.cu
• cpp/tests/stats/cov.cu
• cpp/tests/stats/homogeneity_score.cu
• cpp/tests/stats/mean.cu
• cpp/tests/stats/meanvar.cu
• cpp/tests/stats/minmax.cu
• cpp/tests/stats/stddev.cu
• cpp/tests/stats/sum.cu
• cpp/tests/stats/weighted_mean.cu
• cpp/tests/test_utils.cuh
• cpp/tests/util/dry_run_guards.cu
• cpp/tests/util/dry_run_resources.cpp
• cpp/tests/util/memory_type_dispatcher.cu
• cpp/tests/util/popc.cu
• docs/source/cpp_api/core_resources.rst
• docs/source/developer_guide.md
• docs/source/dry_run_protocol.md

💤 Files with no reviewable changes (1)

• cpp/include/raft/stats/detail/sum.cuh

🚧 Files skipped from review as they are similar to previous changes (223)

• cpp/include/raft/sparse/linalg/detail/utils.cuh
• cpp/include/raft/matrix/detail/select_k-inl.cuh
• cpp/include/raft/matrix/detail/gather_inplace.cuh
• cpp/include/raft/core/copy.hpp
• cpp/include/raft/matrix/init.cuh
• cpp/include/raft/matrix/detail/matrix.cuh
• cpp/include/raft/linalg/reduce_rows_by_key.cuh
• cpp/tests/linalg/coalesced_reduction.cu
• cpp/include/raft/linalg/sqrt.cuh
• cpp/include/raft/matrix/detail/shift.cuh
• cpp/include/raft/linalg/reduce_cols_by_key.cuh
• cpp/include/raft/linalg/matrix_vector_op.cuh
• cpp/include/raft/linalg/strided_reduction.cuh
• cpp/include/raft/linalg/power.cuh
• cpp/tests/linalg/binary_op.cu
• cpp/include/raft/linalg/divide.cuh
• cpp/include/raft/linalg/mean_squared_error.cuh
• cpp/include/raft/matrix/sign_flip.cuh
• cpp/include/raft/linalg/map_reduce.cuh
• cpp/include/raft/linalg/detail/axpy.cuh
• cpp/include/raft/core/sparse_types.hpp
• cpp/include/raft/stats/detail/cov.cuh
• cpp/include/raft/matrix/argmax.cuh
• cpp/include/raft/core/copy.cuh
• docs/source/cpp_api/core_resources.rst
• cpp/include/raft/matrix/slice.cuh
• cpp/include/raft/matrix/diagonal.cuh
• cpp/include/raft/matrix/detail/math.cuh
• cpp/tests/sparse/reduce.cu
• cpp/include/raft/core/bitset.cuh
• cpp/include/raft/matrix/detail/scatter_inplace.cuh
• cpp/include/raft/random/sample_without_replacement.cuh
• cpp/include/raft/linalg/detail/map.cuh
• cpp/tests/linalg/power.cu
• cpp/tests/linalg/eig.cu
• cpp/include/raft/sparse/linalg/detail/sddmm.hpp
• cpp/include/raft/linalg/unary_op.cuh
• cpp/include/raft/matrix/detail/select_warpsort.cuh
• .coderabbit.yaml
• cpp/include/raft/matrix/triangular.cuh
• cpp/tests/core/mdbuffer.cu
• cpp/include/raft/random/make_blobs.cuh
• cpp/include/raft/matrix/col_wise_sort.cuh
• cpp/include/raft/sparse/convert/detail/csr.cuh
• cpp/include/raft/linalg/detail/reduce.cuh
• cpp/include/raft/stats/detail/kl_divergence.cuh
• cpp/include/raft/linalg/multiply.cuh
• cpp/include/raft/stats/rand_index.cuh
• cpp/include/raft/matrix/power.cuh
• cpp/include/raft/core/host_mdarray.hpp
• cpp/include/raft/spectral/detail/partition.hpp
• cpp/tests/stats/stddev.cu
• cpp/tests/sparse/preprocess.cu
• cpp/include/raft/matrix/threshold.cuh
• cpp/include/raft/matrix/linewise_op.cuh
• cpp/tests/core/temporary_device_buffer.cu
• cpp/include/raft/core/managed_mdarray.hpp
• cpp/include/raft/stats/information_criterion.cuh
• cpp/include/raft/matrix/reciprocal.cuh
• cpp/include/raft/linalg/dot.cuh
• cpp/tests/linalg/subtract.cu
• cpp/include/raft/core/pinned_mdarray.hpp
• cpp/tests/linalg/ternary_op.cu
• cpp/include/raft/stats/mean_center.cuh
• cpp/include/raft/matrix/argmin.cuh
• cpp/include/raft/linalg/detail/cholesky_r1_update.cuh
• cpp/tests/sparse/solver/lanczos.cu
• cpp/include/raft/core/detail/copy.hpp
• cpp/include/raft/linalg/detail/svd.cuh
• cpp/tests/core/bitmap.cu
• cpp/tests/sparse/norm.cu
• cpp/include/raft/core/resource/cuda_stream.hpp
• cpp/tests/linalg/transpose.cu
• cpp/tests/linalg/normalize.cu
• cpp/include/raft/random/detail/make_blobs.cuh
• cpp/include/raft/core/coo_matrix.hpp
• cpp/include/raft/linalg/norm.cuh
• cpp/include/raft/linalg/subtract.cuh
• cpp/include/raft/random/detail/multi_variable_gaussian.cuh
• cpp/include/raft/stats/detail/neighborhood_recall.cuh
• cpp/include/raft/sparse/convert/detail/adj_to_csr.cuh
• cpp/include/raft/stats/detail/dispersion.cuh
• cpp/include/raft/sparse/op/sort.cuh
• cpp/include/raft/random/detail/rng_impl_deprecated.cuh
• cpp/tests/util/memory_type_dispatcher.cu
• cpp/tests/core/mdarray.cu
• cpp/tests/linalg/matrix_vector_op.cu
• cpp/tests/linalg/axpy.cu
• cpp/include/raft/linalg/detail/cublaslt_wrappers.hpp
• cpp/tests/stats/meanvar.cu
• cpp/include/raft/stats/accuracy.cuh
• cpp/include/raft/linalg/detail/matrix_vector_op.cuh
• cpp/include/raft/stats/v_measure.cuh
• cpp/include/raft/matrix/gather.cuh
• cpp/include/raft/stats/completeness_score.cuh
• cpp/tests/sparse/select_k_csr.cu
• cpp/include/raft/core/resource/resource_types.hpp
• cpp/tests/sparse/symmetrize.cu
• cpp/tests/linalg/multiply.cu
• cpp/include/raft/stats/detail/rand_index.cuh
• cpp/include/raft/stats/entropy.cuh
• cpp/include/raft/spectral/detail/spectral_util.cuh
• cpp/include/raft/linalg/reduce.cuh
• cpp/include/raft/stats/adjusted_rand_index.cuh
• cpp/tests/core/mdspan_copy.cu
• cpp/tests/linalg/dot.cu
• cpp/include/raft/linalg/detail/coalesced_reduction-inl.cuh
• cpp/include/raft/stats/detail/contingencyMatrix.cuh
• cpp/include/raft/matrix/reverse.cuh
• cpp/include/raft/stats/kl_divergence.cuh
• cpp/tests/sparse/sddmm.cu
• cpp/include/raft/stats/homogeneity_score.cuh
• cpp/include/raft/stats/regression_metrics.cuh
• cpp/include/raft/sparse/op/detail/reduce.cuh
• cpp/include/raft/linalg/detail/transpose.cuh
• cpp/include/raft/linalg/detail/pca.cuh
• cpp/tests/stats/cov.cu
• cpp/include/raft/random/permute.cuh
• cpp/include/raft/linalg/detail/gemv.hpp
• cpp/include/raft/stats/contingency_matrix.cuh
• cpp/include/raft/linalg/detail/rsvd.cuh
• cpp/include/raft/stats/detail/mean_center.cuh
• cpp/include/raft/stats/mutual_info_score.cuh
• cpp/include/raft/matrix/print.cuh
• cpp/include/raft/linalg/coalesced_reduction.cuh
• cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh
• cpp/include/raft/core/device_mdarray.hpp
• cpp/include/raft/stats/histogram.cuh
• cpp/tests/sparse/masked_matmul.cu
• cpp/include/raft/stats/weighted_mean.cuh
• cpp/tests/sparse/diagonal.cu
• cpp/tests/stats/mean.cu
• cpp/include/raft/random/detail/make_regression.cuh
• cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh
• cpp/tests/core/mdspan_utils.cu
• cpp/include/raft/stats/mean.cuh
• cpp/include/raft/sparse/matrix/preprocessing.cuh
• cpp/include/raft/label/classlabels.cuh
• cpp/include/raft/stats/r2_score.cuh
• cpp/include/raft/sparse/linalg/detail/spmm.hpp
• cpp/tests/sparse/filter.cu
• cpp/tests/util/popc.cu
• cpp/tests/linalg/add.cu
• cpp/include/raft/sparse/solver/detail/cholesky_qr.cuh
• cpp/tests/linalg/gemm_basic.cpp
• cpp/include/raft/stats/meanvar.cuh
• cpp/include/raft/sparse/linalg/detail/symmetrize.cuh
• cpp/include/raft/matrix/detail/select_radix.cuh
• cpp/include/raft/stats/detail/mutual_info_score.cuh
• cpp/include/raft/matrix/copy.cuh
• cpp/include/raft/core/device_container_policy.hpp
• cpp/include/raft/stats/minmax.cuh
• cpp/tests/linalg/rsvd.cu
• cpp/include/raft/linalg/add.cuh
• docs/source/developer_guide.md
• cpp/tests/stats/minmax.cu
• cpp/include/raft/matrix/detail/gather.cuh
• cpp/tests/stats/sum.cu
• cpp/tests/linalg/unary_op.cu
• cpp/tests/linalg/map.cu
• cpp/tests/linalg/strided_reduction.cu
• cpp/include/raft/random/detail/rmat_rectangular_generator.cuh
• cpp/include/raft/stats/sum.cuh
• cpp/tests/linalg/norm.cu
• cpp/tests/core/bitset.cu
• docs/source/dry_run_protocol.md
• cpp/include/raft/stats/detail/batched/silhouette_score.cuh
• cpp/include/raft/sparse/matrix/detail/preprocessing.cuh
• cpp/include/raft/stats/detail/entropy.cuh
• cpp/include/raft/stats/neighborhood_recall.cuh
• cpp/include/raft/linalg/detail/lstsq.cuh
• cpp/tests/linalg/pca.cu
• cpp/include/raft/stats/detail/homogeneity_score.cuh
• cpp/include/raft/core/bitset.hpp
• cpp/tests/matrix/sample_rows.cu
• cpp/tests/linalg/reduce.cu
• cpp/include/raft/matrix/detail/columnWiseSort.cuh
• cpp/tests/linalg/mean_squared_error.cu
• cpp/include/raft/sparse/op/detail/filter.cuh
• cpp/include/raft/stats/detail/trustworthiness_score.cuh
• cpp/include/raft/sparse/solver/detail/randomized_svds.cuh
• cpp/tests/sparse/csr_transpose.cu
• cpp/include/raft/label/detail/classlabels.cuh
• cpp/include/raft/core/csr_matrix.hpp
• cpp/tests/CMakeLists.txt
• cpp/tests/linalg/svd.cu
• cpp/tests/core/sparse_matrix.cu
• cpp/tests/linalg/divide.cu
• cpp/include/raft/matrix/sqrt.cuh
• cpp/include/raft/matrix/norm.cuh
• cpp/include/raft/stats/detail/mean.cuh
• cpp/include/raft/stats/dispersion.cuh
• cpp/include/raft/stats/stddev.cuh
• cpp/include/raft/linalg/detail/norm.cuh
• cpp/include/raft/stats/detail/adjusted_rand_index.cuh
• cpp/include/raft/stats/detail/silhouette_score.cuh
• cpp/tests/stats/weighted_mean.cu
• cpp/include/raft/mr/dry_run_resource.hpp
• cpp/include/raft/linalg/detail/qr.cuh
• cpp/include/raft/sparse/solver/detail/lanczos.cuh
• cpp/tests/matrix/gather.cu
• cpp/tests/linalg/gemm_layout.cu
• cpp/include/raft/sparse/linalg/spmm.hpp
• cpp/include/raft/linalg/detail/eig.cuh
• cpp/include/raft/sparse/op/detail/sort.h
• cpp/include/raft/core/resource/dry_run_flag.hpp
• cpp/include/raft/solver/detail/lap_functions.cuh
• cpp/include/raft/random/detail/rng_impl.cuh
• cpp/include/raft/random/rng.cuh
• cpp/include/raft/stats/detail/scores.cuh
• cpp/include/raft/sparse/solver/detail/svds_sign_correction.cuh
• cpp/include/raft/stats/detail/weighted_mean.cuh
• cpp/tests/test_utils.cuh
• cpp/include/raft/spectral/detail/matrix_wrappers.hpp
• cpp/include/raft/stats/detail/meanvar.cuh
• cpp/tests/sparse/spmm.cu
• cpp/include/raft/stats/detail/stddev.cuh
• cpp/include/raft/linalg/normalize.cuh
• cpp/include/raft/sparse/linalg/detail/laplacian.cuh
• cpp/tests/sparse/convert_csr.cu
• cpp/include/raft/stats/detail/v_measure.cuh
• cpp/include/raft/spectral/detail/modularity_maximization.hpp
• cpp/tests/stats/homogeneity_score.cu

[coderabbitai]

🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win

Reset bytesize_ when you release the old buffer.

Line 122 deallocates the buffer and Line 123 clears data_, but bytesize_ keeps the old value. If the allocation inside host_container{count, mr_} throws, *this is left with bytesize_ > 0 and data_ == nullptr. The destructor is safe because it checks both fields, but resize() then computes cur_count from the stale bytesize_ and calls std::copy from a null pointer.

🛠️ Proposed fix

     if (data_ != nullptr) {
       mr_.deallocate_sync(data_, bytesize_);
       data_ = nullptr;
     }
+    bytesize_ = 0;
     auto tmp = host_container{count, mr_};
     std::swap(tmp, *this);

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change

	void reallocate(size_type count)
	{
	if (bytesize_ == sizeof(value_type) * count) { return; }
	if (data_ != nullptr) {
	mr_.deallocate_sync(data_, bytesize_);
	data_ = nullptr;
	}
	auto tmp = host_container{count, mr_};
	std::swap(tmp, *this);
	}
	void reallocate(size_type count)
	{
	if (bytesize_ == sizeof(value_type) * count) { return; }
	if (data_ != nullptr) {
	mr_.deallocate_sync(data_, bytesize_);
	data_ = nullptr;
	}
	bytesize_ = 0;
	auto tmp = host_container{count, mr_};
	std::swap(tmp, *this);
	}

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@cpp/include/raft/core/host_container_policy.hpp` around lines 118 - 127,
Update host_container::reallocate so bytesize_ is reset to zero immediately when
the existing data_ buffer is deallocated and data_ is cleared, before
constructing host_container{count, mr_}; preserve the existing allocation and
swap behavior.