[BUG] `cuvsNNDescentBuild` fails for a high-dimensional float32 dataset passed through the cuVS C API as a DLPack host tensor.

[tkcaccia]

Describe the bug
Failing shape and parameters:

COIL20 dataset

• n_rows = 1440
• n_features = 16384
• dtype = float32
• DLPack device: kDLCPU
• Layout: contiguous row-major, strides = nullptr
• Metric: L2Expanded
• graph_degree = 10
• intermediate_graph_degree = 20
• max_iterations = 20
• return_distances = true

A lower-dimensional float32 dataset, for example 70000 x 784, works through the same C API path.

Actual Error

CUDA error encountered at:
cpp/src/neighbors/detail/nn_descent.cuh line=1564
call='cudaPeekAtLastError()'
Reason=cudaErrorInvalidValue: invalid argument

Expected behavior
cuvsNNDescentBuild succeeds for this valid float32 host tensor shap

Environment details (please complete the following information):

• GPU: NVIDIA GeForce RTX 5060 Ti, 16 GB
• Driver: 595.71.05
• CUDA: 13.2
• libcuvs: 26.06.00, CUDA 13 build
• libraft: 26.06.00, CUDA 13 build
• OS: Linux x86_64

Observed Behavior
The build fails inside cuVS with cudaErrorInvalidValue after cudaPeekAtLastError() in nn_descent.cuh.

Additional context
The same installation can run cuVS NN-descent on lower-dimensional float32 input such as 70000 x 784, so the failure appears related to high feature count rather than general CUDA/cuVS setup.

Steps/Code to reproduce bug

#include <cstdint>
#include <vector>

#include <dlpack/dlpack.h>
#include <cuvs/core/c_api.h>
#include <cuvs/neighbors/nn_descent.h>

static DLManagedTensor make_tensor(void* data, int64_t* shape) {
  DLManagedTensor tensor{};
  tensor.dl_tensor.data = data;
  tensor.dl_tensor.device.device_type = kDLCPU;
  tensor.dl_tensor.device.device_id = 0;
  tensor.dl_tensor.ndim = 2;
  tensor.dl_tensor.dtype.code = kDLFloat;
  tensor.dl_tensor.dtype.bits = 32;
  tensor.dl_tensor.dtype.lanes = 1;
  tensor.dl_tensor.shape = shape;
  tensor.dl_tensor.strides = nullptr;
  tensor.dl_tensor.byte_offset = 0;
  tensor.manager_ctx = nullptr;
  tensor.deleter = nullptr;
  return tensor;
}

int main() {
  const int64_t n_rows = 1440;
  const int64_t n_features = 16384;

  std::vector<float> x(static_cast<size_t>(n_rows) * n_features);
  // Fill x with finite float32 values.

  int64_t dataset_shape[2] = {n_rows, n_features};
  DLManagedTensor dataset_tensor = make_tensor(x.data(), dataset_shape);

  cuvsResources_t res;
  cuvsResourcesCreate(&res);

  cuvsNNDescentIndexParams_t params;
  cuvsNNDescentIndexParamsCreate(&params);
  params->metric = L2Expanded;
  params->graph_degree = 10;
  params->intermediate_graph_degree = 20;
  params->max_iterations = 20;
  params->return_distances = true;

  cuvsNNDescentIndex_t index;
  cuvsNNDescentIndexCreate(&index);

  auto status = cuvsNNDescentBuild(
    res,
    params,
    &dataset_tensor,
    nullptr,
    index
  );

  cuvsNNDescentIndexDestroy(index);
  cuvsNNDescentIndexParamsDestroy(params);
  cuvsResourcesDestroy(res);

  return static_cast<int>(status);
}

[tkcaccia]

cuVS NN-descent FP32 L2 build can fail with cudaErrorInvalidValue for high-dimensional inputs unless the L2-norm kernel opts into larger dynamic shared memory

Summary

While testing direct RAPIDS cuVS NN-descent through the cuVS C API, a
high-dimensional FP32 Euclidean/L2 dataset failed during cuvsNNDescentBuild
with cudaErrorInvalidValue: invalid argument. The failure happened before any
nearest-neighbour result was produced.

The failing shape was COIL20:

n = 1440
dim = 16384
k = 10
graph_degree = 10
intermediate_graph_degree = 20
max_iterations = 1
input dtype = float32
metric = L2Expanded

The same build path succeeds after patching cuVS to opt the
compute_l2_norms_kernel launch into the larger device-supported dynamic
shared-memory limit with
cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize, ...).

Environment

Observed with:

cuVS/libcuvs: 26.06.00
Source tag used for local rebuild: NVIDIA/cuVS v26.06.00
Source commit: 2bd7cd71d31e39eb9ee00c2a250085a1afb84977
CUDA runtime/toolkit family: CUDA 13.x
API path: cuVS C API, cuvsNNDescentBuild
GPU tested: NVIDIA GPU with cudaDevAttrMaxSharedMemoryPerBlockOptin = 101376

