Re: MIOpen package LFS files
Hi Cory and Christian,
> I suppose the next question is whether we can filter out these assembly kernels and still have a working MIOpen library.
As discussed earlier, we’re going to test MIOpen availability when the ASM codes are excluded.
Here I’ll first briefly introduce the MIOpen workflow, then highlight my changes to MIOpen source code to accommodate with this decision, and present the next possible steps last.
MIOpen Workflow
~~~~~~~~~~~~~~~
+-------+ +-------------+ +---------------+ +----------+
|Problem+-->|Find solution+-->|Load/JIT kernel+-->|Run on GPU|
+-------+ +-------------+ +---------------+ +----------+
Users consult MIOpen with the problem (conv, pooling, activation, etc.) and input specifications (filter size, stride, padding, etc.).
MIOpen then internally iterates all possible “solutions” for that problem and specs, where each solution will be materialized with a kernel written in HIP, OpenCL or AMD GPU native asm code.
As we discussed in this thread before, there exists pre-compiled kernel caches providing the compiled binary code objects for each GPU architecture.
If the cache is unavailable, or the demanding kernel is missing, MIOpen will try to compile the kernel at runtime (JIT) to produce the target GPU code.
Finally, the code will be sent to GPU for computing.
Notable changes
~~~~~~~~~~~~~~~
* Disabling of GCN ASM codes
Although we exclude all the asm codes, MIOpen itself still tries to load and compile the asm code based on the solution.
There is a environment variable MIOPEN_DEBUG_GCN_ASM_KERNELS [1] for MIOpen debugging intending to skip kernels from specific sources (e.g., ASM, HIP, OpenCL).
After searching for its usage in the codebase, I patch the library initialization to disable MIOpen to load the asm codes, by making the solution “not applicable” if its kernel is built from ASM code [2].
* Integration of hipRTC
Note that src/kernels contains some HIP source codes, which will be JIT-ed during runtime.
This is achieved by hipRTC [3], which requires C++ dev environment.
So I explicitly add them to the runtime dependencies of package libmiopen1 [4].
And MIOpen now doesn’t specify the target GPU architecture of the generated code, which leads to runtime errors like this:
hip Error: Please provide architecture for which code is to be generated.
So I patch the JIT process to add backend architecture option [5].
Next directions
~~~~~~~~~~~~~~~
test directory contains some functional tests for MIOpen.
They are built with gtest and each unit corresponds to one executable.
But they are not packaged into the libmiopen-tests package as for now, maybe I should resolve this and integrate them into the -tests package.
Besides, there are some permission issues of ROCm podman container environment, I’ll work with Christian to cope with them and enable autopkgtest for MIOpen.
Lastly, sorry for the delayed reply as there are some personal works this week.
Best,
Xuanteng
[1]: https://rocm.docs.amd.com/projects/MIOpen/en/latest/how-to/debug-log.html#filtering-by-build-method
[2]: https://salsa.debian.org/rocm-team/miopen/-/merge_requests/1/diffs#1dcea28fd7822b07a1284c0ec755b65b7575c28c
[3]: https://salsa.debian.org/rocm-team/miopen/-/blob/master/src/hipoc/hipoc_program.cpp?ref_type=heads#L285-290
[4]: https://salsa.debian.org/rocm-team/miopen/-/merge_requests/1/diffs#58ef006ab62b83b4bec5d81fe5b32c3b4c2d1cc2
[5]: https://salsa.debian.org/rocm-team/miopen/-/merge_requests/1/diffs#9ecfa1f1dc4536389a3f9e705160a00e07881aad
Reply to: