windows support

[SimonDanisch]

This is a very hacky work in progress for windows support!
This finds most libraries, but I'm not sure which dll contains hsa_init...
From googling, i got the impression that it may be in amdhip64.dll (https://docs.amd.com/projects/HIP/en/latest/how_to_guides/debugging.html), but that doesn't find the symbols.
Seems like AMD hip/hsa works quite differently on windows ;)
Any pointers appreciated what to try next!

[vchuravy]

HSA is Linux only, the equivalent tech on the windows side is called amdpal (I think?)

[vchuravy]

E.g. https://github.com/GPUOpen-Drivers/pal

The question is: Can we avoid HSA completely and just use HIP?

[SimonDanisch]

Crazy... That seems like a bigger project then, since the current AMDGPU.jl seems to heavily rely on HSA for lots of functionality, or does it just look like that?

HSA runtime is unavailable, compilation and runtime functionality will be disabled.

That seems pretty essential to the package :D

[pxl-th]

Actually the only place where we directly use HSA is the HostCall implementation, which is not needed for the most part.
So just relying on HIP should work

[pxl-th]

Also, did you manage to get anything working? Last time I tried ROCm 5.5+ it was complaining about LLVM opaque pointers, so I was going to deal with that first and then move on to Windows support

[SimonDanisch]

I'm really unfamiliar with the whole stack & code base, so not sure where to start if init fails without HSA... If you give me an MWE without HSA I can try if it works with hip only

[radudiaconu0]

dont't we need miopen for this to work? As far as i know miopen isn't avalable on windows

[vchuravy]

dont't we need miopen for this to work? As far as i know miopen isn't avalable on windows

For the first layer of support? No. If you want to use Flux on AMD then yes.

[vchuravy]

I'm really unfamiliar with the whole stack & code base, so not sure where to start if init fails without HSA... If you give me an MWE without HSA I can try if it works with hip only

The first think to look at is src/runtime/hsa_device.jl We need to obtain the ISA using HIP instead of HSA. Other than that I think you can comment out everything with host call for now.

The first important step is to compile "a kernel", and to be able to allocate memory and then launch the kernel.
Most of that was moved to HIP from HSA

[radudiaconu0]

any updates here?

[pxl-th]

I have a somewhat working private branch which I'll push a bit later.

There are couple things left to do though.
Mainly we need to recompile artifact device libraries for Windows (to support mixed-mode).
We currently have them only for Linux, otherwise we get incompatibility errors due to LLVM 15 vs 16 until Julia supports LLVM 16.

Also I don't see if there is an equivalent to HSA_OVERRIDE_GFX_VERSION which would allow us to work with gfx 10xx GPUs.

[radudiaconu0]

@pxl-th i don't know for windows is neccesary. On windows rx 7900xtx is officialy supported

[pxl-th]

Superseded by #542.
Still needs more work to work on Windows though.