The failure is not specific to R. A standalone C/C++ repro using the cuVS C API
with the shape above produced the same cuVS error. faissR only exposed the
error because it calls the cuVS C API directly.

Actual Error

The unpatched cuVS library failed with:

cuvsNNDescentBuild failed: CUDA error encountered at:
cpp/src/neighbors/detail/nn_descent.cuh line=1564
call='cudaPeekAtLastError()'
Reason=cudaErrorInvalidValue: invalid argument

The line number can vary slightly by build, but the stack points to
neighbors/detail/nn_descent.cuh immediately after launching the L2-norm kernel.

Expected Behavior

If the GPU supports enough opt-in dynamic shared memory, cuVS should set the
kernel attribute and launch successfully.

If the GPU cannot support the required dynamic shared memory, cuVS should fail
with a clear error explaining the required bytes and the device limit, rather
than surfacing a raw cudaErrorInvalidValue after launch.

Diagnosis

For FP32 L2 input with dim = 16384, the L2-norm kernel dynamic shared-memory
request is:

ceil(dim / warp_size) * warp_size * sizeof(float)
= ceil(16384 / 32) * 32 * 4
= 65536 bytes

That is larger than CUDA's usual default dynamic shared-memory launch limit
around 48 KiB per block. The test GPU reports an opt-in maximum of 101376
bytes, so the launch is legal on the hardware if cuVS opts in before launching
the kernel.

The local patch confirmed that the missing opt-in is the issue.

Local Patch That Fixed The Reproducer

In cpp/src/neighbors/detail/nn_descent.cuh, add a helper equivalent to:

template <typename Data_t>
size_t compute_l2_norms_smem_size(size_t dim)
{
  return sizeof(Data_t) *
         raft::ceildiv(dim, static_cast<size_t>(raft::warp_size())) *
         static_cast<size_t>(raft::warp_size());
}

template <typename Data_t>
size_t configure_compute_l2_norms_kernel_smem(size_t dim)
{
  size_t smem = compute_l2_norms_smem_size<Data_t>(dim);

  int dev_id = 0;
  RAFT_CUDA_TRY(cudaGetDevice(&dev_id));

  int max_smem = 0;
  RAFT_CUDA_TRY(cudaDeviceGetAttribute(
    &max_smem,
    cudaDevAttrMaxSharedMemoryPerBlockOptin,
    dev_id
  ));

  RAFT_EXPECTS(
    smem <= static_cast<size_t>(max_smem),
    "NN_DESCENT L2 norm kernel requires %zu bytes of shared memory "
    "(dim=%zu), but the device supports at most %d bytes per block.",
    smem,
    dim,
    max_smem
  );

  auto kernel = compute_l2_norms_kernel<Data_t>;
  RAFT_CUDA_TRY(cudaFuncSetAttribute(
    kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    static_cast<int>(smem)
  ));

  return smem;
}

Then replace the raw dynamic shared-memory expression used for
compute_l2_norms_kernel<<<...>>> with:

auto smem =
  configure_compute_l2_norms_kernel_smem<float>(build_config_.dataset_dim);

or, in the templated direct-data path:

auto smem =
  configure_compute_l2_norms_kernel_smem<input_t>(build_config_.dataset_dim);

The kernel launch should then use smem as the dynamic shared-memory argument.

Validation After Patch

After rebuilding libcuvs.so and libcuvs_c.so with the local patch and
preloading the patched libraries before the conda-installed cuVS libraries, the
same calls succeeded:

COIL20 full matrix: 1440 x 16384, k = 10
backend: cuda_cuvs_nndescent
elapsed: 0.696 s
result: OK, indices 1440 x 10, distances 1440 x 10

MNIST70k float32: 70000 x 784, k = 10
backend: cuda_cuvs_nndescent
elapsed: 1.162 s
result: OK, indices 70000 x 10, distances 70000 x 10

A standalone cuVS C/C++ reproducer for the COIL20-like shape also returned:

shape=1440x16384 k=10 graph_degree=10 intermediate=20 iters=1
max_optin_smem=101376
OK

Request

Please consider adding the dynamic shared-memory opt-in around the
NN-descent L2-norm kernel launch, or an equivalent implementation that:

1. Computes the dynamic shared-memory requirement from the dataset dimension and
   input type.
2. Queries cudaDevAttrMaxSharedMemoryPerBlockOptin.
3. Calls cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...)
   when the requirement exceeds the default limit but is within the opt-in
   device limit.
4. Raises a clear cuVS/RAFT error if the requirement exceeds the hardware limit.

This should make direct cuVS NN-descent reliable for high-dimensional FP32 L2
datasets on GPUs that support the required opt-in dynamic shared memory.

[jinsolp]

Hi @tkcaccia , thanks for sharing the investigation and patch!

Would you be open to contributing this patch by opening a PR?

Thank you : )