Re: MIOpen package LFS files
Hi Xuanteng,
excellent work! I only added minor comments to the MR. But the result so
far is already really good.
On 2024-05-25 18:00, Xuanteng Huang wrote:
> 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.
Thank you, this was quite helpful.
> Notable changes
> ~~~~~~~~~~~~~~~
>
> * Disabling of GCN ASM codes
> * 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].
LGTM, though I look forward to Cory's feedback on this last change.
> 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.
You already did this in the meantime, looks good. Let me know if you'd
like me to run tests on other architectures.
I think we're getting pretty close to an upload to experimental.
Best,
Christian
Reply to: