Skip to content

Use CandidateSelectionModel heuristic for conv_hip_implicit_gemm_3d_grouped_*_xdlops solvers #3923

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

Open
wants to merge 85 commits into
base: develop_deprecated
Choose a base branch
from
Open

Use CandidateSelectionModel heuristic for conv_hip_implicit_gemm_3d_grouped_*_xdlops solvers #3923

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
85 commits
Select commit Hold shift + click to select a range
49a0bbb
Added first attempt at fdeep versions of 3d ktn models.
amd-bartgips Jul 25, 2025
d015559
Added co-pilot generated new model class for selecting the best candi…
amd-bartgips Jul 25, 2025
289a1f2
performed dot product directly in cpp code, removed candidate_selecto…
amd-bartgips Jul 25, 2025
2c155cc
Added the machinery to use the new candidate selection heuristics for…
amd-bartgips Jul 25, 2025
10722a3
added split_k functionality
amd-bartgips Jul 25, 2025
fe19eb9
refactored by moving some helper functions outside the main function
amd-bartgips Jul 25, 2025
2e000f3
improved loading of metadata for candidate selection model
amd-bartgips Jul 28, 2025
dbf021b
split off own metadata class to keep it distinct from legacy version
amd-bartgips Jul 28, 2025
11ec8d5
Added new versions of CS models (+new metadata)
amd-bartgips Jul 28, 2025
05871e3
Added methods for preprocessing input and kernel_convigs for CS model…
amd-bartgips Jul 28, 2025
b007d9e
fixed naming of CandidateSelectionMetadata and Ptrs variables
amd-bartgips Jul 28, 2025
7574ca7
Added new model and metadata class to header file
amd-bartgips Jul 28, 2025
e795096
removed unused function
amd-bartgips Jul 28, 2025
c0a95a8
removed duplicate definition of CandidateSelectionMetadata
amd-bartgips Jul 28, 2025
b89cae2
* altered SelectBestCandidate to avoid errors.
amd-bartgips Jul 28, 2025
b09b18f
removed unused function
amd-bartgips Jul 28, 2025
cb846ec
Build now works.
amd-bartgips Jul 28, 2025
efeece1
moved new candidateSelection code to its own files (and sub namespace)
amd-bartgips Jul 29, 2025
33244fc
refactored general 3D conv kernel tuning functions into their own fil…
amd-bartgips Jul 29, 2025
25b25a8
Cleaned up unused include
amd-bartgips Jul 29, 2025
279b87a
removed superfluous includes
amd-bartgips Jul 29, 2025
3a97cec
3D conv heuristics (KTN part) (#3918)
amd-bartgips Jul 29, 2025
2cbb65d
made model methods public. Made output vectors instead of fdeep tensors.
amd-bartgips Jul 29, 2025
5473d54
started on 3D tuning heuristics test function
amd-bartgips Jul 29, 2025
de3b965
Reorganised declaration of classes and function (in .hpp vs .cpp)
amd-bartgips Jul 30, 2025
ece42d0
added ai_candidate_selection.cpp to cmakelists
amd-bartgips Jul 30, 2025
af37d7a
added encodekernelconfig test. Changed assert calls into explicit if …
amd-bartgips Jul 30, 2025
f482b50
fdeep does not support batch processing, so loop over candidates
amd-bartgips Jul 30, 2025
5926d4a
implemented multi-thread prediction for candidate condig encoding by …
amd-bartgips Jul 30, 2025
ebaf508
added more (edge case) test functions
amd-bartgips Jul 30, 2025
cd2e1f8
renamed test function to make clear this focusses on candidate_selection
amd-bartgips Jul 30, 2025
7672acb
renamed kernel_tuning_utils to reflect that this is for 3D conv
amd-bartgips Jul 30, 2025
fa0e796
fixed typo
amd-bartgips Jul 30, 2025
5dcd20a
fixed includes
amd-bartgips Jul 30, 2025
c76fc0b
fixed GetFeatures3D to match current model (but should probably be do…
amd-bartgips Jul 30, 2025
d5d42e1
Started with test functions for the utilities and running the full mo…
amd-bartgips Jul 30, 2025
3bc4428
Working CandidateSelectionModel + corresponding test functions (#3922)
amd-bartgips Jul 31, 2025
7329093
Merge branch 'miopenff/3d-heuristic' into bartgips/3d_conv_test_funcs
amd-bartgips Jul 31, 2025
aa955ca
removed leftover old test function for candidateselectionmodel. Added…
amd-bartgips Jul 31, 2025
4143633
added passing of ctx to HeuristicInit
amd-bartgips Aug 1, 2025
29724c9
solver_name was renamed to DeviceGroupedConvBwdWeight. TODO: investig…
amd-bartgips Aug 4, 2025
c7c6470
added some try/miopen_throw statements for easier debugging
amd-bartgips Aug 4, 2025
d0d1826
wip tests
amd-bartgips Aug 4, 2025
3056d84
wip, such that Jeonghyun can pull files to complete build
amd-bartgips Aug 4, 2025
beb900a
removed reference to non-existing HasValidStream method for ctx
amd-bartgips Aug 4, 2025
d3a288f
Switched to improved GetKernelAsTokens functions to parse kernel config
amd-bartgips Aug 4, 2025
b6190e5
better split up whole kernel name as a single token. Start of kernel …
amd-bartgips Aug 4, 2025
316c1b2
updated new models + metadata format
amd-bartgips Aug 5, 2025
93923f6
fixed some build warnings:
amd-bartgips Aug 5, 2025
cc35e51
switched to using a map for GetFeatures3D so we can more easily drop …
amd-bartgips Aug 5, 2025
3bef719
updated test functions to follow new way to deal with model input in …
amd-bartgips Aug 5, 2025
3248446
improved encodeKernelParams to perform proper conversion to float vec…
amd-bartgips Aug 5, 2025
7b1ac26
removed superfluous logging
amd-bartgips Aug 5, 2025
ed3b15f
Removed dealing with constants from Encode*WithFdeep (should be done …
amd-bartgips Aug 5, 2025
372a791
implemented new kernel_str_mapping attribute for metadata
amd-bartgips Aug 5, 2025
b4007cb
Added new metadata files
amd-bartgips Aug 5, 2025
93bc21b
fixed tests that deal with input features (map instead of vector)
amd-bartgips Aug 5, 2025
1bd3f0b
updated metadata files
amd-bartgips Aug 5, 2025
e7a6354
WIP altering test functions to deal with new ways of encoding input a…
amd-bartgips Aug 5, 2025
57f02b3
Replaced references of K0PerBlock by KPerBlock
amd-bartgips Aug 6, 2025
58f1f3e
Added support for NaN tokens;
amd-bartgips Aug 6, 2025
d1e3cf5
removed debug print statements
amd-bartgips Aug 6, 2025
ccfad4c
fixed all current test functions
amd-bartgips Aug 6, 2025
606e67e
Added some extra tests suggested by copilot
amd-bartgips Aug 6, 2025
98fe338
Improved the encoding of kernel configs
amd-bartgips Aug 6, 2025
321d8e1
All tests are now passing.
amd-bartgips Aug 6, 2025
5f359f2
Added test for full MIOpen pathway/api
amd-bartgips Aug 6, 2025
9ead9a1
Replaced Filter* function with FillHeuristicKernels function
amd-bartgips Aug 6, 2025
19570bd
removed old non-gtest test functions
amd-bartgips Aug 6, 2025
cd39f00
copied functionality to other 2 direction solvers (Bwd, Fwd). Added c…
amd-bartgips Aug 6, 2025
5141c18
Implemented working heuristics for all 3 conv 3d solvers.
amd-bartgips Aug 7, 2025
657aaf8
converted std::cerr into MIOPEN_THROW. Made metadata const
amd-bartgips Aug 7, 2025
6466043
Added latest model files, added TODO statement
amd-bartgips Aug 8, 2025
7e06292
undid erroneous additions of "const" prefix to member variables
amd-bartgips Aug 8, 2025
4755e4a
removed hard-coded datatype, now properly deduced form problem descri…
amd-bartgips Aug 8, 2025
ebed5cc
Add 3D convolution AI heuristics implementation (#3925)
JH-Leon-KIM-AMD Aug 12, 2025
3cccb91
Merge remote-tracking branch 'origin/miopenff/3d-heuristic' into bart…
amd-bartgips Aug 12, 2025
344b27b
replaced funciton-local static dummy_ctx with a GetDummyCtx function
amd-bartgips Aug 12, 2025
cfc273e
fixed discrepancy in name of GetDummyCtx helper function
amd-bartgips Aug 12, 2025
7012c6c
fixed mismatch in declaration of GetNanToken
amd-bartgips Aug 12, 2025
d4af913
fixed discrepancy in name of GetDummyCtx helper function
amd-bartgips Aug 12, 2025
4b1035d
move to range-based for loop to remove warning
amd-bartgips Aug 12, 2025
9e138ca
cache env var to make it thread safe
amd-bartgips Aug 12, 2025
a2de4cd
Added a mutex for thread-safe initialization of MIOPEN_AI_FDEEP_USE_S…
amd-bartgips Aug 12, 2025
0b203de
changed year in copyright statement
amd-bartgips Aug 12, 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
3 changes: 3 additions & 0 deletions src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -349,6 +349,9 @@ set( MIOpen_Source

if(MIOPEN_ENABLE_AI_KERNEL_TUNING OR MIOPEN_ENABLE_AI_IMMED_MODE_FALLBACK)
list(APPEND MIOpen_Source conv/heuristics/ai_heuristics.cpp)
list(APPEND MIOpen_Source conv/heuristics/metadata_3d.cpp)
list(APPEND MIOpen_Source conv/heuristics/ai_candidate_selection.cpp)
list(APPEND MIOpen_Source conv/heuristics/ai_conv_3d_kernel_tuning_utils.cpp)
list(APPEND MIOpen_Source anyramdb.cpp)
endif()

Expand Down
Loading