Navi 31 support (7900 XTX / XT)

[jellersby]

I just installed hip-runtime-amd 5.4.1-2 (and friends) from the community-testing repo. To test, I tried running both rocminfo and the small sample C++ program provided. Below are the results. Any ideas? Thank you.

$ /opt/rocm/bin/rocminfo
ROCk module is loaded
rocminfo: /build/hsa-rocr/src/ROCR-Runtime-rocm-5.4.1/src/core/runtime/amd_gpu_agent.cpp:339: void rocr::AMD::GpuAgent::AssembleShader(const char*, AssembleTarget, void*&, size_t&) const: Assertion `code_buf != NULL && "Code buffer allocation failed"' failed.
zsh: IOT instruction (core dumped)  /opt/rocm/bin/rocminfo

$ ./hiptest
hiptest: /build/hsa-rocr/src/ROCR-Runtime-rocm-5.4.1/src/core/runtime/amd_gpu_agent.cpp:339: void rocr::AMD::GpuAgent::AssembleShader(const char*, AssembleTarget, void*&, size_t&) const: Assertion `code_buf != NULL && "Code buffer allocation failed"' failed.
zsh: IOT instruction (core dumped)  ./hiptest

[jellersby]

Update: I think it's because I've got a Navi 31 card (gfx1100). It doesn't seem to be supported yet by the entire ROCm stack (including MIOpen). Will continue to experiment with HSA_OVERRIDE_GFX_VERSION env var and see how far I can get.

More info:

ROCm/MIOpen#1904

ROCm/ROCm#1813

[wsippel]

rocminfo should work on RDNA3, without HSA_OVERRIDE_GFX_VERSION. It works for me with the AMD provided binary ROCm packages (opencl-amd from the AUR). If I had to guess, maybe @tpkessler simply didn't enable the gfx1100 targets for the current builds?

More advances stuff like MIOpen, PyTorch or Tensorflow will require ROCm 5.5, but HIP definitely already works. I ran a bunch of Blender benchmarks no problem.

[tpkessler]

Packages are built with the default flags from AMD. Maintaining our own list of supported architectures (separately for every package) is not practical as it may change from one release to another. Furthermore, I can't test it as I don't have access to all possibly supported GPUs (not even a Navi 31).

[jellersby]

I just did a fresh install of hip-runtime-amd (via the shiny new miopen-hip; thanks @tpkessler!) and I am indeed able to run rocminfo successfully without setting HSA_OVERRIDE_GFX_VERSION. It looks like gfx1100 support is baked in after all, at least in HIP. I'm not entirely sure why it worked this time instead of a few weeks ago, but I'll take the win.