ci: Properly install rocwmma for hip builds

[IMbackK]

No description provided.

[IMbackK]

We require rocwmma to be properly installed for #16221 see #16221 (comment)

Unfortunately i dont have a windows environment to work in, and have been having a hard time getting the rocwmma buildsystem to work on the windows ci, see cb59db2 for the attempt.
For this reason i have disabled rocwmma on the windows builds for now.

[slaren]

rocwmma is also used in the binary releases in release.yml and rocm.Dockerfile. However, if I understand correctly, this would reduce performance, and is only necessary to workaround an issue with ROCM 7, which is still in preview, and is not used in the binary releases either.

[IMbackK]

rocm 7.0 has been released. disabeling rocwmma on windows will reduce performance yes, but the way rocwmma was used on windows previously was just plain wrong, you cant use the library from its source directory without installing it, even though it is header implemented.

[slaren]

Ok, so I guess the question for the binary releases is whether using ROCM 7 or rocwmma is more important.

[IMbackK]

ideally someone with a windows machine would come and fix the install - but as i say i dont have such a machine.

[IMbackK]

Ok, so I guess the question for the binary releases is whether using ROCM 7 or rocwmma is more important.

the changes in the pr #16221 make llamacpp build against rocm 7 at all, and break the windows ci (beacuse the rocwmma installation there is incomplete). So this is not just about the binary releases.

[slaren]

Ok, but we still cannot break or degrade the performance of the windows binaries just for the sake of supporting ROCM 7. This needs to be fixed.

[IMbackK]

I find this a bit difficult i explicitly merged rocwmma support with it disabled by default, which it still is, as i knew that supporting it across architectures and os's is quite difficult as its very unstable - which has proven true often, with rdna4 being broken on rocwmma < 2 and > 1.4 and CDNA being broken on rocwmma 2.0 and rocwmma having no official windows support at all with it being excluded from windows distributions of rocm.

Others then decided to enable rocwmma on ci, i would not have done so.

[IMbackK]

I will give it another try, but ultimately this should not block #16221

[slaren]

I have been giving it a try, but it seems more trouble than it is worth. Considering that this is a header-only library and the only thing that the cmake script does (for us) is generate rocwmma-version.hpp, maybe there is some simpler workaround to determine the version.

The rocmma version seems to be solely stored in CMakeLists.txt, so parsing this file as such should do it:

$env:ROCWMMA_VERSION = if ((Get-Content -Path "CMakeLists.txt" -Raw) -match 'set \( VERSION_STRING\s+"?([0-9.]+)[^0-9.]+') { $matches[1] }

[IMbackK]

I have been giving it a try, but it seems more trouble than it is worth. Considering that this is a header-only library and the only thing that the cmake script does (for us) is generate rocwmma-version.hpp, maybe there is some simpler workaround to determine the version.

The rocmma version seems to be solely stored in CMakeLists.txt, so parsing this file as such should do it:

$env:ROCWMMA_VERSION = if ((Get-Content -Path "CMakeLists.txt" -Raw) -match 'set \( VERSION_STRING\s+"?([0-9.]+)[^0-9.]+') { $matches[1] }

I dont really like this hack much at all, even though this solves the intimidate problem, ultimately what we are doing here on ci is just plain wrong and there is no guarantee that problem remains confined to the missing rocwmma-version.hpp generation as rocwmma evolves.

IMO we should just accept that rocwmma is not supported on windows by amd atm, disable it for now and then petition amd to support rocwmma on windows and add it to the rocm distribution.

[slaren]

IMO we cannot make suboptimal releases for Windows just because it is not convenient, because otherwise the result will be that people will be forced to look for the releases elsewhere. It would fracture the community and damage the reputation of llama.cpp. I don't think it is too concerning to add this hack to the HIP backend, and in any case my understanding is that @JohannesGaessler plans to replace WMMA entirely with mma, so at worst it would be temporary.

[JohannesGaessler]

I have already purchased and installed an RX 9060 XT, I've also ordered a MI100 that is set to arrive in a few weeks. I intend to add support for the AMD "WMMA" instructions to the kernel in fattn-mma-f16.cuh. To be clear, the naming is confusing since AMD has "WMMA" instructions and NVIDIA has a "WMMA" interface (which is bad and I therefore want to phase it out). The end result will likely be that rocWMMA (which enables the use of NVIDIA's WMMA interface on AMD) will be obsolete on RDNA.

On CDNA (MI100) there are "MFMA" instructions which work a bit differently but the end result should be the same as with the AMD WMMA instructions.

So yes, if you give me a few weeks I will likely be able to make it so that rocWMMA can be removed entirely. The WMMA kernel will still be needed for V100s, I'm in the process of purchasing one of those for development. I don't think it would make sense to maintain the WMMA kernel only for MUSA.