[HIP][OpenCL] Do not pass GPU type from the host code, use predefined macros provided by hip-clang instead

[atamazov]

Many of our kernels need to know the target GPU type. Right now we define macros (like MIO_BN_GFX110X) at the host side and pass to the kernel during compilation, via -D options. Therefore, compilation options are changed whenever we introduce support for new GPUs (or deprecate support for the old ones).

🔴 But compilation options are used to build KDB key. Therefore, when we add support for new GPUs, the precompiled binary cache becomes obsolete (and needs to be regenerated!) for currently supported GPUs.

However passing GPU type via options is not required because hip-clang knows the target GPU type, see https://clang.llvm.org/docs/AMDGPUSupport.html. IIRC we already use this approach in the assembly kernels.

Proposal: Remove passing macros from the host code to HIP/OCL compiler. Instead, use the predefined macros provided by hip-clang.

Additional info:

• Example of introducing the new GPU type:
  • [gfx12] add support of gfx12 platforms #3109
• Some discussion (plus proposal to use intermediate macros to represent target features in the kernel sources):
  • Fix issue with multiple targets and remove smfmac tests from unsupported test targets composable_kernel#1372 (comment)

---

[Attribution] @junliume @JehandadKhan @CAHEK7

• quality
• value_middle
• Proposed assignees: None.