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

Improve RBC eps-neighborhood query performance #2211

Merged
merged 13 commits into from
Mar 11, 2024
Merged

Improve RBC eps-neighborhood query performance #2211

Changes from 1 commit
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
2179ed6
optimize rbc eps-NN query
mfoerste4 Mar 4, 2024
1782558
remove L2 cache flush
mfoerste4 Mar 5, 2024
b3e998a
review suggestion constexpr
mfoerste4 Mar 5, 2024
0129748
utilize reordered index for rbc knn
mfoerste4 Mar 5, 2024
ab34c25
Merge branch 'branch-24.04' into rbc_eps_nn_performance
mfoerste4 Mar 5, 2024
fde845b
minor optimizations / suggestions
mfoerste4 Mar 7, 2024
663dabf
add restrict, free 1 register
mfoerste4 Mar 7, 2024
f861403
Merge branch 'rbc_eps_nn_performance' of github.com:mfoerste4/raft in…
mfoerste4 Mar 7, 2024
f417b1a
Merge branch 'branch-24.04' into rbc_eps_nn_performance
mfoerste4 Mar 7, 2024
2be4c17
apply modifications to all kernel variants
mfoerste4 Mar 7, 2024
768eb48
remove one more FLO per landmark visited
mfoerste4 Mar 7, 2024
2315e2f
change mask update
mfoerste4 Mar 8, 2024
2d06d0a
remove warp divergence
mfoerste4 Mar 8, 2024
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
Prev Previous commit
Next Next commit
add restrict, free 1 register
  • Loading branch information
@mfoerste4
mfoerste4 committed Mar 7, 2024
commit 663dabf6bf6f0fd66a8a7d51415a5dac788be302
57 changes: 31 additions & 26 deletions cpp/include/raft/spatial/knn/detail/ball_cover/registers-inl.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -719,21 +719,22 @@ template <typename value_idx = std::int64_t,
int dim = 3,
typename value_int = std::uint32_t,
typename distance_func>
RAFT_KERNEL __launch_bounds__(tpb) block_rbc_kernel_eps_csr_pass_xd(const value_t* X_reordered,
const value_t* X,
const value_int n_queries,
const value_int n_cols,
const value_t* R,
const value_int m,
const value_t eps,
const value_int n_landmarks,
const value_idx* R_indptr,
const value_idx* R_1nn_cols,
const value_t* R_1nn_dists,
const value_t* R_radius,
distance_func dfunc,
value_idx* adj_ia,
value_idx* adj_ja)
RAFT_KERNEL __launch_bounds__(tpb)
block_rbc_kernel_eps_csr_pass_xd(const value_t* __restrict__ X_reordered,
const value_t* __restrict__ X,
const value_int n_queries,
const value_int n_cols,
const value_t* __restrict__ R,
const value_int m,
const value_t eps,
const value_int n_landmarks,
const value_idx* __restrict__ R_indptr,
const value_idx* __restrict__ R_1nn_cols,
const value_t* __restrict__ R_1nn_dists,
const value_t* __restrict__ R_radius,
distance_func dfunc,
value_idx* __restrict__ adj_ia,
value_idx* adj_ja)
{
constexpr int num_warps = tpb / WarpSize;
constexpr int max_lid = WarpSize - 1;
Expand All @@ -748,10 +749,14 @@ RAFT_KERNEL __launch_bounds__(tpb) block_rbc_kernel_eps_csr_pass_xd(const value_
// this is an early out for a full warp
if (query_id >= n_queries) return;

value_idx column_index_offset = write_pass ? adj_ia[query_id] : 0;
uint32_t column_index_offset = 0;

// we have no neighbors to fill for this query
if (write_pass && adj_ia[query_id + 1] == column_index_offset) return;
if constexpr (write_pass) {
value_idx offset = adj_ia[query_id];
// we have no neighbors to fill for this query
if (offset == adj_ia[query_id + 1]) return;
adj_ja += offset;
}

const value_t* x_ptr = X + (dim * query_id);
value_t local_x_ptr[dim];
Expand Down Expand Up @@ -812,11 +817,11 @@ RAFT_KERNEL __launch_bounds__(tpb) block_rbc_kernel_eps_csr_pass_xd(const value_
if constexpr (write_pass) {
const int mask = raft::ballot(in_range);
if (in_range) {
const uint32_t index = R_1nn_cols[R_start_offset + i];
const value_idx row_pos = column_index_offset + __popc(mask & lid_mask);
adj_ja[row_pos] = index;
const uint32_t index = R_1nn_cols[R_start_offset + i];
const uint32_t row_pos = __popc(mask & lid_mask);
adj_ja[row_pos] = index;
}
column_index_offset += __popc(mask);
adj_ja += __popc(mask);
} else {
column_index_offset += (in_range);
}
Expand All @@ -833,11 +838,11 @@ RAFT_KERNEL __launch_bounds__(tpb) block_rbc_kernel_eps_csr_pass_xd(const value_
if constexpr (write_pass) {
const int mask = raft::ballot(in_range);
if (in_range) {
const uint32_t index = R_1nn_cols[R_start_offset + i];
const value_idx row_pos = column_index_offset + __popc(mask & lid_mask);
adj_ja[row_pos] = index;
const uint32_t index = R_1nn_cols[R_start_offset + i];
const uint32_t row_pos = __popc(mask & lid_mask);
adj_ja[row_pos] = index;
}
column_index_offset += __popc(mask);
adj_ja += __popc(mask);
} else {
column_index_offset += (in_range);
}
Expand Down