Replace nested static_for lambdas with compile-time search helper - #6696
Merged
illsilin merged 2 commits intoJun 2, 2026
Merged
Conversation
Add sequence_find_value and find_in_tuple_of_sequences helpers using O(1) template depth via constexpr array lookup and pack expansion. Replace nested static_for lambdas in TensorDescriptor and TensorAdaptor: - GetTransformAndItsUpperDimension: use find_in_tuple_of_sequences - InitializeElementSize: replace generate_tuple lambda with pack expansion This eliminates the bulk of applier::operator() instantiations that previously came from nested lambda closures in the tensor descriptor transform search. Conflict-resolved rebase of ROCm/composable_kernel#3600 onto develop. Supersedes #4287. Generated-by: Claude Code (claude-sonnet-4-6)
4 tasks
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
illsilin
approved these changes
Jun 2, 2026
illsilin
deleted the
users/tenpercent/ck/tensor-descriptor-lambda-elimination
branch
June 2, 2026 23:13
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Jun 2, 2026
Replace nested static_for lambdas with compile-time search helper (#6696) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary - Add `sequence_find_value` and `find_in_tuple_of_sequences` compile-time search helpers with O(1) template depth - Replace nested `static_for` lambdas in `TensorDescriptor::GetTransformAndItsUpperDimension` and `InitializeElementSize` - Apply same optimizations to `TensorAdaptor` Supersedes #4287. Conflict-resolved rebase of #3600 onto current develop. ## Motivation The `TensorDescriptor` and `TensorAdaptor` classes had excessive template instantiation from: 1. Nested `static_for` loops with lambdas creating unique closure types at every call site 2. `generate_tuple` with lambdas causing per-type instantiation overhead The new helpers use constexpr array lookup and pack expansion instead of recursive template patterns, achieving O(1) template depth. ## Results (`example_grouped_conv_fwd_xdl_fp16`, n=10, interleaved, `-j1`, `-ftime-trace`) | TU | Baseline (mean) | New (mean) | Delta | Wilcoxon p | Mann-Whitney p | |----|-----------------|------------|-------|-----------|---------------| | `grouped_conv_fwd_xdl_fp16` (host) | 14,886 ms | 13,353 ms | **-10.3%** | **0.002** | **0.0002** | | `grouped_conv_fwd_xdl_fp16` (device) | 27,762 ms | 25,629 ms | **-7.7%** | **0.002** | **0.0002** | | **Total (all TUs)** | **57,732 ms** | **54,030 ms** | **-6.4%** | | | Unrelated TUs (`device_memory`, `host_tensor`, `convolution_parameter`) show no significant difference (p > 0.3), serving as negative controls. ### Methodology - 10 interleaved runs (baseline₁, new₁, baseline₂, new₂, ...) on the same node to eliminate ordering/warmup bias - Wilcoxon signed-rank test (paired, non-parametric) and Mann-Whitney U test (unpaired) - Built with patched clang (LLVM 22) on ctr2-alola-compile-11, `-j1` for accurate per-TU timing - Raw data available in Slurm job 275230 results ## Test plan - [x] 11 unit tests added (5 for `sequence_find_value`, 6 for `find_in_tuple_of_sequences`) - [x] Compile-time benchmark with statistical significance (p < 0.01) - [ ] Full CI Tracking issue: #4229
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
sequence_find_valueandfind_in_tuple_of_sequencescompile-time search helpers with O(1) template depthstatic_forlambdas inTensorDescriptor::GetTransformAndItsUpperDimensionandInitializeElementSizeTensorAdaptorSupersedes #4287. Conflict-resolved rebase of ROCm/composable_kernel#3600 onto current develop.
Motivation
The
TensorDescriptorandTensorAdaptorclasses had excessive template instantiation from:static_forloops with lambdas creating unique closure types at every call sitegenerate_tuplewith lambdas causing per-type instantiation overheadThe new helpers use constexpr array lookup and pack expansion instead of recursive template patterns, achieving O(1) template depth.
Results (
example_grouped_conv_fwd_xdl_fp16, n=10, interleaved,-j1,-ftime-trace)grouped_conv_fwd_xdl_fp16(host)grouped_conv_fwd_xdl_fp16(device)Unrelated TUs (
device_memory,host_tensor,convolution_parameter) show no significant difference (p > 0.3), serving as negative controls.Methodology
-j1for accurate per-TU timingTest plan
sequence_find_value, 6 forfind_in_tuple_of_sequences)Tracking issue: #4229
Generated-by: Claude Code (claude-sonnet-4-6)