diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index a6438c1a8d..bd31ed6535 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 == "13.3.0")) 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] with: build_type: pull-request - script: "ci/test_cpp.sh ${{ matrix.shard }} ${{ matrix.num_shards }}" + script: "ci/test_cpp.sh" + matrix_type: nightly + matrix_filter: 'map(select(.ARCH == "arm64" and .GPU == "gh200")) | map(.CUDA_VER = "13.3.0" | .LINUX_VER = "ubuntu26.04" | .PY_VER = "3.14") | [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..f6b8e1acaf 100755 --- a/ci/test_cpp.sh +++ b/ci/test_cpp.sh @@ -7,7 +7,9 @@ set -euo pipefail # Usage: ci/test_cpp.sh [SHARD NUM_SHARDS] # With SHARD and NUM_SHARDS (1 <= SHARD <= NUM_SHARDS), only every NUM_SHARDS-th libcuvs test is # run, starting from test number SHARD, so the tests can be split across several CI jobs. +# shellcheck disable=SC2034 SHARD=${1:-1} +# shellcheck disable=SC2034 NUM_SHARDS=${2:-1} . /opt/conda/etc/profile.d/conda.sh @@ -58,16 +60,29 @@ 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})" +# [CI ONLY] Repeat the IVF-RaBitQ tests to reproduce a flaky failure (dim = 1, host input) and +# to verify the empty-list fix. pushd "$CONDA_PREFIX"/bin/gtests/libcuvs -timeout -v --signal=SIGINT --kill-after=60s 100m ctest -j8 --output-on-failure -I "${SHARD},,${NUM_SHARDS}" +rapids-logger "Run the IVF-RaBitQ empty-list regression tests repeatedly" +timeout -v --signal=SIGINT --kill-after=60s 15m ./NEIGHBORS_ANN_IVF_RABITQ_TEST \ + --gtest_filter='*empty_lists*' --gtest_brief=1 --gtest_repeat=50 > empty_lists.log 2>&1 +rapids-logger "empty lists exit code: $?" +grep -c "PASSED" empty_lists.log || true +grep -B2 -A20 "FAILED\|Failure" empty_lists.log | head -200 || true +rapids-logger "Run NEIGHBORS_ANN_IVF_RABITQ_TEST small-dim cases repeatedly" +timeout -v --signal=SIGINT --kill-after=60s 20m ./NEIGHBORS_ANN_IVF_RABITQ_TEST \ + --gtest_filter='IvfRabitq/f32_f32_i64.*/1:IvfRabitq/f32_f32_i64.*/2:IvfRabitq/f32_f32_i64.*/3' \ + --gtest_brief=1 --gtest_repeat=1000 > small_dims.log 2>&1 +rapids-logger "small dims exit code: $?" +grep -c "PASSED" small_dims.log || true +grep -B2 -A40 "FAILED\|Failure\|DEGENERATE\|before subtract" small_dims.log | head -400 || true +grep -oE "dim=[0-9]+ n_lists=[0-9]+ streaming=[0-9] residency=[0-9] on_device=[0-9] min_list=[0-9]+ empty_lists=[0-9]+" small_dims.log | sed -E 's/ n_lists=[0-9]+//; s/ on_device=[0-9]//' | sort | uniq -c | sort -k2 | head -100 || true +rapids-logger "Run full NEIGHBORS_ANN_IVF_RABITQ_TEST repeatedly" +timeout -v --signal=SIGINT --kill-after=60s 30m ./NEIGHBORS_ANN_IVF_RABITQ_TEST \ + --gtest_filter='-*empty_lists*' --gtest_brief=1 --gtest_repeat=30 > full.log 2>&1 +rapids-logger "full exit code: $?" +grep -c "PASSED" full.log || true +grep -B2 -A40 "FAILED\|Failure\|DEGENERATE\|before subtract" full.log | head -400 || true popd rapids-logger "Test script exiting with value: $EXITCODE" diff --git a/cpp/src/neighbors/ivf_rabitq.cu b/cpp/src/neighbors/ivf_rabitq.cu index ad8ea98117..6b39e79f43 100644 --- a/cpp/src/neighbors/ivf_rabitq.cu +++ b/cpp/src/neighbors/ivf_rabitq.cu @@ -12,11 +12,17 @@ #include "detail/ann_utils.cuh" #include +#include #include #include #include #include +#include +#include +#include +#include + namespace cuvs::neighbors::ivf_rabitq { namespace detail { @@ -163,6 +169,70 @@ auto build(raft::resources const& handle, cuvs::cluster::kmeans::predict( handle, kmeans_params, dataset_const_view, centers_const_view, labels_view); } + // [DIAGNOSTIC] Report degenerate k-means results (empty or out-of-range lists). + { + std::vector h_labels(n_rows); + std::vector h_centers(size_t(params.n_lists) * dim); + std::vector h_train(size_t(n_rows_train) * dim); + raft::copy(h_labels.data(), labels.data(), n_rows, stream); + raft::copy(h_centers.data(), cluster_centers.data(), h_centers.size(), stream); + raft::copy(h_train.data(), trainset.data_handle(), h_train.size(), stream); + raft::resource::sync_stream(handle); + std::vector sizes(params.n_lists, 0); + size_t out_of_range = 0; + for (auto l : h_labels) { + if (l < params.n_lists) { + sizes[l]++; + } else { + out_of_range++; + } + } + size_t n_empty = std::count(sizes.begin(), sizes.end(), size_t{0}); + size_t min_sz = *std::min_element(sizes.begin(), sizes.end()); + size_t n_nan_c = std::count_if( + h_centers.begin(), h_centers.end(), [](float x) { return !std::isfinite(x); }); + if (dim <= 8 || n_empty > 0 || out_of_range > 0 || n_nan_c > 0) + fprintf(stderr, + "[RABITQ-DIAG] n_rows=%zu dim=%zu n_lists=%u streaming=%d residency=%d " + "on_device=%d min_list=%zu empty_lists=%zu out_of_range_labels=%zu " + "nonfinite_centers=%zu\n", + size_t(n_rows), + size_t(dim), + params.n_lists, + int(use_streaming), + int(dataset_residency), + int(dataset_on_device), + min_sz, + n_empty, + out_of_range, + n_nan_c); + if (n_empty > 0 || out_of_range > 0 || n_nan_c > 0) { + float tmin = INFINITY, tmax = -INFINITY; + size_t t_nonfinite = 0; + for (auto x : h_train) { + if (!std::isfinite(float(x))) { + t_nonfinite++; + } else { + tmin = std::min(tmin, float(x)); + tmax = std::max(tmax, float(x)); + } + } + fprintf(stderr, + "[RABITQ-DIAG] DEGENERATE: n_train=%zu train_min=%g train_max=%g " + "train_nonfinite=%zu\n", + n_rows_train, + tmin, + tmax, + t_nonfinite); + for (uint32_t c = 0; c < params.n_lists; c++) { + fprintf(stderr, "[RABITQ-DIAG] list %u size=%zu center[0..]=", c, sizes[c]); + for (IdxT j = 0; j < std::min(dim, 4); j++) { + fprintf(stderr, " %.9g", h_centers[size_t(c) * dim + j]); + } + fprintf(stderr, "\n"); + } + } + } } index index(handle, n_rows, dim, params.n_lists, params.bits_per_dim); diff --git a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu index e27d34557a..1290175f4c 100644 --- a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu +++ b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu @@ -719,38 +719,44 @@ void IVFGPU::construct_on_gpu_streaming(const float* host_data, cluster_idx++; } - if (batch_vectors == 0) break; // No more clusters to process + // Every cluster fits into a batch (checked above), so each batch makes progress. + RAFT_EXPECTS(cluster_idx > batch_start_cluster, "IVF-RaBitQ streaming batch is empty"); // ------------------------- // 11. Gather and transfer batch data to GPU using reordered PIDs // ------------------------- - // First, copy the reordered PIDs for this batch to host (reuse preallocated buffer) - raft::copy(batch_pids.data_handle(), d_flat_pids + batch_start_offset, batch_vectors, stream_); - raft::resource::sync_stream(handle_); + // A batch that only contains empty clusters has no data to transfer, but the clusters still + // need their rotated centroids below. + if (batch_vectors > 0) { + // First, copy the reordered PIDs for this batch to host (reuse preallocated buffer) + raft::copy( + batch_pids.data_handle(), d_flat_pids + batch_start_offset, batch_vectors, stream_); + raft::resource::sync_stream(handle_); - // Gather data on host using the reordered PIDs (reuse preallocated buffer) - // Use OpenMP parallel region to leverage persistent thread pool across batches + // Gather data on host using the reordered PIDs (reuse preallocated buffer) + // Use OpenMP parallel region to leverage persistent thread pool across batches #pragma omp parallel num_threads(num_threads) - { - size_t tid = omp_get_thread_num(); - size_t chunk_size = batch_vectors / num_threads; - size_t start = tid * chunk_size; - size_t end = (tid == num_threads - 1) ? batch_vectors : start + chunk_size; - - for (size_t i = start; i < end; ++i) { - std::memcpy(&batch_data.data_handle()[i * num_dimensions], - &host_data[batch_pids.data_handle()[i] * num_dimensions], - num_dimensions * sizeof(float)); + { + size_t tid = omp_get_thread_num(); + size_t chunk_size = batch_vectors / num_threads; + size_t start = tid * chunk_size; + size_t end = (tid == num_threads - 1) ? batch_vectors : start + chunk_size; + + for (size_t i = start; i < end; ++i) { + std::memcpy(&batch_data.data_handle()[i * num_dimensions], + &host_data[batch_pids.data_handle()[i] * num_dimensions], + num_dimensions * sizeof(float)); + } } - } - // OpenMP implicit barrier ensures all threads complete before continuing + // OpenMP implicit barrier ensures all threads complete before continuing - // Transfer batch to GPU (reuse preallocated buffer) - raft::copy(d_batch_data.data_handle(), - batch_data.data_handle(), - batch_vectors * num_dimensions, - stream_); + // Transfer batch to GPU (reuse preallocated buffer) + raft::copy(d_batch_data.data_handle(), + batch_data.data_handle(), + batch_vectors * num_dimensions, + stream_); + } // ------------------------- // 12. Quantize each cluster in the batch @@ -758,8 +764,8 @@ void IVFGPU::construct_on_gpu_streaming(const float* host_data, size_t batch_offset = 0; for (size_t c = batch_start_cluster; c < cluster_idx; ++c) { + // Empty clusters are quantized too, so that their rotated centroids are initialized. size_t cluster_size = h_cluster_meta[c].num; - if (cluster_size == 0) continue; const float* cur_centroid = device_centroids + c * num_dimensions; float* cur_rotated_c = d_rotated_centroids.data_handle() + c * num_padded_dim; diff --git a/cpp/src/neighbors/ivf_rabitq/gpu_index/quantizer_gpu.cu b/cpp/src/neighbors/ivf_rabitq/gpu_index/quantizer_gpu.cu index b35ac9fab9..7361994248 100644 --- a/cpp/src/neighbors/ivf_rabitq/gpu_index/quantizer_gpu.cu +++ b/cpp/src/neighbors/ivf_rabitq/gpu_index/quantizer_gpu.cu @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -520,12 +521,28 @@ void data_transformation_batch_opt(const float* d_data, // 5. Save the rotated centroid: copy CP into d_rotated_c. raft::copy(d_rotated_c, d_CP, D, stream); + // An empty list only needs its rotated centroid; a launch with zero blocks is invalid. + if (num_points == 0) { return; } + // 6. Launch the single FUSED kernel for subtract, normalize, and binarize. const unsigned int FusedBlockSize = 256; // A good default, can be tuned. dim3 gridDim(num_points); dim3 blockDim(FusedBlockSize); size_t sharedMemSize = FusedBlockSize * sizeof(float); + { // [DIAGNOSTIC] + cudaError_t pre = cudaPeekAtLastError(); + if (pre != cudaSuccess || num_points == 0 || num_points > (size_t{1} << 31)) { + fprintf(stderr, + "[RABITQ-DIAG] before subtract_normalize_binarize_Kernel: num_points=%zu D=%zu " + "DIM=%zu pending_error=%d (%s)\n", + num_points, + D, + DIM, + int(pre), + cudaGetErrorName(pre)); + } + } subtract_normalize_binarize_Kernel <<>>( d_XP, // Input: Rotated data @@ -637,6 +654,9 @@ void DataQuantizerGPU::quantize_batch_opt(const float* d_data, D, handle_); + // An empty list has no codes or factors to compute. + if (num_points == 0) { return; } + rabitq_codes_and_factors_fused(d_rotated_c, d_bin_XP.data_handle(), d_XP.data_handle(), @@ -714,12 +734,28 @@ void data_transformation_batch_opt_contiguous(const float* d_contiguous_data, // 5. Save the rotated centroid: copy CP into d_rotated_c. raft::copy(d_rotated_c, d_CP, D, stream); + // An empty list only needs its rotated centroid; a launch with zero blocks is invalid. + if (num_points == 0) { return; } + // 6. Launch the single FUSED kernel for subtract, normalize, and binarize. const unsigned int FusedBlockSize = 256; // A good default, can be tuned. dim3 gridDim(num_points); dim3 blockDim(FusedBlockSize); size_t sharedMemSize = FusedBlockSize * sizeof(float); + { // [DIAGNOSTIC] + cudaError_t pre = cudaPeekAtLastError(); + if (pre != cudaSuccess || num_points == 0 || num_points > (size_t{1} << 31)) { + fprintf(stderr, + "[RABITQ-DIAG] before subtract_normalize_binarize_Kernel: num_points=%zu D=%zu " + "DIM=%zu pending_error=%d (%s)\n", + num_points, + D, + DIM, + int(pre), + cudaGetErrorName(pre)); + } + } subtract_normalize_binarize_Kernel <<>>( d_XP, // Input: Rotated data @@ -758,6 +794,9 @@ void DataQuantizerGPU::quantize_batch_opt_contiguous(const float* d_contiguous_d D, handle_); + // An empty list has no codes or factors to compute. + if (num_points == 0) { return; } + rabitq_codes_and_factors_fused(d_rotated_c, d_bin_XP.data_handle(), d_XP.data_handle(), diff --git a/cpp/tests/neighbors/ann_ivf_rabitq.cuh b/cpp/tests/neighbors/ann_ivf_rabitq.cuh index 2755258600..7296e67ddf 100644 --- a/cpp/tests/neighbors/ann_ivf_rabitq.cuh +++ b/cpp/tests/neighbors/ann_ivf_rabitq.cuh @@ -396,4 +396,123 @@ inline auto var_search_mode_1_bit() -> test_cases_t #define INSTANTIATE(type, vals) \ INSTANTIATE_TEST_SUITE_P(IvfRabitq, type, ::testing::ValuesIn(vals)); /* NOLINT */ +/* Empty lists */ + +enum class dataset_input { device, host, host_streaming }; + +struct ivf_rabitq_empty_lists_inputs { + uint32_t dim; + dataset_input input; +}; + +inline auto operator<<(std::ostream& os, const ivf_rabitq_empty_lists_inputs& p) -> std::ostream& +{ + const char* input = p.input == dataset_input::device ? "device" + : p.input == dataset_input::host ? "host" + : "host_streaming"; + return os << "ivf_rabitq_empty_lists_inputs {.dim = " << p.dim << ", .input = " << input << "}"; +} + +/** + * The dataset has fewer distinct rows than there are lists. Identical rows are always assigned to + * the same list, so k-means is guaranteed to leave some lists empty. The index must still build + * from device and host input, and search must still find the nearest rows. + */ +class ivf_rabitq_empty_lists_test : public ::testing::TestWithParam { + protected: + void run() + { + constexpr int64_t kRows = 4096; + constexpr int64_t kDistinct = 8; + constexpr int64_t kQueries = 2 * kDistinct; + constexpr int64_t kK = 10; + constexpr uint32_t kLists = 32; + constexpr float kSpacing = 0.25f; + constexpr float kOffset = 0.1f; + + auto ps = GetParam(); + auto dim = static_cast(ps.dim); + auto stream = raft::resource::get_cuda_stream(handle_); + + // Row i is the point kSpacing * (i % kDistinct) on the diagonal. Query q is offset from the + // distinct row q % kDistinct by +-kOffset in every dimension, so that row is its unique nearest + // distinct row at squared L2 distance kOffset^2 * dim. + auto h_dataset = raft::make_host_matrix(kRows, dim); + for (int64_t i = 0; i < kRows; i++) { + for (int64_t j = 0; j < dim; j++) { + h_dataset(i, j) = kSpacing * static_cast(i % kDistinct); + } + } + auto h_queries = raft::make_host_matrix(kQueries, dim); + for (int64_t i = 0; i < kQueries; i++) { + for (int64_t j = 0; j < dim; j++) { + h_queries(i, j) = + kSpacing * static_cast(i % kDistinct) + (i < kDistinct ? kOffset : -kOffset); + } + } + auto d_dataset = raft::make_device_matrix(handle_, kRows, dim); + auto d_queries = raft::make_device_matrix(handle_, kQueries, dim); + raft::copy(d_dataset.data_handle(), h_dataset.data_handle(), h_dataset.size(), stream); + raft::copy(d_queries.data_handle(), h_queries.data_handle(), h_queries.size(), stream); + + index_params ipams; + ipams.n_lists = kLists; + if (ps.input == dataset_input::host_streaming) { + ipams.force_streaming = true; + ipams.streaming_batch_size = kRows; + } + auto idx = ps.input == dataset_input::device + ? build(handle_, ipams, raft::make_const_mdspan(d_dataset.view())) + : build(handle_, ipams, raft::make_const_mdspan(h_dataset.view())); + + // Probe every list, because an empty list's centroid can be closer to a query than the + // centroid of the list that holds its neighbors. + search_params sparams; + sparams.n_probes = kLists; + auto neighbors = raft::make_device_matrix(handle_, kQueries, kK); + auto distances = raft::make_device_matrix(handle_, kQueries, kK); + search(handle_, + sparams, + idx, + raft::make_const_mdspan(d_queries.view()), + neighbors.view(), + distances.view()); + + std::vector h_neighbors(kQueries * kK); + std::vector h_distances(kQueries * kK); + raft::copy(h_neighbors.data(), neighbors.data_handle(), h_neighbors.size(), stream); + raft::copy(h_distances.data(), distances.data_handle(), h_distances.size(), stream); + raft::resource::sync_stream(handle_); + + // Every row has kRows / kDistinct > kK copies, so all kK neighbors are copies of the nearest + // distinct row. The next nearest distinct row is (kSpacing - kOffset)^2 * dim away. + const float expected_dist = kOffset * kOffset * static_cast(dim); + for (int64_t q = 0; q < kQueries; q++) { + for (int64_t j = 0; j < kK; j++) { + auto neighbor = h_neighbors[q * kK + j]; + auto dist = h_distances[q * kK + j]; + ASSERT_TRUE(neighbor >= 0 && neighbor < kRows) + << ps << ": query " << q << ", neighbor " << j << " = " << neighbor; + ASSERT_EQ(neighbor % kDistinct, q % kDistinct) + << ps << ": query " << q << ", neighbor " << j << " = " << neighbor; + ASSERT_NEAR(dist, expected_dist, 0.1f * expected_dist) + << ps << ": query " << q << ", neighbor " << j << " = " << neighbor; + } + } + } + + raft::resources handle_; +}; + +inline auto empty_lists_inputs() -> std::vector +{ + std::vector xs; + for (uint32_t dim : {1, 2, 7, 64, 100}) { + for (auto input : {dataset_input::device, dataset_input::host, dataset_input::host_streaming}) { + xs.push_back({dim, input}); + } + } + return xs; +} + } // namespace cuvs::neighbors::ivf_rabitq diff --git a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu index 42bf84e6a3..28771bd785 100644 --- a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu +++ b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu @@ -18,4 +18,9 @@ INSTANTIATE(f32_f32_i64, defaults() + small_dims() + big_dims() + var_n_probes() + var_k() + var_bits_per_dim() + var_search_mode() + var_search_mode_1_bit()); +TEST_P(ivf_rabitq_empty_lists_test, build_search) { this->run(); } // NOLINT +INSTANTIATE_TEST_SUITE_P(IvfRabitq, + ivf_rabitq_empty_lists_test, + ::testing::ValuesIn(empty_lists_inputs())); // NOLINT + } // namespace cuvs::neighbors::ivf_rabitq