Skip to content

[MIOpen] Implement kernel tuning heuristic model for 3D conv ops (two tower model) #1154

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
amd-bartgips merged 79 commits into from
Sep 22, 2025
Merged

[MIOpen] Implement kernel tuning heuristic model for 3D conv ops (two tower model) #1154

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
79 commits
Select commit Hold shift + click to select a range
b39db7f
Copied over all changes from old miopen repo branch
amd-bartgips Aug 12, 2025
63a64f9
renamed "nan_token" into "missing_value_token" to reduce ambiguity
amd-bartgips Aug 12, 2025
f33c7c3
refactored fill_valid_kernel definition and RunParameterPredictionMod…
amd-bartgips Aug 12, 2025
73b8dee
Renamed MIOPEN_AI_FDEEP_USE_SINGLE_THREAD_PREDICT env var to make it …
amd-bartgips Aug 12, 2025
f8391b0
Added split_k values to metadata, refactored such that these values a…
amd-bartgips Aug 12, 2025
452f789
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 12, 2025
5596cd2
added "split_k_values" to metadata for wrw 3d model
amd-bartgips Aug 12, 2025
78c16ec
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 13, 2025
968fd9d
fixed bugs introduced by last refactor
amd-bartgips Aug 13, 2025
eb37904
more fixes after refactor. Also fix related test functions
amd-bartgips Aug 13, 2025
02ccd4a
copied back GetWti function that got lost when copying to monorepo
amd-bartgips Aug 13, 2025
b3e1daf
changed func declaration to avoid warning during build
amd-bartgips Aug 13, 2025
4c906f7
fixed bugs with split_k (for wrw solver)
amd-bartgips Aug 13, 2025
ea4cbac
removed "main" function def from gtest files
amd-bartgips Aug 13, 2025
f8d3cb7
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 13, 2025
3fcaca8
edited tests to follow naming convention
amd-bartgips Aug 13, 2025
190adcb
refactored candidate_selection_model tests for better maintainability
amd-bartgips Aug 13, 2025
7f9fc21
refactored 3d kernel tuning util tests to conform to conventions
amd-bartgips Aug 13, 2025
6382d0d
refactored 3d tunanet heuristics tests to conform to conventions
amd-bartgips Aug 13, 2025
4bc1742
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 13, 2025
39194f9
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 13, 2025
84598bb
Added FP32 as datatype name suffix for GPU test functions
amd-bartgips Aug 13, 2025
d45e92a
simplified GetUseSinglePredictEnv. attempt to avoid thread-safety war…
amd-bartgips Aug 13, 2025
ab95ae8
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 13, 2025
f1ec63a
Used miopen::env to deal with env var. small fix to remove warnings i…
amd-bartgips Aug 14, 2025
7786e65
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 14, 2025
e24c689
added skip if not ai_immed_mode_fallback
amd-bartgips Aug 14, 2025
c87fea3
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 15, 2025
03a1649
Better match conventions: removed skip_tests in indiviual functions
amd-bartgips Aug 15, 2025
179dd2b
removed dangling skip test
amd-bartgips Aug 15, 2025
bd88bf6
Added inline cppcheck suppression for 3d heuristics test function. to…
amd-bartgips Aug 15, 2025
8c2a86e
refactored 3d kernel tuning util tests to work with platforms other t…
amd-bartgips Aug 18, 2025
5a447d2
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 18, 2025
3e13ede
Added fwd heuristic override back in (if present, AI heuristics are s…
amd-bartgips Aug 18, 2025
64a303f
Added InitValidKernels helper function to 3D solvers
amd-bartgips Aug 18, 2025
dccf84d
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 19, 2025
940d893
remove verbose logging for whitespace mismatch
amd-bartgips Aug 21, 2025
6a4f886
Added proper hierarchy for choosing the fallback heuristics
amd-bartgips Aug 25, 2025
61610c6
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 25, 2025
3b5218c
Added MIOPEN_USE_COMPOSABLEKERNEL guard macro to check for CK support…
amd-bartgips Aug 25, 2025
9ed6a60
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 25, 2025
1da8e81
Removed current 3D Tunanet and temporarily disabled it (WTI fallback …
amd-bartgips Aug 26, 2025
4c46652
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 26, 2025
166e006
disable KT heuristics for Fwd 3d conv solver for fp16, bfp16, since w…
amd-bartgips Aug 26, 2025
7e152e5
added no-op stubs for non-ck builds
amd-bartgips Aug 27, 2025
fad931e
explicitly convert fs::path to std::string for compatibility with win…
amd-bartgips Aug 28, 2025
d7228e5
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Aug 28, 2025
48a7392
fixed header to include ai_heuristics.hpp regardless of MIOPEN_ENABLE…
amd-bartgips Sep 1, 2025
88dff31
Added MIOPEN_ENABLE_AI_KERNEL_TUNING guard to test functions so build…
amd-bartgips Sep 2, 2025
56ddf0b
added newline at end of file
amd-bartgips Sep 2, 2025
35e6af6
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 2, 2025
d158e69
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 3, 2025
346140e
Added more compile-time guards for builds without AI heuristics
amd-bartgips Sep 3, 2025
ff9ce90
Update embed.cmake
amd-bartgips Sep 3, 2025
98cf2aa
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 3, 2025
fc7d979
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 8, 2025
24b8ace
Added MIOPEN_INTERNALS_EXPORT to all function definitions used elsewh…
amd-bartgips Sep 8, 2025
60e493c
removed MIOPEN_INTERNALS_EXPORT macros from explicit template instant…
amd-bartgips Sep 8, 2025
296ccda
Also added MIOPEN_INTERNALS_EXPORT macros to relevant header files
amd-bartgips Sep 8, 2025
14e31b1
fixed typo
amd-bartgips Sep 8, 2025
47ef256
added correct includes for MIOPEN_INTERNALS_EXPORT
amd-bartgips Sep 8, 2025
fc54ac8
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 9, 2025
e1a4c05
Added missing MIOPEN_INTERNALS_EXPORT statement
amd-bartgips Sep 9, 2025
b5b2864
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 9, 2025
99cf101
added MIOPEN_INTERNALS_EXPORT to temapletes in proper places in sourc…
amd-bartgips Sep 9, 2025
6e546a4
Removed the MIOPEN_INTERNALS_EXPORT from .cpp source file and moved i…
amd-bartgips Sep 9, 2025
7c0b26f
set and unset env vars in a cross-platform way
amd-bartgips Sep 9, 2025
fa1ac3d
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 9, 2025
bc7502f
fixed typos
amd-bartgips Sep 9, 2025
48941f6
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 9, 2025
49d602e
Added missing endif statement
amd-bartgips Sep 10, 2025
ba2a812
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 10, 2025
5dfda34
fixed calls to new cross-platf utility functions to set env vars
amd-bartgips Sep 10, 2025
4853376
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 10, 2025
cef1d41
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 15, 2025
79d907f
Improved flow of fwd solver and improved/unified logging among the th…
amd-bartgips Sep 19, 2025
e1265d5
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 19, 2025
7b2149e
git commit renamed Disabled test functions to match naming conventions
amd-bartgips Sep 19, 2025
1f731c2
Merge branch 'develop' into silo/feat/3d_conv_heuristics
amd-bartgips Sep 22, 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 projects/miopen/src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -351,6 +351,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