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.

Sign up for GitHub

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

[Feat] find and find_or_insert with return values' addresses #122

Merged
merged 2 commits into from
May 11, 2023
Merged

[Feat] find and find_or_insert with return values' addresses #122

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6aba056
[Feat] find with return values' addresses
rhdong May 10, 2023
03be26a
[Feat] find_or_insert with return values' address
rhdong May 10, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,3 +121,8 @@ add_executable(group_lock_test tests/group_lock_test.cc)
target_compile_features(group_lock_test PUBLIC cxx_std_14)
set_target_properties(group_lock_test PROPERTIES CUDA_ARCHITECTURES OFF)
TARGET_LINK_LIBRARIES(group_lock_test gtest_main)

add_executable(find_or_insert_ptr_test tests/find_or_insert_ptr_test.cc.cu)
target_compile_features(find_or_insert_ptr_test PUBLIC cxx_std_14)
set_target_properties(find_or_insert_ptr_test PROPERTIES CUDA_ARCHITECTURES OFF)
TARGET_LINK_LIBRARIES(find_or_insert_ptr_test gtest_main)
30 changes: 15 additions & 15 deletions README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,29 +103,29 @@ Your environment must meet the following requirements:

* dim = 4, capacity = 64 Million-KV, HBM = 32 GB, HMEM = 0 GB

| load_factor | insert_or_assign | find | find_or_insert | assign | insert_and_evict |
|------------:|:----------------:|---------:|:--------------:|-------:|:----------------:|
| 0.50 | 1.673 | 3.101 | 1.364 | 2.497 | 0.820 |
| 0.75 | 0.909 | 1.907 | 0.828 | 1.536 | 0.620 |
| 1.00 | 0.281 | 0.529 | 0.238 | 0.215 | 0.247 |
| load_factor | insert_or_assign | find | find_or_insert | assign | find* | insert_and_evict |
|------------:|:----------------:|-------:|:--------------:|-------:|-------:|:----------------:|
| 0.50 | 1.686 | 3.191 | 1.379 | 2.562 | 4.005 | 0.790 |
| 0.75 | 0.870 | 1.932 | 0.821 | 1.530 | 2.243 | 0.628 |
| 1.00 | 0.282 | 0.531 | 0.237 | 0.215 | 0.561 | 0.245 |

* dim = 64, capacity = 64 Million-KV, HBM = 16 GB, HMEM = 0 GB

| load_factor | insert_or_assign | find | find_or_insert | assign | insert_and_evict |
|------------:|:----------------:|-------:|:--------------:|-------:|:----------------:|
| 0.50 | 0.801 | 1.606 | 0.833 | 1.128 | 0.590 |
| 0.75 | 0.569 | 1.073 | 0.532 | 0.821 | 0.406 |
| 1.00 | 0.261 | 0.397 | 0.194 | 0.183 | 0.140 |
| load_factor | insert_or_assign | find | find_or_insert | assign | find* | insert_and_evict |
|------------:|:----------------:|-------:|:--------------:|-------:|-------:|:----------------:|
| 0.50 | 0.844 | 1.613 | 0.835 | 1.128 | 3.981 | 0.607 |
| 0.75 | 0.568 | 1.079 | 0.530 | 0.821 | 2.252 | 0.427 |
| 1.00 | 0.261 | 0.398 | 0.194 | 0.182 | 0.561 | 0.141 |

### On HBM+HMEM hybrid mode:

* dim = 64, capacity = 128 Million-KV, HBM = 16 GB, HMEM = 16 GB

| load_factor | insert_or_assign | find | find_or_insert | assign |
|------------:|:----------------:|-------:|:--------------:|-------:|
| 0.50 | 0.121 | 0.148 | 0.122 | 0.150 |
| 0.75 | 0.114 | 0.145 | 0.117 | 0.146 |
| 1.00 | 0.082 | 0.118 | 0.079 | 0.096 |
| load_factor | insert_or_assign | find | find_or_insert | assign | find* |
|------------:|:----------------:|-------:|:--------------:|-------:|-------:|
| 0.50 | 0.123 | 0.149 | 0.123 | 0.150 | 3.688 |
| 0.75 | 0.114 | 0.145 | 0.118 | 0.147 | 1.063 |
| 1.00 | 0.085 | 0.120 | 0.081 | 0.097 | 0.716 |

* dim = 64, capacity = 1024 Million-KV, HBM = 56 GB, HMEM = 200 GB

Expand Down
45 changes: 45 additions & 0 deletions benchmark/benchmark_util.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@
#include <chrono>
#include <cmath>
#include <cstdint>
#include "merlin/utils.cuh"

