[COMgr] Prepare for COMgr 3.x

[atamazov]

This is intended to resolve #3075. By-products are some small fixes and removal of useless code (e.g HIP builds via COMgr; from now on online HIP builds are supported only via HIPRTC). Details:

• [comgr] Added support for AMD COMgr 3.0.
• By-products:
  • [comgr] Removed support for AMD COMgr older than 1.7.
  • [comgr] Removed comgr::BuildHip().
    • Why: The compilation of HIP kernels via COMgr is broken for a long time. Fixing it requires a lot of efforts (because we must guarantee performance of kernels to be on par with kernels built via HIPRTC) but useless because HIPRTC exists on all systems where we have COMgr.
  • [cmake] Do not allow different values of MIOPEN_USE_COMGR and MIOPEN_USE_HIPRTC.
  • [comgr][NFC][quality] Removed support of unused enum members from to_string()
  • [wingorad fury] Fix precompiled binary cache miss for Winograd Fury. Resolves [gfx11][Solvers][Winograd] ConvWinoFuryRxS v2.4 #2778 (comment).
  • [cmake][NFC] Remove leftovers of support of ROCm older than 5.0.

🌀 Performance testing results

Preconditions:

• Navi21 (gfx1030/36)
• ROCm 6.0.2 docker miopen:ci_a9eebf, Ubuntu 20.04.5 LTS
• BUILD_DEV=On, Release build
• branch amdcomgr-3 2f38063 vs. develop fea5579
• 93 known FP32 convolution configs
• 93 known FP16 convolution configs

Tested modes:

• (0) Immediate mode with default (system) find-db
• (1) Find Mode (DYNAMIC_HYBRID)
• (2) Find Mode (NORMAL)
• (3) Immediate mode with user find-db filled by step (2).

No performance or correctness regressions. Detailed logs & csv files are available upon request.

---

[Attribution] @junliume @JehandadKhan

• enhancement
• urgency_high
• value_high
• Proposed reviewers:
  • @JehandadKhan
  • @junliume
  • @lamb-j
  • @averinevg

[lamb-j]

Left one comment, but otherwise Comgr changes LGTM, thanks!

If you're interested, there are some new actions, languages, and data types that have been added that aren't currently included the MIOpen Enums. See here:

https://github.com/ROCm/llvm-project/blob/amd-staging/amd/comgr/docs/ReleaseNotes.md

I assume they could also be added later, as-needed

[lamb-j]

Apparently I can't comment on unchanged lines, but for Line 370 of comgr.cpp:

We may also remove the AMD_COMGR_DATA_KIND_FATBIN language type for 3.0, as we've removed the associated FATBIN action.

It may be convenient to remove it here to avoid a separate patch

[atamazov]

@lamb-j Yep, so it seems that you will likely remove AMD_COMGR_DATA_KIND_FATBIN from amd_comgr_data_kind_t. Which means that we'll have a build error and therefore this data kind needs to be removed for COMgr 3.x. Thanks!

[atamazov]

@lamb-j Done in a5a7606

[atamazov]

@junliume @CAHEK7 @lamb-j Ready for CI testing and another review round. Someone please launch the CI job, thanks!

[averinevg]

CI job launched

[junliume]

@atamazov can you access the Windows build logs? We need to protect it from breaking too.

C:\home\jenkins\agent\workspace\UIF2_MIOpen_PR-3107\MIOpen\src\comgr.cpp:589:20: error: unused function 'IsWave64Enforced' [-Werror,-Wunused-function]
  | 14824 | 08:40:01 AM | 589 \| static inline bool IsWave64Enforced(const OptionList& opts)
  | 14825 | 08:40:01 AM | \|                    ^~~~~~~~~~~~~~~~

[atamazov]

@junliume

can you access the Windows build logs?

Unfortunately, no access

[atamazov]

@junliume @apwojcik

C:\home\jenkins\agent\workspace\UIF2_MIOpen_PR-3107\MIOpen\src\comgr.cpp:589:20: error: unused function 'IsWave64Enforced' [-Werror,-Wunused-function]

FYI this may happen only if HIP runtime does not support precompiled headers (PCH), which seems suspicious.

[atamazov]

@junliume #3107 (comment) is fixed in 9f73a5b

[shurale-nkn]

@atamazov please merge develop into this branch. Otherwise, it will not pass the CI

[atamazov]

@averinevg thanks for e711774!

[lamb-j]

@atamazov @junliume thanks for implementing and merging this!

What's the typical turnaround time for promotion to amd-master for MIOpen?

[atamazov]

@lamb-j

What's the typical turnaround time for promotion to amd-master for MIOpen?

It depends on many things. At https://github.com/ROCm/MIOpen/commits/amd-master I see that amd-master has been updated several times in the last month. But on the other hand, before that it remained unchanged for a month or so. I guess that @junliume maintains some mailing lists dedicated to MIOpen promotions, so you can ask him to add your email there.