Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2726
Old versions of clang (prior to 14) do not generate FMAs on x86 unless this option is used. On the other hand, ARM does not support the needed pragma flag.
Reviewed By: DenisYaroshevskiy
Differential Revision: D43503265
fbshipit-source-id: a024dd221288e44d4e2ade2f5db2c4402e26ff3d
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2725
fvec_inner_product_ref() is no longer needed, using fvec_inner_product() instead in all the cases
fvec_norm_L2sqr_ref() is no longer needed, using fvec_norm_L2sqr() instead in all the cases
Reviewed By: algoriddle
Differential Revision: D43472582
fbshipit-source-id: 2d99a022d420092aed2adfd57e12be5b3e652944
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2721
FAISS_PRAGMA_IMPRECISE_* macros were modified:
* Disabled ones on clang on arm, because it does not support `_Pragma("float_control(precise, off)")`
* Added missing pragma for the GCC compiler.
Reviewed By: alexanderguzhva
Differential Revision: D43437450
fbshipit-source-id: cec8042c3c8c7147ae7e2ffa1ac9e2232c8f1a92
Summary:
Pragmas:
turns out GCC does not support pragma float_control
I have no way to test it locally but I tried on godbolt: https://godbolt.org/z/zTzf7jd7c
I will put these pragmas in folly later but for now - let's just solve the problem.
Test:
The diff: D43353199 that produced slightly different floating point results broke this test.
After looking into it I don't believe that the original test data had any meaning and was just a dump of the output.
Accoding to mdouze the previous test data was questionable and using random numbers would be preferrable, so I did.
Reviewed By: alexanderguzhva
Differential Revision: D43393684
fbshipit-source-id: 691397302b58a98e1ccbe32feffa7aaeee96accd
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2696
Add fvec_L2sqr_ny_y_transposed() call which is used by ProductQuantizer in case of pq.sync_transposed_centroids() is enabled.
Reviewed By: mdouze
Differential Revision: D43057940
fbshipit-source-id: b82a0153fb10f1558c7db04e32d690b2ee7b9144
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2693
A new kernel for fvec_op_ny_D4 for x64 that reorganizes read ops.
This diff also introduces transpose microkernels for D2 and D8 for AVX2 for future needs.
Reviewed By: mdouze
Differential Revision: D43057318
fbshipit-source-id: 37b2e18139905fcdac69e4603038a15725a730dd
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2682
IndexShards normally sees the indexes as opaque, so there is no way to factrorize the coarse quantizer.
This diff introduces IndexIVFShards that handles IVF indexes with a common quantizer so that the quantization is computed only once.
Reviewed By: alexanderguzhva
Differential Revision: D42781513
fbshipit-source-id: 441316eff4c1ba0468501c456af9194ea5f042d6
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2683
Adds a small function and a test to demonstrate how to crop a RCQ to its first quantizers.
Also adds a SearchParameters to set the beam_factor of a RCQ.
Reviewed By: alexanderguzhva
Differential Revision: D42842378
fbshipit-source-id: b522269ce983cddc50c3d0e18b541954c5d7001d
Summary: Added the required terms of use, privacy policy and copyright statement for Meta OSS projects.
Reviewed By: algoriddle
Differential Revision: D43056233
fbshipit-source-id: 8541edd587107bf3129ff7f8b29ec26c016f779f
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2685
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2623
This diff is an attempt at dealing with limitations when I originally wrote the Faiss GPU code back in 2016. Namely, 64 bit integer arithmetic was really slow on the GPU (this was also the impetus for me adding the 32 vs 64 bit indexing logic to PyTorch GPU), and GPUs at the time were very memory limited.
However, there have been enough cases where people are attempting to use Faiss GPU to handle larger datasets, especially including > 2 GB flat indexes, which sometimes worked (total number of elements in a tensor being < 2^31 - 1 but the byte size could be larger) or sometimes didn't work (when considering byte offsets not word sizes). This needed fixing.
This diff changes the `Tensor`, `HostTensor` and `DeviceTensor` classes to default to using `faiss::idx_t` as the index type (64 bit), and also changes many other places in the GPU code to follow the same data type. Much of the code being changed follows from the type propagation coming from the `Tensor` class and sub-classes.
`GpuIndexIVF.setNumProbes` and `GpuIndexIVF.getNumProbes` functions have been removed, just set/get `nprobe` on the object as need (like the CPU side).
The net result is that flat indexes can contain > 2 GB of data, and individual IVF lists can similarity contain > 2 GB of data, as well as any tensor dimension can be over 2G (excepting vector dimension; see limitations below).
Notes:
- performance for large `k` (>= 512) or large `nprobe` (>= 512) deteriorates due to low occupancy in the k-selection kernels due to increased register pressure (we now maintain all indices as int64 rather than int32, thus requiring double the registers). However, based on tests this doesn't hurt more than 10-20% or so at these large sizes. Small `k`/`nprobe` has no significant effect, and in some cases is even faster (depending upon if we needed to produce int64 indices but computations were done internally using int32).
- `k`, like with CPU Faiss, is an int parameter everywhere still.
- `d`, like with CPU Faiss, is an int parameter everywhere except for places where dimensionality is inferred from a tensor dimension (in which case it is naturally `idx_`).
- `nprobe`, like `IndexIVF`, is `size_t` at the top level (`GpuIndexIVF` class member) but internally to the code it is either `int` or `idx_t` depending upon the context (e.g., if `nprobe` is extracted from a tensor dimension, it is already `idx_t`).
In the Faiss GPU code however, `k` and `nprobe` are still limited to 2048 maximum at present.
Reviewed By: alexanderguzhva
Differential Revision: D40494090
fbshipit-source-id: 8581aff86ff25091b47e59f749adebebf2d3b160
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2691
`faiss::gpu::DeviceVector<T>` (the GPU analogue of `std::vector`) is used in Faiss GPU to store IVF list data, among other things. `thrust::device_vector` could not be used at the time of the library's authorship due to lack of certain features (stream usage from what I recall).
Like `std::vector`, it involves a heuristic used to determine what the new vector capacity should be in order to allow for a particular number of elements. It previously has always rounded up allocation sizes to the next highest power of 2, thus incremental appending to a DeviceVector amortizes the cost of memory (re-)allocation and only infrequently results in full re-allocation.
However, this doubling behavior was not limited. While at small sizes of individual IVF lists, this doubling is not that much of an issue (and there is a reclaim memory function on indices to reclaim unused memory resulting from this amortization). However, when operating with data sets close to GPU memory limits, this doubling is dangerous, as attempting to allocate, say, a GpuIndexFlat with 4.x GB of vectors on a 8 GB GPU could run out of memory.
Now we introduce a new heuristic:
- below 4 M list sizes, we always double (as before);
- otherwise, below 128 M list sizes, we only increase by 1.25x;
- otherwise, the allocation made by DeviceVector is exactly sized to the size requested by the user.
Reviewed By: mdouze
Differential Revision: D42951571
fbshipit-source-id: 462d8af582725869ab5642b81e27db0e71774c6e
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2676
This is a cosmetic diff where default values for scalar fields are moved to the .h file where the object is declared rather than in the object constructor. The advantage is that it is easier to read the default value of a parameter from just the class.
Reviewed By: alexanderguzhva
Differential Revision: D42722205
fbshipit-source-id: 9d9bc4641068f6d6233f60f0a3a16ab793c94bb8
Summary:
IVF GPU indexes now support CPU quantizers.
In order for the the auto-tuning to support accessing the parameters of the quantizers, GpuAutoTune needs to be adapted.
Reviewed By: algoriddle
Differential Revision: D42633348
fbshipit-source-id: ef472855aa882ccde9d878ae09782204045e38c5
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2638
This diff is a more streamlined way of searching IVF indexes with precomputed clusters.
This will be used for experiments with hybrid CPU / GPU search.
Reviewed By: algoriddle
Differential Revision: D41301032
fbshipit-source-id: a1d645fd0f2bf806454dfd04971edc0a6200d20d
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2647
The `max_codes` search parameter for IVF indexes limits the number of distance computations that are performed.
Previously, the number of distance computations could exceed max_codes because inverted lists were scanned completely.
This diff changed this to scan the beginning of the last inverted list to reach `max_codes` exactly.
Reviewed By: alexanderguzhva
Differential Revision: D42367593
fbshipit-source-id: 67c88b93a407ab271397e913c5fa17104f4274c3
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2650
LocalSearchQuantizer::icm_encode_step was optimized:
* The loop orders were changes
* Memory allocations were optimized
* Different order for accessing 'binaries' tensor, because it is symmetric
* SIMD for the lookup of the best code.
Results are unchanged.
Also, fixes incorrect test in test_local_search_quantizer.py
Reviewed By: mdouze
Differential Revision: D42352124
fbshipit-source-id: bf7e349f2123e6ee99e0776cf15ad5fc1cf2439a
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2651
Modify approximate heap facility to maintain the correct order as if the input elements were pushed into the heap sequentially
Reviewed By: mdouze
Differential Revision: D42373532
fbshipit-source-id: 477dc8acd2567157e2b99076a566326821021c8c
Summary:
In ```cmp_with_scann.py```, we will save npy file for base and query vector file and gt file. However, we will only do this while the lib is faiss, if we directly run this script with scann lib it will complain that file does not exsit.
Therefore, the code should be refactored to save npy file from the beginning so that nothing will go wrong.
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2573
Reviewed By: mdouze
Differential Revision: D42338435
Pulled By: algoriddle
fbshipit-source-id: 9227f95e1ff79f5329f6206a0cb7ca169185fdb3
Summary:
There are two different names for the same thing and name inconsistency between the function declaration and definition which uses `M2` in the function declaration and uses `nsq` in the function definition. It may cause confusion, so it's better to unify them.
Signed-off-by: zh Wang <[email protected]>
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2542
Reviewed By: mdouze
Differential Revision: D42338110
Pulled By: algoriddle
fbshipit-source-id: d386e00fb8d8904051dd676bae1e7f4702a172d9
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2633
The core idea.
Instead of putting every element of the dataset into MaxHeap, split the dataset into buckets and let every bucket track elements min-1, min-2 or min-3 distances.
Applied to ResidualQuantizer class for vector codec purposes.
An example
```
rq.approx_topk_mode = faiss.APPROX_TOPK_BUCKETS_B16_D2
```
Reviewed By: mdouze
Differential Revision: D42044398
fbshipit-source-id: 43169026476650442806a31d1c1aa2d5d5028e65
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2625
This diff introduces a new abstraction for the code layouts that are not simply flat one after another.
The packed codes are assumed to be packed together in fixed-size blocks. Hence, code `#i` is stored at offset `i % nvec` of block `floor(i / nvec)`. Each block has size `block_size`.
The `CodePacker` object takes care of the translation between packed and flat codes. The packing / unpacking functions are virtual functions now, but they could as well be inlined for performance.
The `CodePacker` object makes it possible to do manipulations onarrays of codes (including inverted lists) in a uniform way, for example merging / adding / updating / removing / converting to&from CPU.
In this diff, the only non-trivial CodePacker implemnted is for the FastScan code. The new functionality supported is merging IVFFastScan indexes.
Reviewed By: alexanderguzhva
Differential Revision: D42072972
fbshipit-source-id: d1f8bdbcf7ab0f454b5d9c37ba2720fd191833d0
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2626
Improves the performance of fused kernels used in exhaustive_L2sqr_blas() call. The code parameters are tweaked to utilize concurrently working ALU and LOAD/STORE units in a single CPU core.
The parameters values were tweaked for AVX2 and AVX512 kernels, but ARM ones remained unchanged because there was no way for me to benchmark the changes. Please feel free to alter ones if you have access to an ARM machine.
Reviewed By: mdouze
Differential Revision: D42079875
fbshipit-source-id: f1f7dc1759dbad57827d1c1b1b2b399322b33df0
Summary:
Pull Request resolved: https://github.com/facebookresearch/faiss/pull/2631
The pytorch in fbcode complains about `storage()` saying it is deprecated and we need to move to UntypedStorage `_storage()`, while github CI is using an older version of pytorch where `_storage()` doesn't exist.
As it is only a warning not an error in fbcode, revert to the old form, but we'll likely have to change to `_storage()` eventually.
Reviewed By: alexanderguzhva
Differential Revision: D42107029
fbshipit-source-id: 699c15932e6ae48cd1c60ebb7212dcd9b47626f6
Summary:
This diff fixes four separate issues:
- Using the pytorch bridge produces the following deprecation warning. We switch to `_storage()` instead.
```
torch_utils.py:51: UserWarning: TypedStorage is deprecated. It will be removed in the future and UntypedStorage will be the only storage class. This should only matter to you if you are using storages directly. To access UntypedStorage directly, use tensor._storage() instead of tensor.storage()
x.storage().data_ptr() + x.storage_offset() * 4)
```
- The `storage_offset` for certain types was wrong, but this would only affect torch tensors that were a view into a storage that didn't begin at the beginning.
- The `reconstruct_n` numpy pytorch bridge function allowed passing `-1` for `ni` which indicated that all vectors should be reconstructed. The torch bridge didn't follow this and throw an error:
```
TypeError: torch_replacement_reconstruct_n() missing 2 required positional arguments: 'n0' and 'ni'
```
- Choosing values in the range (1024, 2048] for `k` or `nprobe` were broken in D37777979; this is now fixed again.
Reviewed By: alexanderguzhva
Differential Revision: D42041239
fbshipit-source-id: c7d9b4aba63db8ac73e271c8ef34e231002963d9
Iurii Makarov