{"id":797,"date":"2026-01-18T12:30:09","date_gmt":"2026-01-18T16:30:09","guid":{"rendered":"https:\/\/www.fitzsim.org\/blog\/?p=797"},"modified":"2026-01-18T19:49:47","modified_gmt":"2026-01-18T23:49:47","slug":"gfx1201-on-power9","status":"publish","type":"post","link":"https:\/\/www.fitzsim.org\/blog\/?p=797","title":{"rendered":"gfx1201 on POWER9"},"content":{"rendered":"\n

AI GoF disclaimer: I don’t expect this blog post to contribute to frontier AI gain-of-function research or I would refrain from publishing it. Please consider supporting Doom Debates<\/a> to improve the quality of discourse around the risks of frontier AI research, and MIRI<\/a> to try to mitigate the risks<\/a>.<\/em><\/p>\n\n\n\n

I have been wanting to experiment with open weights language models on the Talos II<\/em><\/a>.<\/p>\n\n\n\n

I have a gfx803<\/code> card that I always wanted to use for compute, but it is now out-of-support for ROCm<\/em>. I have made progress getting a gfx1201<\/code> card working on this machine and I wanted to write up all the interesting error messages for reference.<\/p>\n\n\n\n

I took a risk and bought a new GPU, the AMD Radeon AI Pro R9700<\/em> (ASRock Creator 32GB<\/a>), without knowing if I could get it working with the Talos II<\/em> mainboard, which is now seven years old.<\/p>\n\n\n\n

First I realized my existing power supply did not have enough free connectors; I needed a new “modular” power supply for the GPU’s new-style 12v-2x6<\/code> power connector (which is actually a 16 pin connector, with an array of 2 x 6 main big pins and 4 little pins at the top). That prerequisite project was nerve-racking but successful. Physically, the card fit fine in the mainboard and EATX<\/em> chassis.<\/p>\n\n\n\n

With the latest Debian Trixie<\/em> kernel driver, the card showed up as a PCIe<\/em> device in lspci<\/code> (validating the physical installation) but without displaying the card’s name. I figured the driver was not new enough to recognize the card’s product identifier. I read online that a Debian<\/em>-derivative’s 6.17<\/code> kernel recognized the card on a different CPU architecture, so I temporarily enabled the Debian<\/em> testing<\/code> repository, installed linux-image<\/code>, and rebooted. Now lspci<\/code> displayed the card’s name, so that was progress. But as a side effect of the kernel upgrade, my virtual machines failed to start up. The libvirtd<\/code> message was:<\/p>\n\n\n\n

qemu-system-ppc64: Can't support 64 kiB guest pages with 4 kiB host pages with this KVM implementation<\/code><\/pre>\n\n\n\n
It turned out Debian<\/em> ppc64le<\/code> had changed the default page size from 64KiB<\/code> to 4KiB<\/code>.  Debian<\/em> though, with its characteristic flexibility, still provided a 64KiB<\/code> page-size linux-image<\/code> variant.  With that the virtual machines worked again and the GPU continued to be recognized.<\/p>\n\n\n\n
Next I shifted to userspace; the Debian<\/em>-packaged rocminfo<\/code> segfaulted early during its initialization, so I looked upstream and found TheRoc<\/em>k<\/a>.<\/p>\n\n\n\n
I had lots of initial trouble with TheRock<\/em>‘s CMake<\/code> monorepo\/subprojects; I am not yet sure what’s up with that, but I suspect it may be ppc64le<\/code>-specific. That said, I was able to make progress by building individual subprojects one-by-one (this is probably a better approach anyway, at this stage of porting).<\/p>\n\n\n\n
Eventually I got amd-llvm<\/code> bootstrapped, built with a minimal configuration with Trixie<\/em>‘s gcc<\/code> 14.2.0<\/code>.  Then I built amd-llvm<\/code> with itself, in the TheRock<\/em>-recommended configuration, except for the PowerPC<\/code> and AMDGPU<\/code> targets.  Next I built rocminfo<\/code>. It segfaulted in the same place as Debian<\/em>‘s package! Some debugging resulted in a patch<\/a> to accommodate ppc64<\/code>‘s vDSO<\/code> naming; that eliminated the segault.<\/p>\n\n\n\n
Then rocminfo<\/code> ran and showed both the CPUs as “Agents” 0 and 1. But no sign of the GPU.<\/p>\n\n\n\n
I further debugged rocminfo<\/code> and found it was traversing sysfs<\/code>, and specifically the AMD<\/em> Kernel Fusion Driver<\/em> (kfd<\/code>) topology. The card did not have an entry there.<\/p>\n\n\n\n
I looked at dmesg<\/code> and noticed:<\/p>\n\n\n\n
[...] amdgpu 0033:03:00.0: amdgpu: Error parsing VCRAT\n[...] kfd kfd: amdgpu: Error adding device to topology\n[...] kfd kfd: amdgpu: Error initializing KFD node\n[...] kfd kfd: amdgpu: device 1002:7551 NOT added due to errors<\/code><\/pre>\n\n\n\n
First I tried building and updating a .deb<\/code> of the linux-firmware<\/code> from its Git<\/code> repository<\/a>, to rule out the parsing error being caused by an outdated binary-only firmware blob. (This is my one disappointment with the ROCm<\/em> stack; it would be great if the firmware and firmware toolchains were free software.) Rebooting with the new firmware produced the same result.<\/p>\n\n\n\n
I looked at the kernel source for that driver, and noticed extra debug printk<\/code>s. Debian<\/em> helpfully enables the CONFIG_DYNAMIC_DEBUG<\/code> kernel option.  I tried dynamically reloading the amdgpu<\/code> driver and various PCIe<\/em> and GPU reset approaches, but I could not get the card back to its after-boot state.  I would have to reboot to test each change.<\/p>\n\n\n\n
I added amdgpu.dyndbg=\"+p\"<\/code> to the kernel command line, and that gave me some extra kfd<\/code> messages; with those I narrowed down the failure to the IO link<\/code> entry of the Virtual Component Resource Association Table<\/em> (VCRAT<\/code>).<\/p>\n\n\n\n
I re-reviewed dmesg<\/code> and, earlier than the parsing error, there was another clue:<\/p>\n\n\n\n
[...] amdgpu: IO link not available for non x86 platforms<\/code><\/pre>\n\n\n\n
That message was printed during the creation of the CPU VCRAT<\/code> (in kfd_create_vcrat_image_cpu<\/code>). That was the #else<\/code> branch of a platform-specific #ifdef<\/code>. kfd_create_vcrat_image_gpu<\/code> which<\/mark> did not have a corresponding #ifdef<\/code>; “this could explain the subsequent parsing failure on the VCRAT<\/code> IO link entry, on ppc64le<\/code>, a non-x86<\/code> platform”, I thought.<\/p>\n\n\n\n
rocminfo<\/code> still failed though, albeit in a new way:<\/p>\n\n\n\n
hsa api call failure at: \/TheRock\/rocm-systems\/projects\/rocminfo\/rocminfo.cc:1329\nCall returned HSA_STATUS_ERROR_OUT_OF_RESOURCES: The runtime failed to allocate the necessary resources. This error may also occur when the core runtime library needs to spawn threads or create internal OS-specific events.<\/code><\/pre>\n\n\n\n
The co-timed dmesg<\/code> messages were:<\/p>\n\n\n\n
[...] amdgpu 0033:03:00.0: amdgpu: bo 00000000bdd46d97 va 0x0ffffffbfe-0x0ffffffc1d conflict with 0x0ffffffc00-0x0ffffffe00\n[...] amdgpu: Failed to map VA 0xffffffbfe000 in vm. ret -22\n[...] amdgpu: Failed to map bo to gpuvm<\/code><\/pre>\n\n\n\n
I analyzed the section of kernel driver code that generated those messages and noticed the use of AMDGPU_GPU_PAGE_SIZE<\/code> in range calculations. It is hard-coded to 4096<\/code>.<\/p>\n\n\n\n
I had a hunch that the driver needed the kernel’s page size to match.  I did a quick side quest to change all my virtual machines to use 4KiB<\/code> pages<\/a>, reconfigured my custom Debian<\/em> kernel for 4KiB<\/code> pages, and rebooted again.<\/p>\n\n\n\n
Now the virtual machines loaded, and finally rocminfo<\/code> showed the card’s information<\/a>!<\/p>\n\n\n\n
[...]\n*******                  \nAgent 3                  \n*******                  \n  Name:                    gfx1201                            \n  Uuid:                    GPU-6413e1798933ffe0               \n  Marketing Name:          AMD Radeon Graphics\n[...]<\/code><\/pre>\n\n\n\n
I think Debian<\/em>‘s decision to use 4KiB<\/code> pages is sensible, likewise amdgpu<\/code>‘s assuming 4KiB<\/code> pages, so I’m happy to have done this reconfiguration. I was only using 64KiB<\/code> pages because it was the default when I first installed the operating system on the Talos II<\/em>.<\/p>\n\n\n\n
The rest of the process was a grind through TheRock<\/em> subprojects with a bunch of build failure workarounds. The hardest one was fixing static_assert<\/code> failures about __bf16<\/code>, reported by clang<\/code>, when building hipblaslt<\/code>:<\/p>\n\n\n\n
In file included from \/TheRock\/rocm-libraries\/projects\/hipblaslt\/tensilelite\/include\/Tensile\/DataTypes.hpp:42:\nIn file included from \/opt\/rocm\/include\/hip\/hip_fp8.h:30:\nIn file included from \/opt\/rocm\/include\/hip\/amd_detail\/amd_hip_fp8.h:67:\n\/opt\/rocm\/include\/hip\/amd_detail\/amd_hip_bf16.h:155:15: error: static assertion failed due to\n      requirement 'sizeof(__bf16) == sizeof(unsigned short)'\n  155 | static_assert(sizeof(__bf16) == sizeof(unsigned short));\n      |               ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~\n\/opt\/rocm\/include\/hip\/amd_detail\/amd_hip_bf16.h:155:30: note: expression evaluates to\n      '0 == 2'\n  155 | static_assert(sizeof(__bf16) == sizeof(unsigned short));\n      |               ~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~<\/code><\/pre>\n\n\n\n
Debugging led me to a workaround<\/a>; this must be a configuration issue with how I built amd-llvm<\/code>. It needs further investigation, but building amd-llvm<\/code> with that hack meant that I could successfully build hipblaslt<\/code>.<\/p>\n\n\n\n
The other TheRock<\/em> dependencies were easier to build; I continued until I had all the dependencies of llama.cpp<\/code> as-configured for ROCm<\/em>, per this wiki entry<\/a>.<\/p>\n\n\n\n
At this point llama.cpp<\/code> built successfully and ran, with good performance!<\/p>\n\n\n\n
Here is an example run with a simple prompt<\/a>:<\/p>\n\n\n\n
$ uname -srom\nLinux 6.17.13 ppc64le GNU\/Linux\n$ build\/bin\/llama-cli -fa 1 -m ~\/models\/chatgpt-oss-20b\/gpt-oss-20b-F16.gguf \nLoadLib(libhsa-amd-aqlprofile64.so) failed: libhsa-amd-aqlprofile64.so: cannot open shared object file: No such file or directory\nggml_cuda_init: GGML_CUDA_FORCE_MMQ:    no\nggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no\nggml_cuda_init: found 1 ROCm devices:\n  Device 0: AMD Radeon Graphics, gfx1201 (0x1201), VMM: no, Wave Size: 32\n\nLoading model... \n\n\n?? ??\n?? ??\n?? ??  ???? ????????  ????    ????? ????? ?????\n?? ?? ????? ?? ?? ?? ?????    ??    ?? ?? ?? ??\n?? ?? ????? ?? ?? ?? ????? ?? ????? ????? ?????\n                                    ??    ??\n                                    ??    ??\n\nbuild      : b7533-c18428423\nmodel      : gpt-oss-20b-F16.gguf\nmodalities : text\n\navailable commands:\n  \/exit or Ctrl+C     stop or exit\n  \/regen              regenerate the last response\n  \/clear              clear the chat history\n  \/read               add a text file\n\n\n> write hello world in common lisp\nwrite hello world in common lisp\n\n<|channel|>analysis<|message|>We need to answer: \"write hello world in common lisp\". The user wants a \"Hello, world\" program in Common Lisp. Likely they want code snippet. We should provide simple Common Lisp code: `(print \"Hello, World!\")` or `(format t \"Hello, World~%\")`. Provide explanation. Also maybe show how to run with SBCL. Let's produce a response.<|end|><|start|>assistant<|channel|>final<|message|>Here\u2019s a minimal \u201cHello, World!\u201d program in **Common?Lisp**.  \nYou can copy it into a file (e.g. `hello.lisp`) and run it with any Common?Lisp implementation such as SBCL, CLisp, or CCL.\n\n```lisp\n;; hello.lisp\n;; Common Lisp \"Hello, World!\" example\n\n;; Print the string to the default output stream.\n(format t \"Hello, World!~%\")\n```\n\n### How to run it\n\n```bash\n# Using SBCL\nsbcl --script hello.lisp\n\n# Using CLisp\nclisp hello.lisp\n\n# Using CCL\nccl hello.lisp\n```\n\n### What each part does\n\n- `format` \u2013 a versatile printing function.\n- `t` \u2013 the target output stream (`*standard-output*`).\n- `\"Hello, World!~%\"` \u2013 the string to print, where `~%` inserts a newline.\n- `--script` (SBCL) or similar options let the interpreter treat the file as a script and exit automatically after execution.\n\nThat\u2019s all you need to see \u201cHello, World!\u201d on your terminal in Common?Lisp!\n\n[ Prompt: 8.7 t\/s | Generation: 103.1 t\/s ]\n\n><\/code><\/pre>\n\n\n\n
I am very happy with this result!  I will see how stable it is in the coming days.<\/p>\n\n\n\n
Now that I have a working proof-of-concept I will try to upstream some patches and ideally make the top-level TheRock<\/em> build “just work” on ppc64le<\/code> Debian<\/em>.<\/p>\n\n\n\n
Thank yous:<\/em><\/p>\n\n\n\n
It was time to recompile the Linux<\/em> kernel. Debian<\/em> makes this surprisingly easy; I followed the official instructions<\/a> to build a custom kernel .deb<\/code> with my attempted fix<\/a> applied to the amdgpu.ko<\/code> module. Another reboot and no more VCRAT<\/code> parsing failure message in dmesg<\/code>. That seemed like more progress. (Perhaps a more proper solution would be to add IO link<\/code> support to ppc64le<\/code> upstream; I don’t know if there is an equivalent POWER9<\/em> capability, hardware-wise. For my purposes, I have not yet needed an IO link<\/code>.)<\/p>\n\n\n\n