[FEA] Codepacking for IVF-flat #1632
Conversation
tarang-jain
added
2 - In Progress
cjnolet
added
non-breaking
|
To resolve style errors, you could install pre-commit hooks. |
There was a problem hiding this comment.
Thanks @tarang-jain for the PR! The new pack/unpack funcions look good, but we shall have test for that. Also, I see a lot of commented out code here, please let me know what is the plan with that.
| * @param[in] offset how many records to skip before writing the data into the list | ||
| */ | ||
| template <typename T> | ||
| __host__ __device__ void pack_1_interleaved( |
There was a problem hiding this comment.
Do we have a reason to add the _interleaved suffix? Faiss has a function called (un)pack_1.
Alternativele, IVF-PQ has functions (un)pack, that can take one or many vectors at once. We shall aim for consistency with one of these.
There was a problem hiding this comment.
Alright, (un)pack_1 sounds like the better naming convention to avoid confusion. I was also told by @cjnolet that the current IVF-PQ CodePacker needs some refactoring and the names of the IVF-PQ functions can be changed to (un)pack_1 then.
There was a problem hiding this comment.
IVF-PQ functions can unpack multiple vectors at a time, I would not name them (un)pack_1
There was a problem hiding this comment.
The way to think of ivf_flat_codepacker.hpp is that the codepacker a public API and is the smallest atomic unit (packing and unpacking one single row). This can now be used in host and device functions alike -- for packing and unpacking multiple records either through a kernel, or through a CPU for loop.
There was a problem hiding this comment.
The helper pack_full_list is just an example of using the CodePacker for packing a full list.
There was a problem hiding this comment.
LGTM. I think we could/should consolidate the kernels into map/map_offset thunks eventually but given that this is going to enable the faiss integration in the meantime, I think we should be good to go for 23.08. Thanks for all your hard work on this @tarang-jain!
| @@ -113,6 +113,7 @@ struct list_spec { | |||
| /** Determine the extents of an array enough to hold a given amount of data. */ | |||
| constexpr auto make_list_extents(SizeT n_rows) const -> list_extents | |||
| { | |||
| // return make_extents<SizeT>(round_up_safe(n_rows, this->align_min), dim); | |||
| dim3 blocks(div_rounding_up_safe<uint32_t>(n_rows, kBlockSize), 1, 1); | ||
| dim3 threads(kBlockSize, 1, 1); | ||
| auto stream = resource::get_cuda_stream(res); | ||
| pack_interleaved_list_kernel<<<blocks, threads, 0, stream>>>( |
There was a problem hiding this comment.
Any reason we couldn't just use a map() or map_offset() for this?
There was a problem hiding this comment.
there is no particular reason. this can indeed be put inside a map_offset.
cjnolet
dismissed
tfeher’s stale review
Tamas is on PTO and Tarang has addressed the review items. We need this functionality to release FAISS in August and we can follow up in 23.10 if needed.
Support reading interleaved lists from a packed non-interleaved dense IVF-list matrix and vice versa.