Skip to content

MI300 radix sort, run length encode, and reduce by key optimizations #674

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

Merged
stanleytsang-amd merged 9 commits into from
Feb 3, 2025
Merged

MI300 radix sort, run length encode, and reduce by key optimizations #674

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
689c730
Resolve "Add hipGraph compatibility for device run length encode"
Beanavil Nov 21, 2024
d30c3ee
Resolve "Investigate RLE (non trivial runs) failure"
Beanavil Nov 21, 2024
a6af335
Resolve "Config tuning and dynamic dispatch for device run length enc…
yungshengtu Nov 21, 2024
dda045a
Resolve "Implement type traits interface for rocPRIM's compiler-exten…
cenxuantian Nov 21, 2024
ddf8d1c
remove ordered block id from run length encode
Snektron Nov 21, 2024
2d5707b
MI200 & MI300 tuning
Snektron Dec 17, 2024
fa09122
Optimizations for MI300 unique performance
Snektron Jan 16, 2025
81bf9cf
improve type traits interface docs
Snektron Jan 20, 2025
594012d
Move changlog items to 6.4.0
NB4444 Feb 3, 2025
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
4 changes: 4 additions & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,10 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
* Added the parallel `search` and `find_end` device functions similar to `std::search` and `std::find_end`, these functions search for the first and last occurrence of the sequence respectively.
* Added a parallel device-level function, `rocprim::search_n`, similar to the C++ Standard Library `std::search_n` algorithm.
* Added new constructors and a `base` function, and added `constexpr` specifier to all functions in `rocprim::reverse_iterator` to improve parity with the C++17 `std::reverse_iterator`.
* Added hipGraph support to device run-length-encode for non trivial runs (`rocprim::run_length_encode_non_trivial_runs`).
* Added configuration autotuning to device run-length-encode for non trivial runs (`rocprim::run_length_encode_non_trivial_runs`) for improved performance on selected architectures.
* Added configuration autotuning to device run-length-encode for trivial runs (`rocprim::run_length_encode`) for improved performance on selected architectures.
* Added a new type traits interface to enable users to provide additional type trait information to rocPRIM, facilitating better compatibility with custom types.

### Changed

Expand Down
1 change: 1 addition & 0 deletions benchmark/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,6 +161,7 @@ add_rocprim_benchmark(benchmark_device_reduce_by_key.cpp)
add_rocprim_benchmark(benchmark_device_reduce_by_key_deterministic.cpp)
add_rocprim_benchmark(benchmark_device_reduce.cpp)
add_rocprim_benchmark(benchmark_device_run_length_encode.cpp)
add_rocprim_benchmark(benchmark_device_run_length_encode_non_trivial_runs.cpp)
add_rocprim_benchmark(benchmark_device_scan.cpp)
add_rocprim_benchmark(benchmark_device_scan_deterministic.cpp)
add_rocprim_benchmark(benchmark_device_scan_by_key.cpp)
Expand Down
14 changes: 11 additions & 3 deletions benchmark/ConfigAutotuneSettings.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -112,13 +112,21 @@ DataType;BlockSize;" PARENT_SCOPE)
set(list_across "${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES};128 192 256 384 512" PARENT_SCOPE)
set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@" PARENT_SCOPE)
elseif(file STREQUAL "benchmark_device_reduce_by_key")
set(list_across_names "KeyType;ValueType;BlockSize;TilesPerBlock" PARENT_SCOPE)
set(list_across "${LIMITED_TUNING_TYPES};${TUNING_TYPES};128 192 256 384 512;1 2" PARENT_SCOPE)
set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@_@TilesPerBlock@" PARENT_SCOPE)
set(list_across_names "KeyType;ValueType;BlockSize" PARENT_SCOPE)
set(list_across "${LIMITED_TUNING_TYPES};${TUNING_TYPES};128 192 256 384 512" PARENT_SCOPE)
set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@" PARENT_SCOPE)
elseif(file STREQUAL "benchmark_device_find_first_of")
set(list_across_names "DataType;BlockSize" PARENT_SCOPE)
set(list_across "${LIMITED_TUNING_TYPES};32 64 128 256 512 1024" PARENT_SCOPE)
set(output_pattern_suffix "@DataType@_@BlockSize@" PARENT_SCOPE)
elseif(file STREQUAL "benchmark_device_run_length_encode")
set(list_across_names "KeyType;BlockSize" PARENT_SCOPE)
set(list_across "${TUNING_TYPES};128 192 256 384 512" PARENT_SCOPE)
set(output_pattern_suffix "@KeyType@_@BlockSize@" PARENT_SCOPE)
elseif(file STREQUAL "benchmark_device_run_length_encode_non_trivial_runs")
set(list_across_names "KeyType;BlockSize;BlockLoadMethod" PARENT_SCOPE)
set(list_across "${TUNING_TYPES};64 128 256 512 1024;block_load_vectorize block_load_warp_transpose" PARENT_SCOPE)
set(output_pattern_suffix "@KeyType@_@BlockSize@_@BlockLoadMethod@" PARENT_SCOPE)
elseif(file STREQUAL "benchmark_device_merge")
set(list_across_names "KeyType;ValueType;BlockSize" PARENT_SCOPE)
set(list_across "${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES};32 64 128 256 512 1024" PARENT_SCOPE)
Expand Down
4 changes: 2 additions & 2 deletions benchmark/benchmark_device_reduce_by_key.parallel.cpp.in
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// MIT License
//
// Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2024-2025 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
Expand All @@ -27,5 +27,5 @@

namespace {
auto benchmarks = config_autotune_register::create_bulk(
device_reduce_by_key_benchmark_generator<@KeyType@, @ValueType@, @BlockSize@, @TilesPerBlock@>::create);
device_reduce_by_key_benchmark_generator<@KeyType@, @ValueType@, @BlockSize@>::create);
}
10 changes: 4 additions & 6 deletions benchmark/benchmark_device_reduce_by_key.parallel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// MIT License
//
// Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2024-2025 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
Expand Down Expand Up @@ -44,8 +44,7 @@ std::string config_name()
{
const rocprim::detail::reduce_by_key_config_params params = Config();
return "{bs:" + std::to_string(params.kernel_config.block_size)
+ ",ipt:" + std::to_string(params.kernel_config.items_per_thread)
+ ",tpb:" + std::to_string(params.tiles_per_block) + "}";
+ ",ipt:" + std::to_string(params.kernel_config.items_per_thread) + "}";
}

template<>
Expand Down Expand Up @@ -224,7 +223,7 @@ struct device_reduce_by_key_benchmark : public config_autotune_interface

#ifdef BENCHMARK_CONFIG_TUNING

template<typename KeyType, typename ValueType, int BlockSize, int TilesPerBlock>
template<typename KeyType, typename ValueType, unsigned int BlockSize>
struct device_reduce_by_key_benchmark_generator
{
template<int ItemsPerThread>
Expand All @@ -237,8 +236,7 @@ struct device_reduce_by_key_benchmark_generator
ItemsPerThread,
rocprim::block_load_method::block_load_transpose,
rocprim::block_load_method::block_load_transpose,
rocprim::block_scan_algorithm::using_warp_scan,
TilesPerBlock>;
rocprim::block_scan_algorithm::using_warp_scan>;
// max segment length argument is irrelevant, tuning overrides segment length
storage.emplace_back(
std::make_unique<
Expand Down
Loading