diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index a6438c1a8d..aad5c1207e 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -11,450 +11,18 @@ jobs: pr-builder: needs: - build-details - - check-nightly-ci - - changed-files - - checks - - check-c-abi - - security-suite - conda-cpp-build - conda-cpp-tests - - conda-cpp-checks - - conda-python-build - - conda-python-tests - - rocky8-clib-standalone-build - - rocky8-clib-tests - - conda-java-build - - conda-java-tests - - conda-lucene-build - - conda-lucene-tests - - rust-build - - go-build - - docs-build - - docs-build-matrix - - wheel-build-libcuvs - - wheel-build-cuvs - - wheel-tests-cuvs - - devcontainer - - telemetry-setup uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@main permissions: contents: read if: always() with: needs: ${{ toJSON(needs) }} - telemetry-setup: - continue-on-error: true - runs-on: ubuntu-latest - permissions: - contents: read - env: - OTEL_SERVICE_NAME: 'pr-cuvs' - steps: - - name: Telemetry setup - if: ${{ vars.TELEMETRY_ENABLED == 'true' }} - uses: rapidsai/shared-actions/telemetry-dispatch-stash-base-env-vars@main build-details: uses: rapidsai/shared-workflows/.github/workflows/compute-build-details.yaml@main - check-nightly-ci: - runs-on: ubuntu-latest - permissions: - actions: read - contents: read - id-token: write - pull-requests: read - env: - GH_TOKEN: ${{ github.token }} - steps: - - name: Get PR Info - id: get-pr-info - uses: nv-gha-runners/get-pr-info@main - - name: Check if nightly CI is passing - uses: rapidsai/shared-actions/check_nightly_success/dispatch@main - with: - repo: ${{ github.repository }} - target-branch: ${{ fromJSON(steps.get-pr-info.outputs.pr-info).base.ref }} - max-days-without-success: 30 - changed-files: - needs: telemetry-setup - permissions: - actions: read - contents: read - packages: read - pull-requests: read - uses: rapidsai/shared-workflows/.github/workflows/changed-files.yaml@main - with: - files_yaml: | - build_docs: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!Dockerfile' - - '!SECURITY.md' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run*.sh' - - '!ci/test*.sh' - - '!ci/validate_wheel.sh' - - '!thirdparty/LICENSES/**' - test_cpp: - - '.github/workflows/pr.yaml' - - 'RAPIDS_BRANCH' - - 'VERSION' - - 'build.sh' - - 'c/**' - - 'ci/build_cpp.sh' - - 'ci/get_test_data.sh' - - 'ci/test_cpp.sh' - - 'cmake/**' - - 'conda/recipes/libcuvs/**' - - 'cpp/**' - - 'dependencies.yaml' - - '!**/*/REVIEW_GUIDELINES.md' - - '!cpp/doxygen/**' - test_standalone_c: - - '.github/workflows/pr.yaml' - - 'RAPIDS_BRANCH' - - 'VERSION' - - 'build.sh' - - 'c/**' - - 'ci/standalone_c/**' - - 'ci/test_standalone_c.sh' - - 'cmake/**' - - 'cpp/**' - - 'dependencies.yaml' - - '!**/*/REVIEW_GUIDELINES.md' - - '!cpp/doxygen/**' - test_java: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_wheel*.sh' - - '!ci/checkstyle.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_python.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!conda/environments/**' - - '!conda/recipes/cuvs/**' - - '!conda/recipes/cuvs-bench/**' - - '!conda/recipes/cuvs-bench-cpu/**' - - '!docs/**' - - '!examples/.clang-format' - - '!examples/.clang-tidy' - - '!examples/.clangd' - - '!examples/README.md' - - '!examples/build.sh' - - '!examples/c/**' - - '!examples/cmake/**' - - '!examples/cpp/**' - - '!examples/go/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!notebooks/**' - - '!pyproject.toml' - - '!python/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_python_conda: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_rust.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_standalone_c.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!cpp/.clang-format' - - '!cpp/.clang-tidy' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_python_wheels: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!ci/build_cpp.sh' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_standalone_c.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_standalone_c.sh' - - '!ci/validate_wheel.sh' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!conda/**' - - '!cpp/.clang-format' - - '!cpp/.clang-tidy' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!java/**' - - '!go/**' - - '!img/**' - - '!notebooks/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - test_rust: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_go.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!docs/**' - - '!examples/**' - - '!fern/**' - - '!go/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!python/**' - - '!thirdparty/LICENSES/**' - test_go: - - '**' - - '!**/*/REVIEW_GUIDELINES.md' - - '!.coderabbit.yaml' - - '!.devcontainer/**' - - '!.github/CODEOWNERS' - - '!.github/ISSUE_TEMPLATE/**' - - '!.github/PULL_REQUEST_TEMPLATE.md' - - '!.github/copy-pr-bot.yaml' - - '!.github/labeler.yml' - - '!.github/ops-bot.yaml' - - '!.github/release.yml' - - '!.github/workflows/build.yaml' - - '!.github/workflows/check-c-abi.yaml' - - '!.github/workflows/labeler.yml' - - '!.github/workflows/publish-rust.yaml' - - '!.github/workflows/security-suite.yml' - - '!.github/workflows/store-c-abi-baseline.yaml' - - '!.github/workflows/test.yaml' - - '!.github/workflows/trigger-breaking-change-alert.yaml' - - '!.github/workflows/update-c-abi-baseline.yaml' - - '!.pre-commit-config.yaml' - - '!.shellcheckrc' - - '!.yamllint.yaml' - - '!CODEOWNERS' - - '!Dockerfile' - - '!README.md' - - '!SECURITY.md' - - '!ci/build_docs.sh' - - '!ci/build_java.sh' - - '!ci/build_lucene.sh' - - '!ci/build_python.sh' - - '!ci/build_rust.sh' - - '!ci/build_wheel*.sh' - - '!ci/check_style.sh' - - '!ci/release/update-version.sh' - - '!ci/run_ctests.sh' - - '!ci/run_cuvs_pytests.sh' - - '!ci/test_cpp.sh' - - '!ci/test_java.sh' - - '!ci/test_lucene.sh' - - '!ci/test_python.sh' - - '!ci/test_standalone_c.sh' - - '!ci/test_wheel_cuvs.sh' - - '!ci/validate_wheel.sh' - - '!docs/**' - - '!examples/.clang-format' - - '!examples/.clang-tidy' - - '!examples/.clangd' - - '!examples/README.md' - - '!examples/build.sh' - - '!examples/c/**' - - '!examples/cmake/**' - - '!examples/cpp/**' - - '!examples/java/**' - - '!fern/**' - - '!img/**' - - '!java/**' - - '!notebooks/**' - - '!python/**' - - '!rust/**' - - '!thirdparty/LICENSES/**' - checks: - needs: telemetry-setup - uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@main - permissions: - contents: read - with: - enable_check_generated_files: false - ignored_pr_jobs: "telemetry-summarize" - check-c-abi: - needs: telemetry-setup - permissions: - actions: read - contents: read - issues: write - pull-requests: read - uses: ./.github/workflows/check-c-abi.yaml - security-suite: - permissions: - actions: read - contents: read - id-token: write # OIDC -> Vault -> nvcr.io image pull - security-events: write # publish redacted SARIF to code scanning - uses: ./.github/workflows/security-suite.yml conda-cpp-build: - needs: [build-details, checks] + needs: [build-details] permissions: actions: read contents: read @@ -468,8 +36,9 @@ jobs: build-datetime: ${{ needs.build-details.outputs.build-datetime }} node_type: cpu16 script: ci/build_cpp.sh + matrix_filter: map(select(.ARCH == "arm64" and (.CUDA_VER | startswith("12.")))) conda-cpp-tests: - needs: [conda-cpp-build, changed-files] + needs: [conda-cpp-build] permissions: actions: read contents: read @@ -478,16 +47,11 @@ jobs: pull-requests: read secrets: inherit # zizmor: ignore[secrets-inherit] uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Split the C++ tests across several jobs for each CUDA/GPU configuration - strategy: - fail-fast: false - matrix: - shard: [1, 2, 3, 4] - num_shards: [4] + if: true with: build_type: pull-request - script: "ci/test_cpp.sh ${{ matrix.shard }} ${{ matrix.num_shards }}" + script: "ci/test_cpp.sh" + matrix_filter: 'map(select(.ARCH == "arm64" and .GPU == "a100")) | [range(4) as $i | .[] | . + {COPY: $i}]' cache-paths: .cache/cuda-jit cache-key-prefix: libcuvs-cudajit-v1 # Caches are immutable. Re-key when the JIT-relevant sources change so a completed run can @@ -518,279 +82,3 @@ jobs: cache-read-only: true cache-environment: | CUDA_CACHE_PATH=.cache/cuda-jit - conda-cpp-checks: - needs: conda-cpp-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@main - with: - build_type: pull-request - symbol_exclusions: (void (thrust::|cub::)) - package_name: libcuvs - conda-python-build: - needs: [build-details, conda-cpp-build] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - script: ci/build_python.sh - # Build a conda package for each CUDA x ARCH x minimum supported Python version - matrix_filter: group_by({CUDA_VER, ARCH}) | map(min_by(.PY_VER | split(".") | map(tonumber))) - conda-python-tests: - needs: [conda-python-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda - with: - build_type: pull-request - script: ci/test_python.sh - rocky8-clib-standalone-build: - needs: checks - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/clib-build.yaml - with: - build_type: pull-request - rocky8-clib-tests: - needs: [rocky8-clib-standalone-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_standalone_c - with: - build_type: pull-request - container-image-repository: "rapidsai/citestwheel" - script: "ci/test_standalone_c.sh" - conda-java-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_java || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - node_type: cpu8 - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_java.sh" - artifact-name-prefix: "cuvs-java" - file_to_upload: "java/cuvs-java/target/" - conda-java-tests: - needs: conda-java-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - with: - build_type: pull-request - script: "ci/test_java.sh" - conda-lucene-build: - # Depends on the Java job for the cuvs-java artifact that cuvs-lucene builds against. - needs: conda-java-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - node_type: cpu8 - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_lucene.sh" - artifact-name-prefix: "cuvs-lucene" - file_to_upload: "java/cuvs-lucene/target/" - conda-lucene-tests: - needs: conda-lucene-build - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - with: - build_type: pull-request - script: "ci/test_lucene.sh" - rust-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_rust || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_rust.sh" - go-build: - needs: [conda-cpp-build, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - uses: ./.github/workflows/matrix-job.yaml - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_go || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp - # Artifacts are not published from these jobs, so it's safe to run for multiple CUDA versions. - # If these jobs start producing artifacts, the names will have to differentiate between CUDA versions. - with: - build_type: pull-request - matrix_filter: map(select(.ARCH == "amd64")) - script: "ci/build_go.sh" - docs-build-matrix: - permissions: - contents: read - uses: rapidsai/shared-workflows/.github/workflows/compute-matrix.yaml@main - with: - build_type: pull-request - matrix_name: docs-build - docs-build: - needs: [docs-build-matrix] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).build_docs - strategy: - fail-fast: false - matrix: ${{ fromJSON(needs.docs-build-matrix.outputs.matrix) }} - with: - build_type: pull-request - node_type: "gpu-${{ matrix.GPU }}-${{ matrix.DRIVER }}-1" - arch: "${{ matrix.arch }}" - container_image: "rapidsai/ci-conda:26.12-cuda${{ matrix.CUDA_VER }}-${{ matrix.LINUX_VER }}-py${{ matrix.PY_VER }}" - script: "ci/build_docs.sh" - wheel-build-libcuvs: - needs: [build-details, checks] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - node_type: cpu16 - script: ci/build_wheel_libcuvs.sh - # build for every combination of arch and CUDA version, but only for the latest Python - matrix_filter: group_by([.ARCH, (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by(.PY_VER|split(".")|map(tonumber))) - package-name: libcuvs - package-type: cpp - wheel-build-cuvs: - needs: [build-details, wheel-build-libcuvs] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main - with: - build_type: pull-request - build-datetime: ${{ needs.build-details.outputs.build-datetime }} - node_type: cpu8 - script: ci/build_wheel_cuvs.sh - package-name: cuvs - package-type: python - # Build a wheel for each CUDA x ARCH x minimum supported Python version - matrix_filter: group_by({CUDA_VER, ARCH}) | map(min_by(.PY_VER | split(".") | map(tonumber))) - wheel-tests-cuvs: - needs: [wheel-build-cuvs, changed-files] - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main - if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels - with: - build_type: pull-request - script: ci/test_wheel_cuvs.sh - devcontainer: - permissions: - actions: read - contents: read - id-token: write - packages: read - pull-requests: read - secrets: inherit # zizmor: ignore[secrets-inherit] - needs: telemetry-setup - uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@main - with: - arch: '["amd64", "arm64"]' - cuda: '["13.3"]' - node_type: "cpu8" - env: | - SCCACHE_DIST_MAX_RETRIES=inf - SCCACHE_SERVER_LOG=sccache=debug - SCCACHE_DIST_FALLBACK_TO_LOCAL_COMPILE=false - build_command: | - sccache --zero-stats; - build-all -j0 --verbose 2>&1 | tee telemetry-artifacts/build.log; - sccache --show-adv-stats | tee telemetry-artifacts/sccache-stats.txt; - telemetry-summarize: - # This job must use a self-hosted runner to record telemetry traces. - runs-on: linux-amd64-cpu4 - needs: pr-builder - permissions: - actions: write - if: ${{ vars.TELEMETRY_ENABLED == 'true' && !cancelled() }} - continue-on-error: true - steps: - - name: Telemetry summarize - uses: rapidsai/shared-actions/telemetry-dispatch-summarize@main - env: - GH_TOKEN: ${{ github.token }} diff --git a/ci/test_cpp.sh b/ci/test_cpp.sh index 0e97a0ddef..694de5783e 100755 --- a/ci/test_cpp.sh +++ b/ci/test_cpp.sh @@ -52,23 +52,58 @@ nvidia-smi # RAPIDS_DATASET_ROOT_DIR is used by test scripts RAPIDS_DATASET_ROOT_DIR=${RAPIDS_TESTS_DIR}/dataset export RAPIDS_DATASET_ROOT_DIR -./ci/get_test_data.sh --NEIGHBORS_ANN_VAMANA_TEST +# skipped for the reproducer: ./ci/get_test_data.sh --NEIGHBORS_ANN_VAMANA_TEST EXITCODE=0 trap "EXITCODE=1" ERR set +e -# Run Python build utilities tests (once, in the first shard) -if [[ "${SHARD}" == "1" ]]; then - rapids-logger "Run libcuvs Python build utilities tests" - pytest cpp/tests/python -fi - -# Run libcuvs gtests from libcuvs-tests package -rapids-logger "Run libcuvs tests (shard ${SHARD} of ${NUM_SHARDS})" +# Flaky spectral clustering reproducer (DO NOT MERGE) +echo "Ignoring shard ${SHARD} of ${NUM_SHARDS}" pushd "$CONDA_PREFIX"/bin/gtests/libcuvs -timeout -v --signal=SIGINT --kill-after=60s 100m ctest -j8 --output-on-failure -I "${SHARD},,${NUM_SHARDS}" +LOGDIR="${RAPIDS_TESTS_DIR}/spectral" +mkdir -p "${LOGDIR}" + +rapids-logger "Full CLUSTER_TEST binary in fresh processes" +FULL_RUNS=${FULL_RUNS:-5} +full_fail=0 +for i in $(seq 1 "${FULL_RUNS}"); do + if ! ./CLUSTER_TEST --gtest_filter='-SpectralClusteringDiag*' > "${LOGDIR}/full_${i}.log" 2>&1; then + full_fail=$((full_fail + 1)) + grep -E "FAILED|Score|Failure|actual|Expected|Actual|eigensolver" "${LOGDIR}/full_${i}.log" || true + fi +done +echo "SUMMARY full CLUSTER_TEST: ${full_fail}/${FULL_RUNS} runs failed" + +rapids-logger "Spectral clustering diagnostics" +SPECTRAL_DIAG_REPEATS=${SPECTRAL_DIAG_REPEATS:-300} ./CLUSTER_TEST --gtest_filter='SpectralClusteringDiag*' 2>&1 | tee "${LOGDIR}/diag.log" | grep -v "graph diff" || true + +rapids-logger "Spectral clustering tests repeated in one process" +./CLUSTER_TEST --gtest_filter='SpectralClusteringTest*' --gtest_repeat=300 --gtest_brief=1 > "${LOGDIR}/repeat.log" 2>&1 || true +grep -E "FAILED|Score|actual" "${LOGDIR}/repeat.log" | sort | uniq -c | sort -rn | head -50 || true +echo "SUMMARY in-process repeat: $(grep -c '^\[ FAILED \] .*ms)$' "${LOGDIR}/repeat.log" || true) failed test instances (300 repeats x 22 tests)" + +rapids-logger "SpectralClusteringTestF.Result/7 in fresh processes" +FRESH_RUNS=${FRESH_RUNS:-200} +fresh_fail=0 +for i in $(seq 1 "${FRESH_RUNS}"); do + if ! ./CLUSTER_TEST --gtest_filter='SpectralClusteringTests/SpectralClusteringTestF.Result/7' --gtest_brief=1 > "${LOGDIR}/fresh_${i}.log" 2>&1; then + fresh_fail=$((fresh_fail + 1)) + grep -E "Score|actual|eigensolver" "${LOGDIR}/fresh_${i}.log" || true + fi +done +echo "SUMMARY fresh-process Result/7: ${fresh_fail}/${FRESH_RUNS} runs failed" + +rapids-logger "Distance and 1-NN tests" +other_fail=0 +./DISTANCE_TEST --gtest_brief=1 || other_fail=$((other_fail + 1)) +./NEIGHBORS_TEST --gtest_filter='NNTest*:Top1nn*' --gtest_brief=1 || other_fail=$((other_fail + 1)) +echo "SUMMARY distance / 1-NN test binaries failed: ${other_fail}" popd +if [[ ${full_fail} -gt 0 || ${fresh_fail} -gt 0 || ${other_fail} -gt 0 ]] || grep -q '^\[ FAILED \]' "${LOGDIR}/repeat.log"; then + EXITCODE=1 +fi + rapids-logger "Test script exiting with value: $EXITCODE" exit ${EXITCODE} diff --git a/cpp/src/distance/detail/distance_ops/l2_exp.cuh b/cpp/src/distance/detail/distance_ops/l2_exp.cuh index c0433c676a..d059836ed1 100644 --- a/cpp/src/distance/detail/distance_ops/l2_exp.cuh +++ b/cpp/src/distance/detail/distance_ops/l2_exp.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -28,6 +28,23 @@ __device__ constexpr AccT get_clamp_precision() } } +/** + * Whether an expanded L2 distance is round-off error of a self-distance and should be clamped to + * zero. Self-neighboring points have (aNorm == bNorm) == accVal, but the dot product (accVal) can + * have round-off errors, which will cause (aNorm == bNorm) ~ accVal instead. + * + * The round-off error scales with the norms, so the tolerance is relative to them for squared + * norms below one. Otherwise, distinct points with equal norms and a small magnitude (for example, + * a spectral embedding) would have their distances incorrectly clamped to zero. + */ +template +DI bool is_self_distance_round_off(AccT dist, AccT aNorm, AccT bNorm) +{ + const AccT norm_product = aNorm * bNorm; + const AccT norm_scale = norm_product < AccT(1) ? norm_product : AccT(1); + return (aNorm == bNorm) && (dist * dist < get_clamp_precision() * norm_scale); +} + // Epilogue operator for CUTLASS based kernel template struct l2_exp_cutlass_op { @@ -38,13 +55,7 @@ struct l2_exp_cutlass_op { inline __device__ AccT operator()(AccT aNorm, AccT bNorm, AccT accVal) const noexcept { AccT outVal = aNorm + bNorm - AccT(2.0) * accVal; - - /** - * Self-neighboring points should have (aNorm == bNorm) == accVal and the dot product (accVal) - * can sometimes have round-off errors, which will cause (aNorm == bNorm) ~ accVal instead. - */ - outVal = - outVal * AccT(!((outVal * outVal < get_clamp_precision()) * (aNorm == bNorm))); + outVal = outVal * AccT(!is_self_distance_round_off(outVal, aNorm, bNorm)); return sqrt ? raft::sqrt(outVal * static_cast(outVal > AccT(0))) : outVal; } @@ -112,15 +123,9 @@ struct l2_exp_distance_op { for (int j = 0; j < Policy::AccColsPerTh; ++j) { AccT accVal = acc[i][j]; AccT val = regxn[i] + regyn[j] - (AccT)2.0 * accVal; - - /** - * Self-neighboring points should have (aNorm == bNorm) == accVal and the dot product - * (accVal) can sometimes have round-off errors, which will cause (aNorm == bNorm) ~ accVal - * instead. - */ - acc[i][j] = val * static_cast((val > AccT(0))) * - static_cast( - !((val * val < get_clamp_precision()) * (regxn[i] == regyn[j]))); + acc[i][j] = + val * static_cast((val > AccT(0))) * + static_cast(!is_self_distance_round_off(val, regxn[i], regyn[j])); } } if (sqrt) { diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index eee4150eed..c60c5dab18 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -141,6 +141,7 @@ ConfigureTest( cluster/linkage.cu cluster/connect_knn.cu cluster/spectral.cu + cluster/spectral_diag.cu GPUS 1 PERCENT 100 ) diff --git a/cpp/tests/cluster/diag_lanczos.cuh b/cpp/tests/cluster/diag_lanczos.cuh new file mode 100644 index 0000000000..8273e362c4 --- /dev/null +++ b/cpp/tests/cluster/diag_lanczos.cuh @@ -0,0 +1,889 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +// Temporary instrumented copy of raft/sparse/solver/detail/lanczos.cuh (DO NOT MERGE). +#pragma once + +// for cmath: +#define _USE_MATH_DEFINES + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace raft { +namespace sparse::solver::diag_detail { +inline std::string& lanczos_trace() +{ + static std::string t; + return t; +} + +template +RAFT_KERNEL kernel_triangular_populate(T* M, const T* beta, int n) +{ + int row = blockIdx.x * blockDim.x + threadIdx.x; + + if (row < n) { + // Upper diagonal: M[row + 1, row] in column-major + if (row < n - 1) { M[(row + 1) * n + row] = beta[row]; } + + // Lower diagonal: M[row - 1, row] in column-major + if (row > 0) { M[(row - 1) * n + row] = beta[row - 1]; } + } +} + +template +RAFT_KERNEL kernel_triangular_beta_k(T* t, const T* beta_k, int k, int n) +{ + int tid = threadIdx.x + blockIdx.x * blockDim.x; + + if (tid < k) { + // Update the k-th column: t[i, k] -> t[k * n + i] in column-major + t[tid * n + k] = beta_k[tid]; + + // Update the k-th row: t[k, j] -> t[j * n + k] in column-major + t[k * n + tid] = beta_k[tid]; + } +} + +template +RAFT_KERNEL kernel_normalize(const T* u, const T* beta, int j, int n, T* v, T* V, int size) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < size) { + if (beta[j] == 0) { + v[i] = u[i] / 1; + } else { + v[i] = u[i] / beta[j]; + } + V[i + (j + 1) * n] = v[i]; + } +} + +template +RAFT_KERNEL kernel_clamp_down(T* value, T threshold) +{ + *value = (fabs(*value) < threshold) ? 0 : *value; +} + +template +RAFT_KERNEL kernel_clamp_down_vector(T* vec, T threshold, int size) +{ + int idx = threadIdx.x + blockIdx.x * blockDim.x; + if (idx < size) { vec[idx] = (fabs(vec[idx]) < threshold) ? 0 : vec[idx]; } +} + +template +void lanczos_solve_ritz( + raft::resources const& handle, + raft::device_matrix_view alpha, + raft::device_matrix_view beta, + std::optional> beta_k, + IndexTypeT k, + LANCZOS_WHICH which, + int ncv, + raft::device_matrix_view eigenvectors, + raft::device_vector_view eigenvalues, + raft::device_matrix_view& eigenvectors_k, + raft::device_vector_view& eigenvalues_k, + raft::device_matrix_view& eigenvectors_k_slice, + raft::device_vector_view sm_eigenvalues, + raft::device_matrix_view sm_eigenvectors) +{ + auto stream = resource::get_cuda_stream(handle); + bool const is_dry_run = resource::get_dry_run_flag(handle); + + ValueTypeT zero = 0; + auto triangular_matrix = + raft::make_device_matrix(handle, ncv, ncv); + raft::matrix::fill(handle, triangular_matrix.view(), zero); + + raft::device_vector_view alphaVec = + raft::make_device_vector_view(alpha.data_handle(), ncv); + raft::matrix::set_diagonal(handle, alphaVec, triangular_matrix.view()); + + int blockSize = 256; + int numBlocks = raft::div_rounding_up_safe(ncv, blockSize); + raft::launch_kernel(handle, + numBlocks, + blockSize, + kernel_triangular_populate, + triangular_matrix.data_handle(), + beta.data_handle(), + ncv); + + if (beta_k) { + int threadsPerBlock = 256; + int blocksPerGrid = raft::div_rounding_up_safe(k, threadsPerBlock); + raft::launch_kernel(handle, + blocksPerGrid, + threadsPerBlock, + kernel_triangular_beta_k, + triangular_matrix.data_handle(), + beta_k.value().data_handle(), + k, + ncv); + } + + auto triangular_matrix_view = + raft::make_device_matrix_view( + triangular_matrix.data_handle(), ncv, ncv); + + raft::linalg::eig_dc(handle, triangular_matrix_view, eigenvectors, eigenvalues); + + IndexTypeT nEigVecs = k; + + auto indices = raft::make_device_vector(handle, ncv); + auto selected_indices = raft::make_device_vector(handle, nEigVecs); + + if (which == LANCZOS_WHICH::SA) { + eigenvectors_k = raft::make_device_matrix_view( + eigenvectors.data_handle(), ncv, nEigVecs); + eigenvalues_k = + raft::make_device_vector_view(eigenvalues.data_handle(), nEigVecs); + eigenvectors_k_slice = raft::make_device_matrix_view( + eigenvectors.data_handle(), ncv, nEigVecs); + } else if (which == LANCZOS_WHICH::LA) { + eigenvectors_k = raft::make_device_matrix_view( + eigenvectors.data_handle() + (ncv - nEigVecs) * ncv, ncv, nEigVecs); + eigenvalues_k = raft::make_device_vector_view( + eigenvalues.data_handle() + (ncv - nEigVecs), nEigVecs); + eigenvectors_k_slice = raft::make_device_matrix_view( + eigenvectors.data_handle() + (ncv - nEigVecs) * ncv, ncv, nEigVecs); + } else if (which == LANCZOS_WHICH::SM || which == LANCZOS_WHICH::LM) { + if (!is_dry_run) { // TODO: we must be missing some allocations in thrust helpers here + thrust::sequence(thrust::device, indices.data_handle(), indices.data_handle() + ncv, 0); + + // Sort indices by absolute eigenvalues (magnitude) using a custom comparator + thrust::sort(thrust::device, + indices.data_handle(), + indices.data_handle() + ncv, + [eigenvalues = eigenvalues.data_handle()] __device__(int a, int b) { + return fabsf(eigenvalues[a]) < fabsf(eigenvalues[b]); + }); + + if (which == LANCZOS_WHICH::SM) { + // Take the first nEigVecs indices (smallest magnitude) + raft::copy(selected_indices.data_handle(), indices.data_handle(), nEigVecs, stream); + } else if (which == LANCZOS_WHICH::LM) { + // Take the last nEigVecs indices (largest magnitude) + raft::copy(selected_indices.data_handle(), + indices.data_handle() + (ncv - nEigVecs), + nEigVecs, + stream); + } + + // Re-sort these indices by algebraic value to maintain algebraic ordering + thrust::sort(thrust::device, + selected_indices.data_handle(), + selected_indices.data_handle() + nEigVecs, + [eigenvalues = eigenvalues.data_handle()] __device__(int a, int b) { + return eigenvalues[a] < eigenvalues[b]; + }); + } + raft::matrix::gather( + handle, + raft::make_device_matrix_view( + eigenvalues.data_handle(), ncv, 1), + raft::make_device_vector_view(selected_indices.data_handle(), nEigVecs), + raft::make_device_matrix_view( + sm_eigenvalues.data_handle(), nEigVecs, 1)); + raft::matrix::gather( + handle, + raft::make_device_matrix_view( + eigenvectors.data_handle(), ncv, ncv), + raft::make_device_vector_view(selected_indices.data_handle(), nEigVecs), + raft::make_device_matrix_view( + sm_eigenvectors.data_handle(), nEigVecs, ncv)); + + eigenvectors_k = raft::make_device_matrix_view( + sm_eigenvectors.data_handle(), ncv, nEigVecs); + eigenvalues_k = + raft::make_device_vector_view(sm_eigenvalues.data_handle(), nEigVecs); + eigenvectors_k_slice = raft::make_device_matrix_view( + sm_eigenvectors.data_handle(), ncv, nEigVecs); + } +} + +template +void lanczos_aux(raft::resources const& handle, + AType A, + raft::device_matrix_view V, + raft::device_matrix_view u, + raft::device_matrix_view alpha, + raft::device_matrix_view beta, + int start_idx, + int end_idx, + int ncv, + raft::device_matrix_view v, + raft::device_matrix_view uu, + raft::device_matrix_view vv, + std::optional seed) +{ + // Deterministic when seed is provided + cusparseSpMVAlg_t spmv_alg; + if (seed.has_value()) { + if constexpr (is_device_coo_matrix_view::value) { + spmv_alg = CUSPARSE_SPMV_COO_ALG2; + } else { + spmv_alg = CUSPARSE_SPMV_CSR_ALG2; + } + } else { + spmv_alg = CUSPARSE_SPMV_ALG_DEFAULT; + } + auto stream = resource::get_cuda_stream(handle).get(); + bool const is_dry_run = resource::get_dry_run_flag(handle); + + IndexTypeT n = A.structure_view().get_n_rows(); + auto v_vector = raft::make_device_vector_view(v.data_handle(), n); + auto u_vector = raft::make_device_vector_view(u.data_handle(), n); + + auto cusparse_h = resource::get_cusparse_handle(handle); + cusparseSpMatDescr_t cusparse_A = raft::sparse::linalg::detail::create_descriptor(A); + + cusparseDnVecDescr_t cusparse_v = raft::sparse::linalg::detail::create_descriptor(v_vector); + cusparseDnVecDescr_t cusparse_u = raft::sparse::linalg::detail::create_descriptor(u_vector); + + ValueTypeT one = 1; + ValueTypeT zero = 0; + size_t bufferSize; + raft::sparse::detail::cusparsespmv_buffersize(cusparse_h, + CUSPARSE_OPERATION_NON_TRANSPOSE, + &one, + cusparse_A, + cusparse_v, + &zero, + cusparse_u, + spmv_alg, + &bufferSize, + stream); + auto cusparse_spmv_buffer = raft::make_device_vector(handle, bufferSize); + + if (!is_dry_run) { + raft::copy( + v.data_handle(), V.data_handle() + start_idx * V.stride(0), n, stream); // V(start_idx, 0) + } + + for (int i = start_idx; i < end_idx; i++) { + if (!is_dry_run) { + raft::sparse::detail::cusparsespmv(cusparse_h, + CUSPARSE_OPERATION_NON_TRANSPOSE, + &one, + cusparse_A, + cusparse_v, + &zero, + cusparse_u, + spmv_alg, + cusparse_spmv_buffer.data_handle(), + stream); + } + + auto alpha_i = + raft::make_device_scalar_view(alpha.data_handle() + i * alpha.stride(1)); // alpha(0, i) + raft::linalg::dot(handle, v_vector, u_vector, alpha_i); + + raft::matrix::fill(handle, vv, zero); + + auto cublas_h = resource::get_cublas_handle(handle); + + ValueTypeT alpha_i_host = 0; + ValueTypeT b = 0; + ValueTypeT mone = -1; + + if (!is_dry_run) { + raft::copy( + &b, beta.data_handle() + ((i - 1 + ncv) % ncv) * beta.stride(1), 1, stream); + raft::copy( + &alpha_i_host, alpha.data_handle() + i * alpha.stride(1), 1, stream); // alpha(0, i) + } + + raft::linalg::axpy(handle, n, &alpha_i_host, v.data_handle(), 1, vv.data_handle(), 1, stream); + raft::linalg::axpy(handle, + n, + &b, + V.data_handle() + (((i - 1 + ncv) % ncv) * V.stride(0)), + 1, + vv.data_handle(), + 1, + stream); + raft::linalg::axpy(handle, n, &mone, vv.data_handle(), 1, u.data_handle(), 1, stream); + + raft::linalg::gemv(handle, + CUBLAS_OP_T, + n, + i + 1, + &one, + V.data_handle(), + n, + u.data_handle(), + 1, + &zero, + uu.data_handle(), + 1, + stream); + + raft::linalg::gemv(handle, + CUBLAS_OP_N, + n, + i + 1, + &mone, + V.data_handle(), + n, + uu.data_handle(), + 1, + &one, + u.data_handle(), + 1, + stream); + + auto uu_i = raft::make_device_scalar_view(uu.data_handle() + uu.stride(1) * i); // uu(0, i) + raft::linalg::add(handle, make_const_mdspan(alpha_i), make_const_mdspan(uu_i), alpha_i); + + raft::launch_kernel(handle, + 1, + 1, + kernel_clamp_down, + alpha_i.data_handle(), + static_cast(1e-9)); + + auto output = raft::make_device_vector_view( + beta.data_handle() + beta.stride(1) * i, 1); + auto input = raft::make_device_matrix_view(u.data_handle(), 1, n); + raft::linalg::norm( + handle, input, output, raft::sqrt_op()); + + int blockSize = 256; + int numBlocks = (n + blockSize - 1) / blockSize; + + raft::launch_kernel(handle, + numBlocks, + blockSize, + kernel_clamp_down_vector, + u.data_handle(), + static_cast(1e-7), + n); + + raft::launch_kernel(handle, + 1, + 1, + kernel_clamp_down, + beta.data_handle() + beta.stride(1) * i, + static_cast(1e-6)); + + if (i >= end_idx - 1) { break; } + + int threadsPerBlock = 256; + int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; + + raft::launch_kernel(handle, + blocksPerGrid, + threadsPerBlock, + kernel_normalize, + u.data_handle(), + beta.data_handle(), + i, + n, + v.data_handle(), + V.data_handle(), + n); + } +} + +template +auto lanczos_smallest(raft::resources const& handle, + AType A, + int nEigVecs, + int maxIter, + int restartIter, + ValueTypeT tol, + LANCZOS_WHICH which, + ValueTypeT* eigVals_dev, + ValueTypeT* eigVecs_dev, + ValueTypeT* v0, + std::optional seed) -> int +{ + // Deterministic when seed is provided + cusparseSpMVAlg_t spmv_alg; + if (seed.has_value()) { + if constexpr (is_device_coo_matrix_view::value) { + spmv_alg = CUSPARSE_SPMV_COO_ALG2; + } else { + spmv_alg = CUSPARSE_SPMV_CSR_ALG2; + } + } else { + spmv_alg = CUSPARSE_SPMV_ALG_DEFAULT; + } + int n = A.structure_view().get_n_rows(); + int ncv = restartIter; + auto stream = resource::get_cuda_stream(handle).get(); + bool is_dry_run = resource::get_dry_run_flag(handle); + + auto V = raft::make_device_matrix(handle, ncv, n); + auto V_0_view = + raft::make_device_matrix_view(V.data_handle(), 1, n); // First Row V[0] + auto v0_view = raft::make_device_matrix_view(v0, 1, n); + + auto u = raft::make_device_matrix(handle, 1, n); + auto u_vector = raft::make_device_vector_view(u.data_handle(), n); + if (!is_dry_run) { raft::copy(u.data_handle(), v0, n, stream); } + + auto cublas_h = resource::get_cublas_handle(handle); + auto v0nrm = raft::make_device_vector(handle, 1); + raft::linalg::norm( + handle, v0_view, v0nrm.view(), raft::sqrt_op()); + + auto v0_vector_const = raft::make_device_vector_view(v0, n); + + raft::linalg::unary_op( + handle, v0_vector_const, V_0_view, [device_scalar = v0nrm.data_handle()] __device__(auto y) { + return y / *device_scalar; + }); + + auto alpha = raft::make_device_matrix(handle, 1, ncv); + auto beta = raft::make_device_matrix(handle, 1, ncv); + ValueTypeT zero = 0; + raft::matrix::fill(handle, alpha.view(), zero); + raft::matrix::fill(handle, beta.view(), zero); + + auto v = raft::make_device_matrix(handle, 1, n); + auto aux_uu = raft::make_device_matrix(handle, 1, ncv); + auto vv = raft::make_device_matrix(handle, 1, n); + + lanczos_aux(handle, + A, + V.view(), + u.view(), + alpha.view(), + beta.view(), + 0, + ncv, + ncv, + v.view(), + aux_uu.view(), + vv.view(), + seed); + + auto eigenvectors = + raft::make_device_matrix(handle, ncv, ncv); + auto eigenvalues = raft::make_device_vector(handle, ncv); + + raft::device_matrix_view eigenvectors_k; + raft::device_vector_view eigenvalues_k; + raft::device_matrix_view eigenvectors_k_slice; + + auto sm_eigenvalues = raft::make_device_vector(handle, nEigVecs); + auto sm_eigenvectors = + raft::make_device_matrix(handle, ncv, nEigVecs); + + lanczos_solve_ritz(handle, + alpha.view(), + beta.view(), + std::nullopt, + nEigVecs, + which, + ncv, + eigenvectors.view(), + eigenvalues.view(), + eigenvectors_k, + eigenvalues_k, + eigenvectors_k_slice, + sm_eigenvalues.view(), + sm_eigenvectors.view()); + + auto ritz_eigenvectors = + raft::make_device_matrix_view(eigVecs_dev, n, nEigVecs); + + auto V_T = + raft::make_device_matrix_view(V.data_handle(), n, ncv); + raft::linalg::gemm( + handle, V_T, eigenvectors_k, ritz_eigenvectors); + + auto s = raft::make_device_vector(handle, nEigVecs); + + auto S_matrix = raft::make_device_matrix_view( + s.data_handle(), 1, nEigVecs); + + raft::matrix::slice_coordinates coords(ncv - 1, 0, ncv, nEigVecs); + raft::matrix::slice(handle, make_const_mdspan(eigenvectors_k_slice), S_matrix, coords); + + auto beta_k = raft::make_device_vector(handle, nEigVecs); + raft::matrix::fill(handle, beta_k.view(), zero); + auto beta_scalar = raft::make_device_scalar_view(beta.data_handle() + + (ncv - 1) * beta.stride(1)); + + raft::linalg::axpy(handle, beta_scalar, raft::make_const_mdspan(s.view()), beta_k.view()); + + ValueTypeT res = 0; + + raft::device_vector output = + raft::make_device_vector(handle, 1); + raft::device_matrix_view input = + raft::make_device_matrix_view(beta_k.data_handle(), 1, nEigVecs); + raft::linalg::norm( + handle, input, output.view(), raft::sqrt_op()); + if (!is_dry_run) { + raft::copy(&res, output.data_handle(), 1, stream); + resource::sync_stream(handle, stream); + } else { + // Force exactly one loop iteration so the MR records all allocations inside the loop body. + res = tol + 1; + maxIter = ncv + (ncv - nEigVecs); + } + + auto uu = raft::make_device_matrix(handle, 1, nEigVecs); + int iter = ncv; + while (res > tol && iter < maxIter) { + auto beta_view = raft::make_device_matrix_view( + beta.data_handle(), 1, nEigVecs); + raft::matrix::fill(handle, beta_view, zero); + + if (!is_dry_run) { + raft::copy(alpha.data_handle(), eigenvalues_k.data_handle(), nEigVecs, stream); + } + + auto x_T = + raft::make_device_matrix_view(ritz_eigenvectors.data_handle(), nEigVecs, n); + + if (!is_dry_run) { raft::copy(V.data_handle(), x_T.data_handle(), nEigVecs * n, stream); } + + ValueTypeT one = 1; + ValueTypeT mone = -1; + + raft::linalg::gemv(handle, + CUBLAS_OP_T, + n, + nEigVecs, + &one, + V.data_handle(), + n, + u.data_handle(), + 1, + &zero, + uu.data_handle(), + 1, + stream); + + raft::linalg::gemv(handle, + CUBLAS_OP_N, + n, + nEigVecs, + &mone, + V.data_handle(), + n, + uu.data_handle(), + 1, + &one, + u.data_handle(), + 1, + stream); + + auto V_0_view = + raft::make_device_matrix_view(V.data_handle() + (nEigVecs * n), 1, n); + auto V_0_view_vector = + raft::make_device_vector_view(V_0_view.data_handle(), n); + auto unrm = raft::make_device_vector(handle, 1); + raft::linalg::norm( + handle, raft::make_const_mdspan(u.view()), unrm.view(), raft::sqrt_op()); + + raft::linalg::unary_op( + handle, + raft::make_const_mdspan(u_vector), + V_0_view, + [device_scalar = unrm.data_handle()] __device__(auto y) { return y / *device_scalar; }); + + auto cusparse_h = resource::get_cusparse_handle(handle); + cusparseSpMatDescr_t cusparse_A = raft::sparse::linalg::detail::create_descriptor(A); + + cusparseDnVecDescr_t cusparse_v = + raft::sparse::linalg::detail::create_descriptor(V_0_view_vector); + cusparseDnVecDescr_t cusparse_u = raft::sparse::linalg::detail::create_descriptor(u_vector); + + ValueTypeT zero = 0; + size_t bufferSize; + raft::sparse::detail::cusparsespmv_buffersize(cusparse_h, + CUSPARSE_OPERATION_NON_TRANSPOSE, + &one, + cusparse_A, + cusparse_v, + &zero, + cusparse_u, + spmv_alg, + &bufferSize, + stream); + auto cusparse_spmv_buffer = raft::make_device_vector(handle, bufferSize); + + if (!is_dry_run) { + raft::sparse::detail::cusparsespmv(cusparse_h, + CUSPARSE_OPERATION_NON_TRANSPOSE, + &one, + cusparse_A, + cusparse_v, + &zero, + cusparse_u, + spmv_alg, + cusparse_spmv_buffer.data_handle(), + stream); + } + + auto alpha_k = raft::make_device_scalar_view(alpha.data_handle() + nEigVecs); + + raft::linalg::dot( + handle, make_const_mdspan(V_0_view_vector), make_const_mdspan(u_vector), alpha_k); + + raft::linalg::binary_op(handle, + make_const_mdspan(u_vector), + make_const_mdspan(V_0_view_vector), + u_vector, + [device_scalar_ptr = alpha_k.data_handle()] __device__( + ValueTypeT u_element, ValueTypeT V_0_element) { + return u_element - (*device_scalar_ptr) * V_0_element; + }); + + auto temp = raft::make_device_vector(handle, n); + + auto V_k = raft::make_device_matrix_view( + V.data_handle(), nEigVecs, n); + auto V_k_T = + raft::make_device_matrix(handle, n, nEigVecs); + + raft::linalg::transpose(handle, V_k, V_k_T.view()); + + raft::linalg::gemv(handle, + CUBLAS_OP_N, + n, + nEigVecs, + &one, + V_k.data_handle(), + n, + beta_k.data_handle(), + 1, + &zero, + temp.data_handle(), + 1, + stream); + + auto one_scalar = raft::make_device_scalar(handle, 1); + raft::linalg::binary_op(handle, + make_const_mdspan(u_vector), + make_const_mdspan(temp.view()), + u_vector, + [device_scalar_ptr = one_scalar.data_handle()] __device__( + ValueTypeT u_element, ValueTypeT temp_element) { + return u_element - (*device_scalar_ptr) * temp_element; + }); + + auto output1 = raft::make_device_vector_view( + beta.data_handle() + beta.stride(1) * nEigVecs, 1); + raft::linalg::norm( + handle, raft::make_const_mdspan(u.view()), output1, raft::sqrt_op()); + + auto V_kplus1 = + raft::make_device_vector_view(V.data_handle() + V.stride(0) * (nEigVecs + 1), n); + + raft::linalg::unary_op( + handle, + make_const_mdspan(u_vector), + V_kplus1, + [device_scalar = (beta.data_handle() + beta.stride(1) * nEigVecs)] __device__(auto y) { + return y / *device_scalar; + }); + + lanczos_aux(handle, + A, + V.view(), + u.view(), + alpha.view(), + beta.view(), + nEigVecs + 1, + ncv, + ncv, + v.view(), + aux_uu.view(), + vv.view(), + seed); + iter += ncv - nEigVecs; + lanczos_solve_ritz(handle, + alpha.view(), + beta.view(), + beta_k.view(), + nEigVecs, + which, + ncv, + eigenvectors.view(), + eigenvalues.view(), + eigenvectors_k, + eigenvalues_k, + eigenvectors_k_slice, + sm_eigenvalues.view(), + sm_eigenvectors.view()); + + auto ritz_eigenvectors = raft::make_device_matrix_view( + eigVecs_dev, n, nEigVecs); + + auto V_T = + raft::make_device_matrix_view(V.data_handle(), n, ncv); + raft::linalg::gemm( + handle, V_T, eigenvectors_k, ritz_eigenvectors); + + auto S_matrix = raft::make_device_matrix_view( + s.data_handle(), 1, nEigVecs); + + raft::matrix::slice_coordinates coords(ncv - 1, 0, ncv, nEigVecs); + raft::matrix::slice(handle, make_const_mdspan(eigenvectors_k_slice), S_matrix, coords); + + raft::matrix::fill(handle, beta_k.view(), zero); + + auto beta_scalar = raft::make_device_scalar_view( + beta.data_handle() + beta.stride(1) * (ncv - 1)); // &((beta.view())(0, ncv - 1)) + + raft::linalg::axpy(handle, beta_scalar, raft::make_const_mdspan(s.view()), beta_k.view()); + + raft::device_vector output2 = + raft::make_device_vector(handle, 1); + raft::device_matrix_view input2 = + raft::make_device_matrix_view(beta_k.data_handle(), 1, nEigVecs); + raft::linalg::norm( + handle, input2, output2.view(), raft::sqrt_op()); + if (!is_dry_run) { + raft::copy(&res, output2.data_handle(), 1, stream); + resource::sync_stream(handle, stream); + } + { + std::vector ev(nEigVecs); + raft::copy(ev.data(), eigenvalues_k.data_handle(), nEigVecs, stream); + std::vector bt(ncv); + raft::copy(bt.data(), beta.data_handle(), ncv, stream); + resource::sync_stream(handle, stream); + char buf[512]; + int len = snprintf(buf, + sizeof(buf), + "[lanczos] iter=%d res=%.6g beta_last=%.6g ritz=[", + iter, + (double)res, + (double)bt[ncv - 1]); + for (auto e : ev) + len += snprintf(buf + len, sizeof(buf) - len, " %.6g", (double)e); + snprintf(buf + len, sizeof(buf) - len, " ]\n"); + lanczos_trace() += buf; + } + } + + if (!is_dry_run) { + raft::copy(eigVals_dev, eigenvalues_k.data_handle(), nEigVecs, stream); + raft::copy(eigVecs_dev, ritz_eigenvectors.data_handle(), n * nEigVecs, stream); + } + + { + char buf[256]; + snprintf(buf, sizeof(buf), "[lanczos] done: iter=%d res=%.6g\n", iter, (double)res); + lanczos_trace() += buf; + } + return 0; +} + +template +auto lanczos_compute_eigenpairs( + raft::resources const& handle, + lanczos_solver_config const& config, + AType A, + std::optional> v0, + raft::device_vector_view eigenvalues, + raft::device_matrix_view eigenvectors) -> int +{ + if (v0.has_value()) { + return lanczos_smallest(handle, + A, + config.n_components, + config.max_iterations, + config.ncv, + config.tolerance, + config.which, + eigenvalues.data_handle(), + eigenvectors.data_handle(), + v0->data_handle(), + config.seed); + } else { + // Handle the optional v0 initial Lanczos vector if nullopt is used + auto n = A.structure_view().get_n_rows(); + auto temp_v0 = raft::make_device_vector(handle, n); + uint64_t seed = config.seed.value_or(std::random_device{}()); + raft::random::RngState rng_state(seed); + raft::random::uniform(handle, rng_state, temp_v0.view(), ValueTypeT{0.0}, ValueTypeT{1.0}); + return lanczos_smallest(handle, + A, + config.n_components, + config.max_iterations, + config.ncv, + config.tolerance, + config.which, + eigenvalues.data_handle(), + eigenvectors.data_handle(), + temp_v0.data_handle(), + config.seed); + } +} + +} // namespace sparse::solver::diag_detail +} // namespace raft diff --git a/cpp/tests/cluster/spectral_diag.cu b/cpp/tests/cluster/spectral_diag.cu new file mode 100644 index 0000000000..a7b08dce5c --- /dev/null +++ b/cpp/tests/cluster/spectral_diag.cu @@ -0,0 +1,367 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +// Temporary diagnostics for the flaky SpectralClusteringTestF.Result/7 (DO NOT MERGE). + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "diag_lanczos.cuh" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cuvs { +namespace { + +template +auto to_host(const T* d, size_t n, cudaStream_t s) -> std::vector +{ + std::vector h(n); + raft::update_host(h.data(), d, n, s); + RAFT_CUDA_TRY(cudaStreamSynchronize(s)); + return h; +} + +template +auto fnv(const std::vector& v) -> uint64_t +{ + uint64_t h = 1469598103934665603ULL; + const auto* bytes = reinterpret_cast(v.data()); + for (size_t i = 0; i < v.size() * sizeof(T); i++) { + h ^= bytes[i]; + h *= 1099511628211ULL; + } + return h; +} + +auto n_components(int n, const std::vector& rows, const std::vector& cols) -> int +{ + std::vector p(n); + std::iota(p.begin(), p.end(), 0); + std::function find = [&](int x) { return p[x] == x ? x : p[x] = find(p[x]); }; + for (size_t i = 0; i < rows.size(); i++) { + int a = find(rows[i]), b = find(cols[i]); + if (a != b) p[a] = b; + } + int c = 0; + for (int i = 0; i < n; i++) + if (find(i) == i) c++; + return c; +} + +struct counter { + std::map m; + int add(uint64_t h) + { + auto [it, inserted] = m.emplace(h, static_cast(m.size())); + return it->second; + } +}; + +} // namespace + +void run_spectral_diag() +{ + int repeats = 50; + if (const char* e = std::getenv("SPECTRAL_DIAG_REPEATS")) repeats = std::atoi(e); + + const int n_samples = 1000, n_features = 30, n_clusters = 5, n_comp = 5, n_neighbors = 20, + n_init = 3; + const float cluster_std = 0.3f; + uint64_t seed = 444ULL; + if (const char* e = std::getenv("SPECTRAL_DIAG_SEED")) seed = std::strtoull(e, nullptr, 10); + + raft::resources handle; + cudaStream_t stream = raft::resource::get_cuda_stream(handle).get(); + + counter c_x, c_graph, c_lap, c_eig, c_emb, c_kmeans_fixed, c_labels, c_fp; + int n_bad = 0, n_bad_fp = 0, n_missed = 0; + std::vector first_embedding; + std::vector first_r, first_c; + std::vector first_v; + + for (int r = 0; r < repeats; r++) { + auto X = raft::make_device_matrix(handle, n_samples, n_features); + auto true_labels = raft::make_device_vector(handle, n_samples); + raft::random::make_blobs(X.data_handle(), + true_labels.data_handle(), + n_samples, + n_features, + n_clusters, + stream, + true, + nullptr, + nullptr, + cluster_std, + false, + -10.0f, + 10.0f, + seed); + int ix = c_x.add(fnv(to_host(X.data_handle(), X.size(), stream))); + + auto graph = raft::make_device_coo_matrix(handle, n_samples, n_samples); + cuvs::preprocessing::spectral_embedding::params embed_params; + embed_params.n_neighbors = n_neighbors; + embed_params.seed = seed; + cuvs::preprocessing::spectral_embedding::helpers::create_connectivity_graph( + handle, embed_params, X.view(), graph); + int nnz = graph.structure_view().get_nnz(); + auto g_r = to_host(graph.structure_view().get_rows().data(), nnz, stream); + auto g_c = to_host(graph.structure_view().get_cols().data(), nnz, stream); + auto g_v = to_host(graph.get_elements().data(), nnz, stream); + int ig = c_graph.add(fnv(g_r) ^ (fnv(g_c) * 3) ^ (fnv(g_v) * 7)); + int ncomp = n_components(n_samples, g_r, g_c); + if (r == 0) { + first_r = g_r; + first_c = g_c; + first_v = g_v; + } else if (g_r != first_r || g_c != first_c || g_v != first_v) { + int nd = 0; + for (int i = 0; i < nnz && nd < 5; i++) { + if (g_r[i] != first_r[i] || g_c[i] != first_c[i] || g_v[i] != first_v[i]) { + printf("[diag] graph diff at %d: (%d,%d,%g) vs first (%d,%d,%g)\n", + i, + g_r[i], + g_c[i], + g_v[i], + first_r[i], + first_c[i], + first_v[i]); + nd++; + } + } + } + + // Replicate the spectral embedding internals so that the eigenvalues can be inspected. + auto diagonal = raft::make_device_vector(handle, n_samples); + auto laplacian = + raft::sparse::linalg::laplacian_normalized(handle, graph.view(), diagonal.view()); + int lnnz = laplacian.structure_view().get_nnz(); + auto lap_vals = + raft::make_device_vector_view(laplacian.get_elements().data(), lnnz); + raft::linalg::map( + handle, lap_vals, [] __device__(float x) { return -x; }, raft::make_const_mdspan(lap_vals)); + int il = c_lap.add(fnv(to_host(laplacian.get_elements().data(), lnnz, stream)) ^ + (fnv(to_host(diagonal.data_handle(), n_samples, stream)) * 5)); + + auto config = raft::sparse::solver::lanczos_solver_config(); + config.n_components = n_comp; + config.max_iterations = 10 * n_samples; + config.ncv = std::min(n_samples - n_comp, std::max(2 * n_comp + 1, 20)); + config.tolerance = 0.0f; + config.which = raft::sparse::solver::LANCZOS_WHICH::LA; + config.seed = seed; + auto evals = raft::make_device_vector(handle, n_comp); + auto evecs = + raft::make_device_matrix(handle, n_samples, n_comp); + raft::sparse::solver::diag_detail::lanczos_trace().clear(); + raft::sparse::solver::diag_detail::lanczos_compute_eigenpairs( + handle, + config, + laplacian.view(), + std::optional>{}, + evals.view(), + evecs.view()); + auto h_evals = to_host(evals.data_handle(), n_comp, stream); + bool missed_null = std::abs(h_evals[0]) > 1e-4f; + if (r == 0 || missed_null) { + printf("[diag] r=%d lanczos trace%s:\n%s", + r, + missed_null ? " (MISSED NULL VECTOR)" : "", + raft::sparse::solver::diag_detail::lanczos_trace().c_str()); + } + if (missed_null) n_missed++; + int ie = + c_eig.add(fnv(h_evals) ^ (fnv(to_host(evecs.data_handle(), evecs.size(), stream)) * 11)); + + if (r == 0 && std::getenv("SPECTRAL_DIAG_SEED_SWEEP")) { + int nsweep = std::atoi(std::getenv("SPECTRAL_DIAG_SEED_SWEEP")); + int n_not_null = 0; + for (int sd = 0; sd < nsweep; sd++) { + auto cfg = config; + cfg.seed = 1000 + sd; + raft::sparse::solver::lanczos_compute_eigenpairs( + handle, cfg, laplacian.view(), std::nullopt, evals.view(), evecs.view()); + auto ev = to_host(evals.data_handle(), n_comp, stream); + bool bad = std::abs(ev[0]) > 1e-4f; + if (bad) n_not_null++; + printf("[diag-sweep] seed=%d evals=[%.6g %.6g %.6g %.6g %.6g]%s\n", + 1000 + sd, + ev[0], + ev[1], + ev[2], + ev[3], + ev[4], + bad ? " BAD" : ""); + } + printf( + "[diag-sweep-summary] %d/%d seeds did not find a 5-dim null space\n", n_not_null, nsweep); + raft::sparse::solver::lanczos_compute_eigenpairs( + handle, config, laplacian.view(), std::nullopt, evals.view(), evecs.view()); + } + + // The public embedding + kmeans path. + cuvs::preprocessing::spectral_embedding::params se; + se.n_components = n_comp; + se.n_neighbors = n_neighbors; + se.norm_laplacian = true; + se.drop_first = false; + se.seed = seed; + se.tolerance = 0.0f; + auto emb_cm = raft::make_device_matrix(handle, n_samples, n_comp); + cuvs::preprocessing::spectral_embedding::transform(handle, se, graph.view(), emb_cm.view()); + auto h_emb = to_host(emb_cm.data_handle(), emb_cm.size(), stream); + int iemb = c_emb.add(fnv(h_emb)); + if (r == 0) first_embedding = h_emb; + + auto run_kmeans = [&](const std::vector& emb_col_major, std::vector& out_labels) { + auto e_cm = raft::make_device_matrix(handle, n_samples, n_comp); + raft::update_device(e_cm.data_handle(), emb_col_major.data(), emb_col_major.size(), stream); + auto e_rm = raft::make_device_matrix(handle, n_samples, n_comp); + raft::linalg::transpose( + handle, e_cm.data_handle(), e_rm.data_handle(), n_samples, n_comp, stream); + cuvs::cluster::kmeans::params kp; + kp.n_clusters = n_clusters; + kp.rng_state = raft::random::RngState(seed); + kp.n_init = n_init; + kp.oversampling_factor = 0.0; + auto labels = raft::make_device_vector(handle, n_samples); + float inertia; + int n_iter; + cuvs::cluster::kmeans::fit_predict(handle, + kp, + e_rm.view(), + std::nullopt, + std::nullopt, + labels.view(), + raft::make_host_scalar_view(&inertia), + raft::make_host_scalar_view(&n_iter)); + out_labels = to_host(labels.data_handle(), n_samples, stream); + return raft::stats::adjusted_rand_index( + true_labels.data_handle(), labels.data_handle(), n_samples, stream); + }; + + std::vector lab_fixed, lab; + run_kmeans(first_embedding, lab_fixed); + int ikf = c_kmeans_fixed.add(fnv(lab_fixed)); + double ari = run_kmeans(h_emb, lab); + int ilab = c_labels.add(fnv(lab)); + + // Full public fit_predict exactly as the original test does it. + cluster::spectral::params params; + params.n_clusters = n_clusters; + params.n_components = n_comp; + params.n_neighbors = n_neighbors; + params.n_init = n_init; + params.tolerance = 0.0f; + params.rng_state = raft::random::RngState(seed); + auto fp_labels = raft::make_device_vector(handle, n_samples); + cluster::spectral::fit_predict(handle, params, graph.view(), fp_labels.view()); + double ari_fp = raft::stats::adjusted_rand_index( + true_labels.data_handle(), fp_labels.data_handle(), n_samples, stream); + int ifp = c_fp.add(fnv(to_host(fp_labels.data_handle(), n_samples, stream))); + + if (ari < 0.7 || ari_fp < 0.7 || r == 0) { + // Summarize the embedding per ground-truth cluster: centroid and max deviation from it. + auto h_true = to_host(true_labels.data_handle(), n_samples, stream); + for (int k = 0; k < n_clusters; k++) { + std::vector mean(n_comp, 0.0); + int cnt = 0; + for (int i = 0; i < n_samples; i++) { + if (h_true[i] != k) continue; + cnt++; + for (int j = 0; j < n_comp; j++) + mean[j] += h_emb[j * n_samples + i]; + } + for (auto& m : mean) + m /= std::max(cnt, 1); + double maxdev = 0; + for (int i = 0; i < n_samples; i++) { + if (h_true[i] != k) continue; + double d = 0; + for (int j = 0; j < n_comp; j++) { + double t = h_emb[j * n_samples + i] - mean[j]; + d += t * t; + } + maxdev = std::max(maxdev, std::sqrt(d)); + } + printf( + "[diag] r=%d true cluster %d (n=%d): centroid=[%.4g %.4g %.4g %.4g %.4g] maxdev=%.4g\n", + r, + k, + cnt, + mean[0], + mean[1], + mean[2], + mean[3], + mean[4], + maxdev); + } + } + if (ari < 0.7) n_bad++; + if (ari_fp < 0.7) n_bad_fp++; + printf( + "[diag] r=%d X=%d graph=%d(nnz=%d%s) lap=%d eig=%d evals=[%.6g %.6g %.6g %.6g %.6g] emb=%d " + "kmeans_fixed=%d labels=%d ari=%.4f fp=%d ari_fp=%.4f\n", + r, + ix, + ig, + nnz, + ncomp >= 0 ? (std::string(",ncomp=") + std::to_string(ncomp)).c_str() : "", + il, + ie, + h_evals[0], + h_evals[1], + h_evals[2], + h_evals[3], + h_evals[4], + iemb, + ikf, + ilab, + ari, + ifp, + ari_fp); + } + printf( + "[diag-summary] repeats=%d distinct: X=%zu graph=%zu lap=%zu eig=%zu emb=%zu kmeans_fixed=%zu " + "labels=%zu fp=%zu; ari<0.7: %d, ari_fp<0.7: %d, lanczos missed null vector: %d\n", + repeats, + c_x.m.size(), + c_graph.m.size(), + c_lap.m.size(), + c_eig.m.size(), + c_emb.m.size(), + c_kmeans_fixed.m.size(), + c_labels.m.size(), + c_fp.m.size(), + n_bad, + n_bad_fp, + n_missed); +} + +TEST(SpectralClusteringDiag, Result7) { run_spectral_diag(); } + +} // namespace cuvs diff --git a/cpp/tests/distance/dist_l2_exp.cu b/cpp/tests/distance/dist_l2_exp.cu index a91daff1b3..1c9dbe4e28 100644 --- a/cpp/tests/distance/dist_l2_exp.cu +++ b/cpp/tests/distance/dist_l2_exp.cu @@ -1,11 +1,17 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "../test_utils.cuh" #include "distance_base.cuh" +#include +#include + +#include +#include + namespace cuvs { namespace distance { @@ -162,5 +168,46 @@ INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucExpTestD, ::testing::ValuesIn( class BigMatrixEucExp : public BigMatrixDistanceTest {}; TEST_F(BigMatrixEucExp, Result) {} + +// Distinct points with equal norms must keep their (small) distances. The clamp of self-distance +// round-off to zero used to apply an absolute tolerance, which collapsed such points when the data +// has a small magnitude (e.g. a spectral embedding). +template +void test_small_scale_equal_norms(T scale) +{ + raft::resources handle; + auto stream = raft::resource::get_cuda_stream(handle); + constexpr int64_t n = 8; + std::vector h_x(n * n, T{0}); + for (int64_t i = 0; i < n; i++) { + h_x[i * n + i] = scale; + } + auto x = raft::make_device_matrix(handle, n, n); + auto d = raft::make_device_matrix(handle, n, n); + raft::update_device(x.data_handle(), h_x.data(), h_x.size(), stream); + + for (auto metric : {DistanceType::L2Expanded, DistanceType::L2SqrtExpanded}) { + cuvs::distance::pairwise_distance(handle, + raft::make_const_mdspan(x.view()), + raft::make_const_mdspan(x.view()), + d.view(), + metric); + std::vector h_d(n * n); + raft::update_host(h_d.data(), d.data_handle(), h_d.size(), stream); + raft::resource::sync_stream(handle, stream); + + const T expected = + metric == DistanceType::L2Expanded ? T(2) * scale * scale : std::sqrt(T(2)) * scale; + for (int64_t i = 0; i < n; i++) { + for (int64_t j = 0; j < n; j++) { + EXPECT_NEAR(h_d[i * n + j], i == j ? T(0) : expected, T(1e-3) * expected) + << "metric=" << static_cast(metric) << " i=" << i << " j=" << j; + } + } + } +} + +TEST(DistanceEucExpSmallScale, EqualNormsF) { test_small_scale_equal_norms(1e-2f); } +TEST(DistanceEucExpSmallScale, EqualNormsD) { test_small_scale_equal_norms(1e-5); } } // end namespace distance } // namespace cuvs diff --git a/cpp/tests/neighbors/distance_nn.cu b/cpp/tests/neighbors/distance_nn.cu index 5ac80fda73..9ab376927c 100644 --- a/cpp/tests/neighbors/distance_nn.cu +++ b/cpp/tests/neighbors/distance_nn.cu @@ -16,8 +16,10 @@ #include #include #include +#include #include +#include namespace cuvs::neighbors { @@ -446,6 +448,77 @@ TEST(Top1nnPlan, RejectsMismatchedLaunch) plan))); } +// Distinct points with equal norms and a small magnitude (as in a spectral embedding) must not be +// treated as self-neighbors when finding the nearest neighbor. +TEST(Top1nnEqualNorms, SmallScale) +{ + using Backend = cuvs::distance::detail::Top1nnBackend; + using Kvp = raft::KeyValuePair; + raft::resources handle; + auto stream = raft::resource::get_cuda_stream(handle); + constexpr int m = 1000; + constexpr int n = 5; + constexpr int k = 5; + constexpr float s = 1e-2f; + std::vector h_x(m * k, 0.0f); + std::vector h_y(n * k, 0.0f); + for (int i = 0; i < m; i++) { + h_x[i * k + i % n] = s; + } + for (int j = 0; j < n; j++) { + h_y[j * k + j] = s; + } + auto x = raft::make_device_matrix(handle, m, k); + auto y = raft::make_device_matrix(handle, n, k); + auto x_norm = raft::make_device_vector(handle, m); + auto y_norm = raft::make_device_vector(handle, n); + auto output = raft::make_device_vector(handle, m); + raft::update_device(x.data_handle(), h_x.data(), h_x.size(), stream.get()); + raft::update_device(y.data_handle(), h_y.data(), h_y.size(), stream.get()); + raft::linalg::rowNorm( + x_norm.data_handle(), x.data_handle(), k, m, stream.get()); + raft::linalg::rowNorm( + y_norm.data_handle(), y.data_handle(), k, n, stream.get()); + + cuvs::distance::detail::Top1nnTuning tuning{}; + for (auto backend : {Backend::Cutlass, Backend::Unfused}) { + if (!cuvs::distance::detail::is_top_1_nn_backend_available( + backend, x.data_handle(), y.data_handle(), m, n, k, DistanceType::L2Expanded)) { + continue; + } + const auto plan = cuvs::distance::probe_top_1_nn( + handle, x.data_handle(), y.data_handle(), m, n, k, tuning, DistanceType::L2Expanded, backend); + ASSERT_TRUE(plan.available); + auto workspace = raft::make_device_vector(handle, plan.workspace_bytes); + cuvs::distance::top_1_nn(handle, + output.data_handle(), + x.data_handle(), + y.data_handle(), + x_norm.data_handle(), + y_norm.data_handle(), + m, + n, + k, + tuning, + workspace.data_handle(), + workspace.size(), + false, + true, + true, + DistanceType::L2Expanded, + 0.0f, + plan); + std::vector h_out(m); + raft::update_host(h_out.data(), output.data_handle(), m, stream.get()); + raft::resource::sync_stream(handle, stream); + int n_wrong = 0; + for (int i = 0; i < m; i++) { + n_wrong += h_out[i].key != i % n; + } + EXPECT_EQ(n_wrong, 0) << "backend=" << static_cast(backend); + } +} + TEST(Top1nnPlan, ExplicitUnavailableBackendIsStrict) { raft::resources handle;