-
Notifications
You must be signed in to change notification settings - Fork 195
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.
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
Adds benchmarks for cub::DeviceMerge
#3529
Conversation
16568b1
to
6bb30c6
Compare
🟩 CI finished in 1h 05m: Pass: 100%/90 | Total: 15h 30m | Avg: 10m 20s | Max: 51m 03s | Hits: 414%/12772
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
Thrust | |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 90)
# | Runner |
---|---|
65 | linux-amd64-cpu16 |
11 | linux-amd64-gpu-v100-latest-1 |
9 | windows-amd64-cpu16 |
4 | linux-arm64-cpu16 |
1 | linux-amd64-gpu-h100-latest-1-testing |
🟨 CI finished in 3h 02m: Pass: 98%/90 | Total: 21h 02m | Avg: 14m 01s | Max: 1h 00m | Hits: 422%/10928
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
Thrust | |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 90)
# | Runner |
---|---|
65 | linux-amd64-cpu16 |
11 | linux-amd64-gpu-v100-latest-1 |
9 | windows-amd64-cpu16 |
4 | linux-arm64-cpu16 |
1 | linux-amd64-gpu-h100-latest-1-testing |
🟨 CI finished in 1d 01h: Pass: 98%/89 | Total: 15h 03m | Avg: 10m 09s | Max: 42m 25s | Hits: 422%/10928
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
Thrust | |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
# | Runner |
---|---|
65 | linux-amd64-cpu16 |
11 | linux-amd64-gpu-v100-latest-1 |
8 | windows-amd64-cpu16 |
4 | linux-arm64-cpu16 |
1 | linux-amd64-gpu-h100-latest-1-testing |
🟩 CI finished in 4h 15m: Pass: 100%/89 | Total: 1d 16h | Avg: 27m 07s | Max: 1h 15m | Hits: 236%/10936
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
Thrust | |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
# | Runner |
---|---|
65 | linux-amd64-cpu16 |
11 | linux-amd64-gpu-v100-latest-1 |
8 | windows-amd64-cpu16 |
4 | linux-arm64-cpu16 |
1 | linux-amd64-gpu-h100-latest-1 |
🟩 CI finished in 1h 57m: Pass: 100%/89 | Total: 16h 36m | Avg: 11m 11s | Max: 1h 48m | Hits: 422%/10936
|
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
Thrust | |
CUDA Experimental | |
python | |
CCCL C Parallel Library | |
Catch2Helper |
Modifications in project or dependencies?
Project | |
---|---|
CCCL Infrastructure | |
libcu++ | |
+/- | CUB |
+/- | Thrust |
CUDA Experimental | |
+/- | python |
+/- | CCCL C Parallel Library |
+/- | Catch2Helper |
🏃 Runner counts (total jobs: 89)
# | Runner |
---|---|
65 | linux-amd64-cpu16 |
8 | windows-amd64-cpu16 |
6 | linux-amd64-gpu-rtxa6000-latest-1 |
4 | linux-arm64-cpu16 |
3 | linux-amd64-gpu-rtx4090-latest-1 |
2 | linux-amd64-gpu-rtx2080-latest-1 |
1 | linux-amd64-gpu-h100-latest-1 |
const auto num_items_lhs = elements / 2; | ||
const auto num_items_rhs = elements - num_items_lhs; | ||
auto counting_it = thrust::make_counting_iterator(offset_t{0}); | ||
thrust::copy_if( | ||
counting_it, | ||
counting_it + elements, | ||
rnd_selector_val.begin(), | ||
thrust::make_tabulate_output_iterator(write_pivot_point_t<offset_t>{ | ||
static_cast<offset_t>(num_items_lhs), thrust::raw_pointer_cast(pivot_point.data())}), | ||
select_lhs_op); | ||
|
||
thrust::device_vector<key_t> keys_lhs(num_items_lhs); | ||
thrust::device_vector<key_t> keys_rhs(num_items_rhs); | ||
thrust::device_vector<key_t> keys_out(elements); | ||
|
||
// Generate increasing input range to sample from | ||
thrust::device_vector<key_t> increasing_input = generate(elements); | ||
thrust::sort(increasing_input.begin(), increasing_input.end()); | ||
|
||
// Select lhs from input up to pivot point | ||
offset_t pivot_point_val = pivot_point[0]; | ||
auto const end_lhs = thrust::copy_if( | ||
increasing_input.cbegin(), | ||
increasing_input.cbegin() + pivot_point_val, | ||
rnd_selector_val.cbegin(), | ||
keys_lhs.begin(), | ||
select_lhs_op); | ||
|
||
// Select rhs items from input up to pivot point | ||
auto const end_rhs = thrust::copy_if( | ||
increasing_input.cbegin(), | ||
increasing_input.cbegin() + pivot_point_val, | ||
rnd_selector_val.cbegin(), | ||
keys_rhs.begin(), | ||
select_rhs_op); | ||
// From pivot point copy all remaining items to rhs | ||
thrust::copy(increasing_input.cbegin() + pivot_point_val, increasing_input.cbegin() + elements, end_rhs); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
remark: this is beautiful 💚
// Retrieve axis parameters | ||
const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}")); | ||
const bit_entropy entropy = str_to_entropy(state.get_string("Entropy")); | ||
|
||
// We generate data distributions in the range [0, 255] that, with lower entropy, get skewed towards 0. | ||
// We use this to generate increasingly large *consecutive* segments of data that are getting selected from the lhs | ||
thrust::device_vector<uint8_t> rnd_selector_val = generate(elements, entropy); | ||
uint8_t threshold = 128; | ||
select_if_less_than_t select_lhs_op{false, threshold}; | ||
select_if_less_than_t select_rhs_op{true, threshold}; | ||
|
||
// The following algorithm only works under the precondition that there's at least 50% of the data in the lhs | ||
// If that's not the case, we simply swap the logic for selecting into lhs and rhs | ||
const auto num_items_selected_into_lhs = | ||
static_cast<offset_t>(thrust::count_if(rnd_selector_val.begin(), rnd_selector_val.end(), select_lhs_op)); | ||
if (num_items_selected_into_lhs < elements / 2) | ||
{ | ||
using ::cuda::std::swap; | ||
swap(select_lhs_op, select_rhs_op); | ||
} | ||
|
||
// We want lhs and rhs to be of equal size. We also want to have skewed distributions, such that we put different | ||
// workloads on the binary search part. For this reason, we identify the index from the input, referred to as pivot | ||
// point, after which the lhs is "full". We compose the rhs by selecting all items up to the pivot point that were not | ||
// selected for lhs and *all* items after the pivot point. | ||
constexpr std::size_t num_pivot_points = 1; | ||
thrust::device_vector<offset_t> pivot_point(num_pivot_points); | ||
const auto num_items_lhs = elements / 2; | ||
const auto num_items_rhs = elements - num_items_lhs; | ||
auto counting_it = thrust::make_counting_iterator(offset_t{0}); | ||
thrust::copy_if( | ||
counting_it, | ||
counting_it + elements, | ||
rnd_selector_val.begin(), | ||
thrust::make_tabulate_output_iterator(write_pivot_point_t<offset_t>{ | ||
static_cast<offset_t>(num_items_lhs), thrust::raw_pointer_cast(pivot_point.data())}), | ||
select_lhs_op); | ||
|
||
thrust::device_vector<key_t> keys_lhs(num_items_lhs); | ||
thrust::device_vector<key_t> keys_rhs(num_items_rhs); | ||
thrust::device_vector<key_t> keys_out(elements); | ||
thrust::device_vector<value_t> values_lhs(num_items_lhs); | ||
thrust::device_vector<value_t> values_rhs(num_items_rhs); | ||
thrust::device_vector<value_t> values_out(elements); | ||
|
||
// Generate increasing input range to sample from | ||
thrust::device_vector<key_t> increasing_input = generate(elements); | ||
thrust::sort(increasing_input.begin(), increasing_input.end()); | ||
|
||
// Select lhs from input up to pivot point | ||
offset_t pivot_point_val = pivot_point[0]; | ||
auto const end_lhs = thrust::copy_if( | ||
increasing_input.cbegin(), | ||
increasing_input.cbegin() + pivot_point_val, | ||
rnd_selector_val.cbegin(), | ||
keys_lhs.begin(), | ||
select_lhs_op); | ||
|
||
// Select rhs items from input up to pivot point | ||
auto const end_rhs = thrust::copy_if( | ||
increasing_input.cbegin(), | ||
increasing_input.cbegin() + pivot_point_val, | ||
rnd_selector_val.cbegin(), | ||
keys_rhs.begin(), | ||
select_rhs_op); | ||
// From pivot point copy all remaining items to rhs | ||
thrust::copy(increasing_input.cbegin() + pivot_point_val, increasing_input.cbegin() + elements, end_rhs); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
suggestion: this looks like a duplicate of workload generation you added in the keys version. Maybe we could simplify the code a bit by having cuda::std::pair<thrust::device_vector<T>, thrust::device_vector<T>> generate(elements, entropy); auto [keys_lhs, keys_rhs] = ...
function.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Sounds good! I'll do that in a follow-up PR 👍
Description
Closes #3528