[FusilliPlugin] Lift and shift fusilli-plugin: shark-ai -> rocm-libraries#2252
Conversation
fusilli-plugin shark-ai -> rocm-libraries
fusilli-plugin shark-ai -> rocm-librariesfusilli-plugin: shark-ai -> rocm-libraries
|
Added |
|
@amd-jnovotny Note that the docs team was pinged because of the *.md file included in this "lift and shift". |
sjain-stanford
left a comment
There was a problem hiding this comment.
LGTM. The main thing is between this PR and PR 3/N (when fusilli-plugin is removed from shark-ai), how do you manage contributions to the codebase - freeze the shark-ai mirror, or cherry-pick in both?
|
I was thinking cherry-pick in both. I don't like it, but hopefully we don't have to spend a lot of time with the sate in superposition. |
amd-jnovotny
left a comment
There was a problem hiding this comment.
I'll approve to unblock this. We'll want to clean up the readme when it's ready for general release.
|
@jayhawk-commits sound good. I approved to unblock it. We'll worry about fixing up the read me later. |
…ROCm#2275) Adds basic CI, CMake build scrips, and Readme.
This PR adds "all the boring bits" of the fusilli plugin. It adds a basic engine plugin that hipDNN can load, which ensures that all of the required API definitions are present, and those APIs called as part of load function. The PR also tests for plugin loading, logging, and basic functionality.
This PR implements [`hipdnnEnginePluginGetApplicableEngineIds`](https://github.com/ROCm/rocm-libraries/blob/7e29d65a8db488a5e038c3616398144d7c2289a5/projects/hipdnn/sdk/include/hipdnn_sdk/plugin/EnginePluginApi.h#L90-L115) and [`hipdnnEnginePluginCreateExecutionContext`](https://github.com/ROCm/rocm-libraries/blob/7e29d65a8db488a5e038c3616398144d7c2289a5/projects/hipdnn/sdk/include/hipdnn_sdk/plugin/EnginePluginApi.h#L170-L191), which together completes `hipDNN` -> `fusilli` graph translation for the initial plugin goal of e2e run of single node ConvFprop graph.
This PR implements `hipdnnEnginePluginExecuteOpGraph` enabling graph/plan execution. We can now complete the initial plugin goal of e2e run of single node ConvFprop graph. Note: a simple e2e integration test is added. This should be expanded to a parameterized matrix of test, once we're on the latest hipDNN [example](https://github.com/ROCm/rocm-libraries/blob/develop/projects/hipdnn/plugins/miopen_legacy_plugin/integration_tests/IntegrationGpuConvForward.cpp). Tickets: [ROCm#2363](nod-ai/amd-shark-ai#2363), [ROCm#2375](nod-ai/amd-shark-ai#2375).
Update build to pull hipdnn from `rocm-libraries` and update clang to version 20 from TheRock. Co-authored-by: Sambhav Jain <sambhav.jain@amd.com>
…ld. (ROCm#2449) This PR fixes `fusilli-plugin` build failure when fusilli benchmarking (`FUSILLI_BUILD_BENCHMARKS`) was enabled, but fusilli testing (`FUSILLI_BUILD_TESTS`) was disabled: [link to CI failure](https://github.com/nod-ai/shark-ai/actions/runs/18206945208/job/51839515287?pr=2410). Both _should_ be disabled, but ideally the build wouldn't break if one wasn't. The issue was: both fusilli tests and fusilli benchmarks depend on `libutils`, but the target was only defined if tests were enabled.
…_agent_enumerator (ROCm#2488) Got this building on my SharkWorkstation with an RX 9070XT (RDNA4) targeting the `gfx1201` arch with minimal changes. The `Backend` enum doesn't need to be specific to the arch, and the iree-compile flag `--iree-hip-target` can just be derived using `rocm_agent_enumerator` (which lists the gfx arch from the GPUs visible) based on https://iree.dev/guides/deployment-configurations/gpu-rocm/#choosing-hip-targets.
… monorepo. (ROCm#2521) Signed-off-by: MaheshRavishankar <mravisha@amd.com>
Add Fusilli+IREE as a kernel provider and JIT engine for hipDNN
When building `hipDNN` from source it handles downloading and configuring transitive dependencies, while an installed version requires the consuming package to provide needed dependencies. In the plugin's current form, the former is more convenient.
…nv3d (ROCm#2560) Compiler fix: iree-org/iree#22320 Issue: iree-org/iree#22312
Adds support for async execution and explicit device/stream selection and updates fusilli plugin to use new API: - Handle creation supports device id and external HIP stream parameters. - Graph execution is async on AMDGPU backend. - Fusilli plugin now uses stream if set by user.
618b1c3 to
05f8da9
Compare
Co-authored-by: Jim <jimguo12@amd.com> [ROCm/composable_kernel commit: 57f4974]
Motivation
This is a direct "lift and shift" of the hipDnn plugin for fusilli (also known as
fusilli-plugin) + git history originally developed inshark-ai.The (current) plan is to move
fusilli-pluginin a few stages:rocm-librariesCI (it also won't be running in TheRock CI).fusilli-pluginshould build and run inrocm-librariesCI. Note: TheRock CI will still not be enabled.fusilli-pluginfromshark-ai.RFC0004-Fusilli-IREE-Kernel-Provider-hipDNN.md) and enable TheRock CITechnical Details
Direct lift and shift using
git subtreeall code has been previously reviewed inshark-ai.Test Plan
Tests of all functionally should be current, however initially CI will not test
fusilli-plugin. CI will come in follow up PR.Test Result
Submission Checklist