namespace benchmark {
enum class TimeUnit {
Expand Down Expand Up @@ -47,4 +48,48 @@ struct Timer {
std::chrono::time_point<std::chrono::steady_clock> startRecord{};
std::chrono::time_point<std::chrono::steady_clock> endRecord{};
};
template <class V>
__global__ void read_from_ptr_kernel(const V* const* __restrict src,
V* __restrict dst, const size_t dim,
size_t N) {
size_t tid = (blockIdx.x * blockDim.x) + threadIdx.x;

for (size_t t = tid; t < N; t += blockDim.x * gridDim.x) {
int vec_index = int(t / dim);
int dim_index = t % dim;
dst[vec_index * dim + dim_index] = src[vec_index][dim_index];
}
}

template <class V>
void read_from_ptr(const V* const* __restrict src, V* __restrict dst,
const size_t dim, size_t n, cudaStream_t stream) {
const size_t block_size = 1024;
const size_t N = n * dim;
const size_t grid_size = nv::merlin::SAFE_GET_GRID_SIZE(N, block_size);

read_from_ptr_kernel<V>
<<<grid_size, block_size, 0, stream>>>(src, dst, dim, N);
}

template <class V>
__global__ void array2ptr_kernel(V** ptr, V* __restrict array, const size_t dim,
size_t N) {
size_t tid = (blockIdx.x * blockDim.x) + threadIdx.x;

for (size_t t = tid; t < N; t += blockDim.x * gridDim.x) {
int vec_index = int(t);
ptr[vec_index] = array + vec_index * dim;
}
}

template <class V>
void array2ptr(V** ptr, V* __restrict array, const size_t dim, size_t n,
cudaStream_t stream) {
const size_t block_size = 1024;
const size_t N = n;
const size_t grid_size = nv::merlin::SAFE_GET_GRID_SIZE(N, block_size);

array2ptr_kernel<V><<<grid_size, block_size, 0, stream>>>(ptr, array, dim, N);
}
} // namespace benchmark
69 changes: 59 additions & 10 deletions benchmark/merlin_hashtable_benchmark.cc.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,6 +54,8 @@ enum class API_Select {
find_or_insert = 2,
assign = 3,
insert_and_evict = 4,
find_ptr = 5,
find_or_insert_ptr = 6,
};

enum class Hit_Mode {
Expand Down Expand Up @@ -310,6 +312,41 @@ float test_one_api(const API_Select api, const size_t dim,
timer.end();
break;
}
case API_Select::find_ptr: {
V** d_vectors_ptr = nullptr;
CUDA_CHECK(cudaMalloc(&d_vectors_ptr, key_num_per_op * sizeof(V*)));
benchmark::array2ptr(d_vectors_ptr, d_vectors, options.dim,
key_num_per_op, stream);

CUDA_CHECK(cudaStreamSynchronize(stream));
timer.start();
table->find(key_num_per_op, d_keys, d_vectors_ptr, d_found, d_metas,
stream);
CUDA_CHECK(cudaStreamSynchronize(stream));
timer.end();
benchmark::read_from_ptr(d_vectors_ptr, d_vectors, options.dim,
key_num_per_op, stream);
CUDA_CHECK(cudaStreamSynchronize(stream));
CUDA_CHECK(cudaFree(d_vectors_ptr));
break;
}
case API_Select::find_or_insert_ptr: {
V** d_vectors_ptr = nullptr;
bool* d_found;
CUDA_CHECK(cudaMalloc(&d_found, key_num_per_op * sizeof(bool)));
CUDA_CHECK(cudaMalloc(&d_vectors_ptr, key_num_per_op * sizeof(V*)));
benchmark::array2ptr(d_vectors_ptr, d_vectors, options.dim,
key_num_per_op, stream);
CUDA_CHECK(cudaStreamSynchronize(stream));
timer.start();
table->find_or_insert(key_num_per_op, d_keys, d_vectors_ptr, d_found,
d_metas, stream);
CUDA_CHECK(cudaStreamSynchronize(stream));
timer.end();
CUDA_CHECK(cudaFree(d_vectors_ptr));
CUDA_CHECK(cudaFree(d_found));
break;
}
default: {
std::cout << "[Unsupport API]\n";
}
Expand Down Expand Up @@ -345,7 +382,9 @@ void print_title() {
<< "| insert_or_assign "
<< "| find "
<< "| find_or_insert "
<< "| assign ";
<< "| assign "
<< "| find* "
<< "| find_or_insert* ";
if (Test_Mode::pure_hbm == test_mode) {
cout << "| insert_and_evict ";
}
Expand All @@ -354,16 +393,20 @@ void print_title() {
//<< "| load_factor "
cout << "|------------:"
//<< "| insert_or_assign "
<< "|:----------------:"
<< "|-----------------:"
//<< "| find "
<< "|-------:"
//<< "| find_or_insert "
<< "|:--------------:"
<< "|---------------:"
//<< "| assign "
<< "|-------:";
<< "|-------:"
//<< "| find* "
<< "|-------:"
//<< "| find_or_insert* "
<< "|----------------:";
if (Test_Mode::pure_hbm == test_mode) {
//<< "| insert_and_evict "
cout << "|:----------------:";
cout << "|-----------------:";
}
cout << "|\n";
}
Expand All @@ -374,11 +417,9 @@ void test_main(const size_t dim,
const size_t hbm4values = 16, const float load_factor = 1.0f) {
std::cout << "|" << rep(8) << fixed << setprecision(2) << load_factor << " ";
std::vector<API_Select> apis{
API_Select::insert_or_assign,
API_Select::find,
API_Select::find_or_insert,
API_Select::assign,
};
API_Select::insert_or_assign, API_Select::find,
API_Select::find_or_insert, API_Select::assign,
API_Select::find_ptr, API_Select::find_or_insert_ptr};
if (Test_Mode::pure_hbm == test_mode) {
apis.push_back(API_Select::insert_and_evict);
}
Expand Down Expand Up @@ -412,6 +453,14 @@ void test_main(const size_t dim,
std::cout << rep(12);
break;
}
case API_Select::find_ptr: {
std::cout << rep(2);
break;
}
case API_Select::find_or_insert_ptr: {
std::cout << rep(11);
break;
}
default: {
std::cout << "[Unsupport API]";
}
Expand Down
Loading