Repository navigation
perf(rocprim) Reduce number of kernels generated for several benchmarks and tests - #13092
Conversation
This change introduces two functors for build time optimization of several benchmarks and tests that use lambda functions in a particular context. Here's an example of the optimization for benchmark_device_select: Several benchmark_device_select benchmarks require a predicate function. Defining that prodice function via a lambda function (within the benchmark's run() implementation) creates this problem: - The benchmark struct takes template args. So if we define a predicate function within run() using a lambda, then for each unique instantiation of the benchmark, the lambda predicate function will be given a unique type. - The predicate passed into rocprim::select, and its type is passed as a template argument to partition_impl. - This means that each unique instantiation of the benchmark will ultimately cause a unique kernel to be generated. - Since the predicate-based benchmarks are instantiated with different probabilities, there is a unique kernel generated for each unique probability value, even though the actual kernel code is identical. Defining the predicates as functors (that take the theshold/flag type as a template parameter) means that the same type is passed into each benchmark instance's call to rocprim::select. This can significantly reduce the number of kernels that need to be generated. This change implements the above optimization for the following targets: - benchmark_device_select - benchmark_device_segmented_radix_sort - benchmark_device_partition - test_device_segmented_topk
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Pre-commit check failed⛔ pre-commit failed Please run locally:
This repo uses |
There seem to be some performance regressions introduced by these build-time optimizations, so revert them.
|
🎉 All checks passed! This PR is ready for review. |
amd-jmahovsky
left a comment
amd-jmahovsky
left a comment
There was a problem hiding this comment.
LGTM but fix the typos.
Claude had a comment about performance that I'll leave as a suggestion:
The huge_1024_f32_f32 benchmarks may now measure something different (benchmark_device_partition.hpp:297, :481; benchmark_device_select.hpp:273).
- Cause: the new functors store the threshold as a full DataType, where the old lambdas captured only a 4-byte float.
- Effect: for custom_huge_type<1024,float,float>, each predicate is now passed into the kernel as a 1 KB argument, most of it the uninitialised data[] array. The three-way
partition benchmark builds two functors (predicate_one and predicate_two, around :557), so it carries about 2 KB where it used to carry 8 bytes.
- Risk: timings for these benchmarks can shift against the baseline even though the predicate logic is unchanged, which weakens the PR's "no performance regression" check.
- Fix: store a float threshold and convert it inside operator():
template struct less_than_op {
float threshold;
bool operator()(const T& v) const { return v < static_cast(threshold); }
};
| } | ||
|
|
||
| // These two functors exist for build time optimization of predicate-based benchmarks. | ||
| // These benchmarks require a predicate function. Defining that prodice function via |
There was a problem hiding this comment.
Thanks, corrected to "predicate". That was a large miscalculation :)
| // - This means that each unique instantiation of the benchmark will ultimately cause a unique kernel to be generated. | ||
| // - Since the predicate-based benchmarks are instantiated with different probabilities, there is a unique kernel | ||
| // generated for each unique probability value, even though the actual kernel code is identical. | ||
| // Defining the predicates as functors here (that take the theshold/flag type as a template parameter) means that the same |
There was a problem hiding this comment.
typo: theshold -> threshold
This is a good point, thanks. I will hold off on merging this change until I can run the benchmark before and after the change to see exactly how much of a runtime regression we're talking about. |
To follow up, benchmarking shows that the runtime differences before and after this change are pretty minimal. With this in mind, I think we can go ahead and merge it. |
TheRock Submodule Bump ActivityNewest first
|
Motivation
We are looking for opportunities to to reduce rocPRIM build time. This change attempts to do that for several benchmark targets.
Jira ID: AIPRIMS-263, EXSWSTRHPC-338
Technical Details
This change introduces two functors for build time optimization of several benchmarks that use lambda functions in a particular context. Here's an example of the optimization for benchmark_device_select:
Several benchmark_device_select benchmarks require a predicate function. Defining that predicate function via a lambda function (within the benchmark's
run()implementation) creates this problem:rocprim::select, and its type is passed as a template argument topartition_impl.Defining the predicates as functors (that take the theshold/flag type as a template parameter) means that the same type is passed into each benchmark instance's call to rocprim::select. This can significantly reduce the number of kernels that need to be generated.
This change implements the above optimization for the following targets:
benchmark_device_selectbenchmark_device_partitionTest Plan
Test Result
Submission Checklist