Fall back to cudaMalloc on devices without stream-ordered allocation

cpp/src/neighbors/detail/cagra/cagra_filter_payload.hpp

@@ -10,6 +10,7 @@
 
 #include <cuvs/core/bloom_filter.hpp>
 #include <raft/core/error.hpp>
+#include <util/stream_alloc_compat.hpp>  // malloc_async_compat
 
 #include <cuda_runtime_api.h>
 
@@ -53,7 +54,7 @@ struct cagra_device_payload_owner {
     ~state() noexcept
     {
       if (device_payload != nullptr) {
-        RAFT_CUDA_TRY_NO_THROW(cudaFreeAsync(device_payload, stream));
+        RAFT_CUDA_TRY_NO_THROW(cuvs::util::free_async_compat(device_payload, stream));
       }
       if (ready_event != nullptr) { RAFT_CUDA_TRY_NO_THROW(cudaEventDestroy(ready_event)); }
     }
@@ -63,7 +64,7 @@ struct cagra_device_payload_owner {
       std::lock_guard<std::mutex> lock(mutex);
       if (device_payload == nullptr) {
         RAFT_CUDA_TRY(cudaGetDevice(&device));
-        RAFT_CUDA_TRY(cudaMallocAsync(
+        RAFT_CUDA_TRY(cuvs::util::malloc_async_compat(
           reinterpret_cast<void**>(&device_payload), sizeof(PayloadT), cuda_stream));
         RAFT_CUDA_TRY(cudaMemcpyAsync(
           device_payload, &host_payload, sizeof(PayloadT), cudaMemcpyHostToDevice, cuda_stream));

cpp/src/neighbors/detail/cagra/compute_distance.hpp

@@ -13,6 +13,7 @@
 #include <cuvs/neighbors/cagra.hpp>
 #include <cuvs/neighbors/common.hpp>
 #include <raft/core/logger.hpp>
+#include <util/stream_alloc_compat.hpp>  // malloc_async_compat
 #include <raft/core/operators.hpp>
 
 // TODO: This shouldn't be invoking spatial/knn
@@ -230,7 +231,7 @@ struct dataset_descriptor_host {
     {
       if (std::holds_alternative<ready_t>(value)) {
         auto& [ptr, stream] = std::get<ready_t>(value);
-        RAFT_CUDA_TRY_NO_THROW(cudaFreeAsync(ptr, stream));
+        RAFT_CUDA_TRY_NO_THROW(cuvs::util::free_async_compat(ptr, stream));
       }
       RAFT_CUDA_TRY_NO_THROW(cudaEventDestroy(ready_event));
     }
@@ -241,7 +242,8 @@ struct dataset_descriptor_host {
       if (std::holds_alternative<init_f>(value)) {
         auto& [fun, size]     = std::get<init_f>(value);
         dev_descriptor_t* ptr = nullptr;
-        RAFT_CUDA_TRY(cudaMallocAsync(&ptr, size, stream));
+        RAFT_CUDA_TRY(
+          cuvs::util::malloc_async_compat(reinterpret_cast<void**>(&ptr), size, stream));
         fun(ptr, stream);
         RAFT_CUDA_TRY(cudaEventRecord(ready_event, stream));
         value = std::make_tuple(ptr, stream);

cpp/src/util/stream_alloc_compat.hpp

@@ -0,0 +1,78 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+/**
+ * @file stream_alloc_compat.hpp
+ * @brief Stream-ordered allocation with a fallback for devices without memory-pool support.
+ *
+ * `cudaMallocAsync` / `cudaFreeAsync` are backed by the CUDA stream-ordered memory allocator, which
+ * is optional hardware/driver functionality: on devices that report
+ * `cudaDevAttrMemoryPoolsSupported == 0` -- Maxwell-generation GPUs among them -- every call fails
+ * with `cudaErrorNotSupported`.
+ *
+ * The few places in cuVS that allocate directly from a stream rather than through RMM therefore go
+ * through the helpers below, which fall back to plain `cudaMalloc` / `cudaFree`. The fallback is
+ * semantically stronger than what the caller asked for, not weaker:
+ *
+ *   - `cudaMalloc` is synchronous with respect to the whole device, so the returned memory is
+ *     trivially usable by any work subsequently issued into the stream, which is exactly the
+ *     guarantee `cudaMallocAsync` gives for that stream;
+ *   - `cudaFree` blocks until all previously issued device work has finished, which subsumes
+ *     `cudaFreeAsync`'s "free once the preceding work in this stream completes".
+ *
+ * The price is the loss of pool reuse and the implicit synchronisation, so these helpers are only
+ * appropriate for allocations that happen a handful of times per search, not per launch.
+ */
+
+#pragma once
+
+#include <raft/util/cuda_rt_essentials.hpp>  // RAFT_CUDA_TRY
+
+#include <cuda_runtime.h>
+
+#include <cstddef>
+
+namespace cuvs::util {
+
+/** Whether the current device supports the CUDA stream-ordered memory allocator. */
+inline bool device_supports_memory_pools()
+{
+  int device = 0;
+  RAFT_CUDA_TRY(cudaGetDevice(&device));
+  int supported = 0;
+  RAFT_CUDA_TRY(cudaDeviceGetAttribute(&supported, cudaDevAttrMemoryPoolsSupported, device));
+  return supported != 0;
+}
+
+/**
+ * @brief `cudaMallocAsync`, or `cudaMalloc` where memory pools are unsupported.
+ *
+ * @param[out] ptr    allocated device pointer
+ * @param[in]  size   allocation size in bytes
+ * @param[in]  stream stream the allocation is ordered against
+ */
+inline cudaError_t malloc_async_compat(void** ptr, std::size_t size, cudaStream_t stream)
+{
+  if (device_supports_memory_pools()) { return cudaMallocAsync(ptr, size, stream); }
+  return cudaMalloc(ptr, size);
+}
+
+/**
+ * @brief `cudaFreeAsync`, or `cudaFree` where memory pools are unsupported.
+ *
+ * @param[in] ptr    pointer previously returned by @ref malloc_async_compat
+ * @param[in] stream stream the deallocation is ordered against
+ */
+inline cudaError_t free_async_compat(void* ptr, cudaStream_t stream)
+{
+  if (device_supports_memory_pools()) { return cudaFreeAsync(ptr, stream); }
+  // `cudaFree` is not stream-ordered, so make sure the work that may still be reading this
+  // allocation has completed before the memory is handed back.
+  auto status = cudaStreamSynchronize(stream);
+  if (status != cudaSuccess) { return status; }
+  return cudaFree(ptr);
+}
+
+}  // namespace cuvs::util