Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Forward-merge branch-23.06 to branch-23.08 #1545

Merged
merged 9 commits into from
May 23, 2023
Merged

Conversation

GPUtester
Copy link
Contributor

Forward-merge triggered by push to branch-23.06 that creates a PR to keep branch-23.08 up-to-date. If this PR is unable to be immediately merged due to conflicts, it will remain open for the team to manually merge.

mdoijade and others added 9 commits May 19, 2023 16:59
-- as the kernel arch given by the cudaFuncAttribute ptxVersion depends on what archs the kernel was compiled for
we should renam kernel_runtime_arch() as kernel_virtual_arch().
-- accordingly update comments to reflect this.

Authors:
  - Mahesh Doijade (https://github.com/mdoijade)
  - Corey J. Nolet (https://github.com/cjnolet)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1536
The input validation code wasn't being triggered for the python bfknn api, causing invalid output when passed col-major inputs. Fix.

Authors:
  - Ben Frederickson (https://github.com/benfred)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1537
This PR updates CAGRA::search to support `uint64_t` in the index data type. This update is required to implement the RAFT ANN benchmark for CAGRA.

Authors:
  - tsuki (https://github.com/enp1s0)
  - Ben Frederickson (https://github.com/benfred)
  - Tamas Bela Feher (https://github.com/tfeher)

Approvers:
  - Tamas Bela Feher (https://github.com/tfeher)

URL: #1514
A prototype that introduces a per-sample filtering for IVFPQ search. Please feel free to use it as a foundation for the future change, if appropriate, because the code is functional, but is not super clean-and-neat.

The diff introduces a template parameter called `SampleFilterT`. An instance is expected 
* to be `SampleFilterT()` constructible (which was mostly needed to define a default behavior in the form of `SampleFilterT sample_filter=SampleFilterT()`, see below)
* to provide a `inline __device__ bool operator(...)`  that returns `true` is a given sample is valid for being used against a given query in IVFPQ search

The default filter (that I set as a default one in certain facilities in the form of `typename SampleFilterT = NoneSampleFilter` in order not to modify way to many files) allows all samples to be used:
```
struct NoneSampleFilter {
  inline __device__ __host__ bool operator()(
    // query index
    const uint32_t query_ix,
    // the current inverted list index
    const uint32_t cluster_ix,
    // the index of the current sample inside the current inverted list
    const uint32_t sample_ix
  ) const {
    return true;
  }
};
```
Here `__host__` is needed for a CPU-based testing only.
Also, I've provided an implementation of `BitMaskSampleFilter` that allows to filter samples based on a bit mask, as an example. The implementation was tested in the semi-production environment.

All the filter-related files were added to `cpp/include/raft/neighbors/detail/sample_filter.cuh`.

I did not change the default `ivf_pq_search()` method remains unchanged, but one more `ivf_pq_search_with_filtering()` method with an additional template argument `SampleFilterT` and one more input parameter was introduced.
```
template <typename T, typename IdxT, typename SampleFilterT>
void search_with_filtering(raft::device_resources const& handle,
            const raft::neighbors::ivf_pq::search_params& params,
            const index<IdxT>& idx,
            const T* queries,
            uint32_t n_queries,
            uint32_t k,
            IdxT* neighbors,
            float* distances,
            rmm::mr::device_memory_resource* mr = nullptr,
            SampleFilterT sample_filter = SampleFilterT());
```

All the current instantiations use `NoneSampleFilter` only.

I've used `SampleFilterT sample_filter` parameters passing instead of `const SampleFilterT sample_filter` in the function calls in order to be able to add some debugging facilities to a filter and with the hope that the compiler is smart enough to understand the de-facto constness if needed.

The filter does not take a computed distance score into account by design, thus the current implementation cannot have a distance threshold. This can be easily changed, if appropriate.

It is still questionable to me whether this filtering needs to be injected right inside the search kernel instead of doing post-processing, please let me know if you have any thoughts on the topic.

I'm happy to address the comments.

Thanks.

Authors:
  - Alexander Guzhva (https://github.com/alexanderguzhva)
  - Artem M. Chirkin (https://github.com/achirkin)
  - Corey J. Nolet (https://github.com/cjnolet)

Approvers:
  - Artem M. Chirkin (https://github.com/achirkin)
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1513
…uments (#1539)

closes #1357

breaking change: the type of argument mdspans has slightly changed (second template parameter fixed to `uint32_t`)

Authors:
  - Artem M. Chirkin (https://github.com/achirkin)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1539
Remove the device_resources.cuh include from linalg::map.

For the implicit integration, I don't have libraries like cusolver in the CI environment, and the build is currently failing with errors like

```
/project/_skbuild/linux-x86_64-3.8/cmake-build/_deps/raft-src/cpp/include/raft/core/device_resources.hpp:31:10:
       fatal error: cusolverDn.h: No such file or directory
```

(from https://github.com/benfred/implicit/actions/runs/5033022104/jobs/9026999533?pr=656)

Fix by not including device_resources here

Authors:
  - Ben Frederickson (https://github.com/benfred)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1540
@GPUtester GPUtester requested review from a team as code owners May 23, 2023 18:03
@GPUtester GPUtester merged commit 99655ea into branch-23.08 May 23, 2023
@GPUtester
Copy link
Contributor Author

SUCCESS - forward-merge complete.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

8 participants