Commit a994ea5
[rocm-libraries] ROCm/rocm-libraries#1154 (commit 422e872)
[MIOpen] Implement kernel tuning heuristic model for 3D conv
ops (two tower model) (#1154)
Copied over form branch in old repo:
#3923
## Motivation
With this PR MIOpen should be able to use heuristics for 3D convolutions
on gfx942:
* the parameter selection/kernel tuning of the three
`conv_hip_implicit_gemm_3d_grouped_*_xdlops` solvers.
We have added some `ai_*` files housing the new models and helper
functions, and made some changes to existing files where required.
* ~~the solver selection (i.e. a "3D tunanet")~~ At the moment the 3D
solver selection model is inadequate and the already existing WTI
fallback is preferred. Improving the 3D solver selection heuristics will
be the focus of a future PR
## Technical Details
* `ai_heuristics.cpp` contains all the code already there beforehand,
plust new code that relies on fdeep includes, since fdeep can only be
imported in a single file
* `ai_candidate_selection` contains the actual two-towers (aka
CandidateSelection) "model" and "metadata" classes that do the
computation using floating point vectors.
* `ai_conv_3d_kernel_tuning_utils` contains the machinery one level
higher. That is, how to convert kernel configs and fdb_key input to
float vectors and how to fetch and call the relevant
CandidateSelectionModel. This is shared for all three solvers, so it
made sense to centralise it in this file.
* `kernels/gfx942....` model and metadatafiles. Perhaps these should be
committed using git lfs, but I have not seen this done to other model
files (e.g. Tunanet or KTN), so have not done so here.
* `solver/conv/...cpp` solver_specific files, they ultimately contain
the solver-specific machinery. In this case our 3 solvers rely heavily
on `ai_conv_3d_kernel_tuning_utils` and through that on
`ai_candidate_selection`
* gtest files: should speak for themselves, please have a look. They
test all the new machinery
* solvers.hpp a huge header file that contains declarations for all
solvers (i.e. for the `solver/conv/...cpp` files), so this needed to be
altered as well
## Test Plan
3 new gtest .cpp files are added:
* ~~`test/gtest/conv_ai_3d_heuristics.cpp`
This aims to test all new functionality related to the 3D tunanet (model
+ metadata).~~ While the tests are still there for future use, they will
be skipped since 3D Tunanet model data is no longer included.
* `test/gtest/conv_ai_3d_kernel_tuning_utils.cpp`
Aims to test all new machinery in `ai_conv_3d_kernel_tuning_utils.cpp`
(preprocessing and handling of inputs to the CandidateSelectionModel for
the 3D solvers).
* `test/gtest/conv_ai_candidate_selection_model.cpp`
Test interenal code related to the CandidateSelectionModel and its
metadata.
## Test Result
The `./bin/test_conv_ai_*` tests all succeed without errors when
building and running them on a conductor MI300 node.
Besides manually running all other `./bin/test_*` functions, is there a
better way to perform a full test?
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.1 parent b8eeb30 commit a994ea5
File tree
25 files changed
+4654
-214
lines changed- src
- conv/heuristics
- include/miopen/conv
- heuristics
- kernels
- solver/conv
- test/gtest
25 files changed
+4654
-214
lines changed| Original file line number | Diff line number | Diff line change | |
|---|---|---|---|
| |||
351 | 351 | | |
352 | 352 | | |
353 | 353 | | |
| 354 | + | |
| 355 | + | |
| 356 | + | |
354 | 357 | | |
355 | 358 | | |
356 | 359 | | |
| |||
0 commit comments