Allow changing the reserved keys #191
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
jq
force-pushed
the
r3
branch
5 times, most recently
from
aafcbd2
to
1789f38
Compare
Documentation previewhttps://nvidia-merlin.github.io/HierarchicalKV/review/pr-191 |
rhdong
reviewed
May 14, 2024
include/merlin/types.cuh
Outdated
| __constant__ uint64_t LOCKED_KEY = UINT64_C(0xFFFFFFFFFFFFFFFD); | ||
| template <class K> | ||
| __forceinline__ __device__ bool IS_RESERVED_KEY(K key) { | ||
| return key == EMPTY_KEY || key == RECLAIM_KEY || key == LOCKED_KEY; |
There was a problem hiding this comment.
Considering the performance(IS_RESERVED_KEY will be called anytime on each threads), it would be better to use & operator, so we can use one reserved mask to indicate the reserved keys space like RESERVED_KEY_MASK = UINT64_C(0xFFFFFFFFFFFFFFFC);, by default, it should be 0x3, if we wanna reserve 0xFFFFFFFFFFFFFFxF(x = E, C, D), we can let the user give a mask valued 0x30.
jq
force-pushed
the
r3
branch
5 times, most recently
from
a1b3116
to
1fc4abb
Compare
jq
force-pushed
the
r3
branch
13 times, most recently
from
63f502e
to
c950300
Compare
jq
force-pushed
the
r3
branch
4 times, most recently
from
27f39f7
to
767c84f
Compare
jq
commented
May 16, 2024
include/merlin/types.cuh
Outdated
| uint64_t lockedKey = emptyKey & ~(UINT64_C(2) << index); | ||
| EMPTY_KEY_CPU = emptyKey; | ||
|
|
||
| printf("reserved keys are emptyKey, reclaimKey, lockedKey and reservedKeyMask2\n"); |
rhdong
reviewed
May 16, 2024
include/merlin/types.cuh
Outdated
| printf("lockedKey: %lx, reservedKeyMask1: %lx\n", lockedKey, reservedKeyMask1); | ||
| printf("reservedKeyMask2: %lx, vacantKeyMask1: %lx\n", reservedKeyMask2, vacantKeyMask1); | ||
|
|
||
| cudaMemcpyToSymbol(EMPTY_KEY, &emptyKey, sizeof(uint64_t)); |
rhdong
reviewed
May 16, 2024
include/merlin_hashtable.cuh
Outdated
| @@ -98,6 +98,7 @@ struct HashTableOptions { | |||
| int device_id = -1; ///< The ID of device. | |||
| bool io_by_cpu = false; ///< The flag indicating if the CPU handles IO. | |||
| bool use_constant_memory = false; ///< reserved | |||
| int reserved_key_index = 0; ///< The binary index of reserved key. | |||
There was a problem hiding this comment.
Need more detail comment which allows user know the configuration effect exactly. For example: index=0 means which keys will be reserved.
rhdong
reviewed
May 16, 2024
tests/merlin_hashtable_test.cc.cu
Outdated
| @@ -3437,6 +3440,12 @@ void test_bucket_size(bool load_scores = true) { | |||
| CudaCheckError(); | |||
| } | |||
|
|
|||
|
|
|||
| TEST(MerlinHashTableTest, test_basic) { | |||
rhdong
reviewed
May 16, 2024
tests/merlin_hashtable_test.cc.cu
Outdated
| @@ -145,6 +145,8 @@ void test_basic(size_t max_hbm_for_vectors) { | |||
| options.dim = DIM; | |||
| options.max_bucket_size = BUCKET_MAX_SIZE; | |||
| options.max_hbm_for_vectors = nv::merlin::GB(max_hbm_for_vectors); | |||
| options.reserved_key_index = 2; | |||
There was a problem hiding this comment.
Please add test cases like this for the other API.
rhdong
reviewed
May 16, 2024
README.md
Outdated
| - The keys of `0xFFFFFFFFFFFFFFFC`, `0xFFFFFFFFFFFFFFFD`, `0xFFFFFFFFFFFFFFFE`, and `0xFFFFFFFFFFFFFFFF` are reserved for internal using. | ||
|
|
||
| - The keys of `0xFFFFFFFFFFFFFFFD`, `0xFFFFFFFFFFFFFFFE`, and `0xFFFFFFFFFFFFFFFF` are reserved for internal using. | ||
| - Call set options.reserved_key_index to change the reserved keys if the default one conflicted with your keys. |
There was a problem hiding this comment.
Maybe 'reserved_key_start_index' would be better?
jq
force-pushed
the
r3
branch
6 times, most recently
from
9367c46
to
7733b72
Compare
|
/blossom-ci |
rhdong
reviewed
May 17, 2024
| @@ -865,6 +876,7 @@ class HashTable : public HashTableBase<K, V, S> { | |||
| return; | |||
| } | |||
| options_ = options; | |||
| CUDA_CHECK(init_reserved_keys(options.reserved_key_start_bit)); | |||
There was a problem hiding this comment.
Pls use MERLIN_CHECK to check if the value falls in [0, 61]
rhdong
reviewed
May 17, 2024
include/merlin_hashtable.cuh
Outdated
| * reserved_key_start_bit = 0, is the default behavior, HKV reserves `0xFFFFFFFFFFFFFFFD`, | ||
| * `0xFFFFFFFFFFFFFFFE`, and `0xFFFFFFFFFFFFFFFF` for internal using. | ||
| * if the default one conflicted with your keys, change the reserved_key_start_bit | ||
| * value to a numbers between 1 and 61, reserved_key_start_bit = 1 means using the |
There was a problem hiding this comment.
62 meaning using 62 and 63, then it is no longer neg number, what do you think?
jq
force-pushed
the
r3
branch
3 times, most recently
from
d9b9047
to
7e7d0af
Compare
|
/blossom-ci |
11 tasks
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Allow changing the reserved keys in the confliction, call
init_reserved_keysto change it