[Feat]: find API with missing parameters #183
Conversation
Documentation previewhttps://nvidia-merlin.github.io/HierarchicalKV/review/pr-183 |
rhdong
changed the title
| : founds(founds_) {} | ||
|
|
||
| __forceinline__ __device__ void operator()(const int idx, const K /*key*/, | ||
| const bool found) { |
There was a problem hiding this comment.
Please format the code by referring the guidance
| const K* __restrict keys, V** __restrict values, S* __restrict scores, | ||
| K* __restrict missed_keys, int* __restrict missed_indices, | ||
| int* __restrict missed_size, int* __restrict dst_offset, size_t N) { | ||
| FoundFunctorV2<K> found_functor(missed_keys, missed_indices, missed_size); |
There was a problem hiding this comment.
atomicadd in FoundFunctorV2 runs on all warps when the key is not found, which can cause performance degradation when there are lots of missed key. Isn't it better to split the missed_keys size used by found_functor from n to the bucket size, and then merge according to missed_size at last kernel finishing?
There was a problem hiding this comment.
Good point, I'll work on refining this code.
There was a problem hiding this comment.
Good point, I'll work on refining this code.
What was the result? Does this optimization show up in performance tests?
There was a problem hiding this comment.
I apologize for the late reply, but I have verified your suggestion. Due to the additional calculations brought by the bucket separation, the benchmark results did not show performance improvement.
|
/blossom-ci |
4 participants
In our inference service, we found that the find API will have better performance by directly writing the missed keys when the hit rate is high, which it usually is.