Fix NN-descent L2 norm shared memory opt-in

[tkcaccia]

Fixes #2267.

Summary

This PR fixes a CUDA launch failure in cuVS NN-descent when the L2 norm kernel requires more dynamic shared memory than the default per-block launch limit.

For high-dimensional FP32 L2 input, for example a 1440 x 16384 host DLPack tensor passed through the cuVS C API, cuvsNNDescentBuild can fail after launching compute_l2_norms_kernel with:

cudaErrorInvalidValue: invalid argument
cpp/src/neighbors/detail/nn_descent.cuh
call='cudaPeekAtLastError()'

The dynamic shared-memory request for dim = 16384 and float input is 65536 bytes, which is larger than the usual default dynamic shared-memory launch limit but within the opt-in limit on the tested GPU.

Change

The patch adds a small helper around compute_l2_norms_kernel that:

• computes the dynamic shared-memory requirement from the dataset dimension and input type;
• queries cudaDevAttrMaxSharedMemoryPerBlockOptin for the current device;
• raises a clear RAFT/cuVS error if the request exceeds the device opt-in limit;
• calls cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, ...) before launching the kernel.

The helper is used at both NN-descent L2 norm launch sites.

Validation

Locally in this PR checkout:

git diff --check

I could not build CUDA/cuVS on the local macOS machine used to prepare the PR. The same fix was previously validated by rebuilding cuVS 26.06.00 on a Linux CUDA machine and rerunning the reproducer described in #2267:

COIL20-shaped input: 1440 x 16384 float32, k = 10
Result: cuvsNNDescentBuild returned OK

MNIST70k input: 70000 x 784 float32, k = 10
Result: cuVS NN-descent returned OK

The change does not alter the NN-descent algorithm and does not introduce a fallback to another method.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[tkcaccia]

Thanks for the automated checks. I do not have permission to add labels on this repository. Could a maintainer please add the required Label Checker labels?

Suggested labels:

• bug as the category label
• non-breaking as the breaking-status label

This PR fixes the NN-descent CUDA launch failure reported in #2267 and does not change public API or intended behavior.

[jinsolp]

/ok to test 5ec35c4

[jinsolp]

Thanks for the PR @tkcaccia ! It looks like the style checks are failing. Could you run the pre-commit hooks? It'll fix the style for you.