-
Notifications
You must be signed in to change notification settings - Fork 75
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
* Make use of HIP vector types in rocprim * Format device_merge_sort_mergepath.hpp to simplify future changes The files has indents not compatible with the current formatting style, small changes in the code lead to incorrect auto-formatting. * Support large indices in device_merge_sort, add tests * Use BlockSize instead of rocprim::flat_block_size() in device_merge_sort BlockSize is known at compile-time, while flat_block_size() computes a flat index from 3 dimensions. * Fix a rare bug in reduce_by_key with large indices For some sizes d_global_head_count may be updated twice in reduce_by_key_init_kernel: by the 0-th thread and by the save_index-th thread. This is possible when for grid_size * block_size > number_of_tiles - 1, i.e. number_of_tiles is sligthtly smaller than number_of_tiles_launch so the last block has a thread with id = number_of_tiles - 1. * Only use LDS bank conflict resolution in block_exchange when size per thread is higher than 128 bits Changed bank_conflict_padding for smaller types * test(test_*_iterator): fix warnings from the deprecation of std::ostream returning << operator for iterators * ci: set up sccache * ci(.gitlab-ci.yml): allow passing benchmark seed as variable * allow const inputs for rocprim::partial_sort_copy * Add new generate_limit tool to fix numeric_limits not working on custom type * Changes to fix clang-format * Added documentation for (non-)stable sorts * Fixed block_radix_sort test to check stablity * updated 3 examples of block_sort by specifying items per thread * corrected grammar in examples * Apply 1 suggestion(s) to 1 file(s) Co-authored-by: Nick Breed <[email protected]> * kept one example with the default behaviour * key-value example, corrected key type * fix: fixed incorrect definition and use of rocprim specific type traits for 128-bit integers This fixes an issue where in certain situations where using radix codec on 128-bit integers would not compile due to ambiguity. * Remove hcc path from masked_bit_count * Fix data generation in tests: tailing items may be uninitialized * find_first_of: Add a naive implementation * find_first_of: Add a specialized kernel with early-exit * find_first_of: Extend tests, add a test for large indices * find_first_of: Use ordered_block_id for better load balancing between CUs * find_first_of: Implement autotuning * find_first_of: Add docs * Add test_ordered_block_id * Update copyright date and code format * Add a test for 10000 blocks and a timer * Added find_first_of to rocprim header * Fix build error after adding headers * fix: fix doxygen warning due to __launch_bounds__ --------- Co-authored-by: Nick Breed <[email protected]> Co-authored-by: Anton Gorenko <[email protected]> Co-authored-by: Nara Prasetya <[email protected]> Co-authored-by: Beatriz Navidad Vilches <[email protected]> Co-authored-by: milo <[email protected]> Co-authored-by: Cenxuan Tian <[email protected]>
- Loading branch information
7 people
authored
Oct 25, 2024
1 parent
0ce43c0
commit 81f6073
Showing
55 changed files
with
3,225 additions
and
885 deletions.
There are no files selected for viewing
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
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
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,150 @@ | ||
| // MIT License | ||
| // | ||
| // Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. | ||
| // | ||
| // Permission is hereby granted, free of charge, to any person obtaining a copy | ||
| // of this software and associated documentation files (the "Software"), to deal | ||
| // in the Software without restriction, including without limitation the rights | ||
| // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell | ||
| // copies of the Software, and to permit persons to whom the Software is | ||
| // furnished to do so, subject to the following conditions: | ||
| // | ||
| // The above copyright notice and this permission notice shall be included in all | ||
| // copies or substantial portions of the Software. | ||
| // | ||
| // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | ||
| // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | ||
| // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | ||
| // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | ||
| // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | ||
| // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE | ||
| // SOFTWARE. | ||
|
|
||
| #include "benchmark_device_find_first_of.parallel.hpp" | ||
| #include "benchmark_utils.hpp" | ||
|
|
||
| // CmdParser | ||
| #include "cmdparser.hpp" | ||
|
|
||
| // Google Benchmark | ||
| #include <benchmark/benchmark.h> | ||
|
|
||
| // HIP API | ||
| #include <hip/hip_runtime.h> | ||
|
|
||
| #include <cstddef> | ||
| #include <string> | ||
|
|
||
| #ifndef DEFAULT_BYTES | ||
| constexpr size_t DEFAULT_BYTES = size_t{1} << 27; // 128 MiB | ||
| #endif | ||
|
|
||
| #define CREATE_BENCHMARK_FIND_FIRST_OF(TYPE, KEYS_SIZE, FIRST_OCCURENCE) \ | ||
| { \ | ||
| const device_find_first_of_benchmark<TYPE> instance(KEYS_SIZE, FIRST_OCCURENCE); \ | ||
| REGISTER_BENCHMARK(benchmarks, size, seed, stream, instance); \ | ||
| } | ||
|
|
||
| // clang-format off | ||
| #define CREATE_BENCHMARK0(TYPE, KEYS_SIZE) \ | ||
| { \ | ||
| CREATE_BENCHMARK_FIND_FIRST_OF(TYPE, KEYS_SIZE, 0.1) \ | ||
| CREATE_BENCHMARK_FIND_FIRST_OF(TYPE, KEYS_SIZE, 0.5) \ | ||
| CREATE_BENCHMARK_FIND_FIRST_OF(TYPE, KEYS_SIZE, 1.0) \ | ||
| } | ||
|
|
||
| #define CREATE_BENCHMARK(TYPE) \ | ||
| { \ | ||
| CREATE_BENCHMARK0(TYPE, 1) \ | ||
| CREATE_BENCHMARK0(TYPE, 10) \ | ||
| CREATE_BENCHMARK0(TYPE, 100) \ | ||
| CREATE_BENCHMARK0(TYPE, 1000) \ | ||
| CREATE_BENCHMARK0(TYPE, 10000) \ | ||
| } | ||
| // clang-format on | ||
|
|
||
| int main(int argc, char* argv[]) | ||
| { | ||
| cli::Parser parser(argc, argv); | ||
| parser.set_optional<size_t>("size", "size", DEFAULT_BYTES, "number of bytes"); | ||
| parser.set_optional<int>("trials", "trials", -1, "number of iterations"); | ||
| parser.set_optional<std::string>("name_format", | ||
| "name_format", | ||
| "human", | ||
| "either: json,human,txt"); | ||
| parser.set_optional<std::string>("seed", "seed", "random", get_seed_message()); | ||
| #ifdef BENCHMARK_CONFIG_TUNING | ||
| // optionally run an evenly split subset of benchmarks, when making multiple program invocations | ||
| parser.set_optional<int>("parallel_instance", | ||
| "parallel_instance", | ||
| 0, | ||
| "parallel instance index"); | ||
| parser.set_optional<int>("parallel_instances", | ||
| "parallel_instances", | ||
| 1, | ||
| "total parallel instances"); | ||
| #endif | ||
| parser.run_and_exit_if_error(); | ||
|
|
||
| // Parse argv | ||
| benchmark::Initialize(&argc, argv); | ||
| const size_t size = parser.get<size_t>("size"); | ||
| const int trials = parser.get<int>("trials"); | ||
| bench_naming::set_format(parser.get<std::string>("name_format")); | ||
| const std::string seed_type = parser.get<std::string>("seed"); | ||
| const managed_seed seed(seed_type); | ||
|
|
||
| // HIP | ||
| hipStream_t stream = 0; // default | ||
|
|
||
| // Benchmark info | ||
| add_common_benchmark_info(); | ||
| benchmark::AddCustomContext("size", std::to_string(size)); | ||
| benchmark::AddCustomContext("seed", seed_type); | ||
|
|
||
| // Add benchmarks | ||
| std::vector<benchmark::internal::Benchmark*> benchmarks{}; | ||
| #ifdef BENCHMARK_CONFIG_TUNING | ||
| const int parallel_instance = parser.get<int>("parallel_instance"); | ||
| const int parallel_instances = parser.get<int>("parallel_instances"); | ||
| config_autotune_register::register_benchmark_subset(benchmarks, | ||
| parallel_instance, | ||
| parallel_instances, | ||
| size, | ||
| seed, | ||
| stream); | ||
| #else // BENCHMARK_CONFIG_TUNING | ||
| CREATE_BENCHMARK(int8_t) | ||
| CREATE_BENCHMARK(int16_t) | ||
| CREATE_BENCHMARK(int32_t) | ||
| CREATE_BENCHMARK(float) | ||
| CREATE_BENCHMARK(int64_t) | ||
| CREATE_BENCHMARK(double) | ||
|
|
||
| using custom_int2 = custom_type<int, int>; | ||
| using custom_longlong_double = custom_type<long long, double>; | ||
|
|
||
| CREATE_BENCHMARK(custom_int2) | ||
| CREATE_BENCHMARK(custom_longlong_double) | ||
| #endif // BENCHMARK_CONFIG_TUNING | ||
|
|
||
| // Use manual timing | ||
| for(auto& b : benchmarks) | ||
| { | ||
| b->UseManualTime(); | ||
| b->Unit(benchmark::kMillisecond); | ||
| } | ||
|
|
||
| // Force number of iterations | ||
| if(trials > 0) | ||
| { | ||
| for(auto& b : benchmarks) | ||
| { | ||
| b->Iterations(trials); | ||
| } | ||
| } | ||
|
|
||
| // Run benchmarks | ||
| benchmark::RunSpecifiedBenchmarks(); | ||
| return 0; | ||
| } |
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,30 @@ | ||
| // MIT License | ||
| // | ||
| // Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. | ||
| // | ||
| // Permission is hereby granted, free of charge, to any person obtaining a copy | ||
| // of this software and associated documentation files (the "Software"), to deal | ||
| // in the Software without restriction, including without limitation the rights | ||
| // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell | ||
| // copies of the Software, and to permit persons to whom the Software is | ||
| // furnished to do so, subject to the following conditions: | ||
| // | ||
| // The above copyright notice and this permission notice shall be included in all | ||
| // copies or substantial portions of the Software. | ||
| // | ||
| // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | ||
| // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | ||
| // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | ||
| // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | ||
| // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | ||
| // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE | ||
| // SOFTWARE. | ||
|
|
||
| #include "benchmark_device_find_first_of.parallel.hpp" | ||
| #include "benchmark_utils.hpp" | ||
|
|
||
| namespace | ||
| { | ||
| auto benchmarks = config_autotune_register::create_bulk( | ||
| device_find_first_of_benchmark_generator<@DataType@, @BlockSize@>::create); | ||
| } // namespace |
Oops, something went wrong.