Experiments with using ROCM on the FW16 AMD

Framework Laptop 16 Linux
62

I use my C++ program to imitate James’s Rust program, in which I can easily change to different methods of memory allocation (Memory management — HIP 7.1.0 Documentation). James is using hipMalloc(), which allocates memory in the device (probably it falls back to one of the host ones). This hipMalloc() is not safe when we directly write into device memory from the host; instead, if it is used, we should use a copy operation to copy a vector allocated in host memory space. But, using hipMallocManaged() or hipHostMalloc() eliminates the non-recoverable crash during the allocation. So, I used hipMallocManaged() and never had a non-recoverable crash again.

I used rocgdb to debug the program. With 50000 matrix sizes, like James, I also get sometimes gpu reset, sometimes not.

I also tried to employ recommendations from HIP Debugging — HIP 6.0.0 Documentation , in particular, I use HIP_LAUNCH_BLOCKING=1 to serialize kernel enqueue. It somehow helped, but not always. So I can’t say the reset is because of a parallel kernel issue.

But one thing I noticed is that something happens either at the calling of rocblas_cgemm or after at hipDeviceSynchronize function. Unfortunately, I was not able to step into rocblas_cgemm() with rocgdb to check the point where the reset occurs. At least, currently, I don’t know how to do it.

I would say that memory allocations were always without the problem.

A successful 50k multiplication within rocgdb:

$ printenv | grep -i block
$ rocgdb --args ./matmulblas 50000
GNU gdb (GDB) 18.0.50.20250927-git
Copyright (C) 2025 Free Software Foundation, Inc.
License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html>
This is free software: you are free to change and redistribute it.
There is NO WARRANTY, to the extent permitted by law.
Type "show copying" and "show warranty" for details.
This GDB was configured as "x86_64-pc-linux-gnu".
Type "show configuration" for configuration details.
For bug reporting instructions, please see:
<https://www.gnu.org/software/gdb/bugs/>.
Find the GDB manual and other documentation resources online at:
    <http://www.gnu.org/software/gdb/documentation/>.

For help, type "help".
Type "apropos word" to search for commands related to "word"...
Reading symbols from ./matmulblas...
(gdb) r
Starting program: /home/xyz/test/amdgpu_test/matmulblas 50000
warning: os_agent_id 16142 (`Device 1002:15bf.dd'): architecture gfx1103 not supported.
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
matrix rows: 50000 cols: 50000 size: 2500000000 ram: 20000 MB
[New Thread 0x7fffeefff6c0 (LWP 73660)]
[New Thread 0x7fffee7fe6c0 (LWP 73661)]
[Thread 0x7fffee7fe6c0 (LWP 73661) exited]
Matrix A (input):
(2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) 
00 00 00 40 00 00 00 00 
Matrix B (input):
(2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) (2 + 0i) 
00 00 00 40 00 00 00 00 
Matrix C (input):
(0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) (0 + 0i) 
[New Thread 0x7fffed3ff6c0 (LWP 73820)]
❌️ Error while mapping shared library sections:
'memory://73646#offset=0x555555f22b00&size=47528': AMDGCN architecture 0x44 is not supported.
[New Thread 0x7fffe79ff6c0 (LWP 73821)]
[Thread 0x7fffe79ff6c0 (LWP 73821) exited]
❌️ Error while mapping shared library sections:
'memory://73646#offset=0x555555cfdc60&size=261704': AMDGCN architecture 0x44 is not supported.
❌️ Error while mapping shared library sections:
'memory://73646#offset=0x55555605c0a0&size=55992': AMDGCN architecture 0x44 is not supported.
Matrix C (output):
(200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) (200000 + 0i) 
calc1:       8972 us
sync:        1863986983 us
[Thread 0x7fffed3ff6c0 (LWP 73820) exited]
[Thread 0x7fffeefff6c0 (LWP 73660) exited]
[Inferior 1 (process 73646) exited normally]

Look at the warning and errors (not affecting the program run).

1 Like
63

@Mahboob_Karimian

Hi. Regarding the details about hipMalloc().

My program was specifically to test with an APU, the 7840HS in the FW16.

As such, the GPU has no device memory, it is 100% shared with the CPU memory, so if you write values to the CPU memory, you don’t have to transfer anything to the GPU.

I assumed that rocblas_cgemm() was doing any pointer magic, if it needed to.

64

Yes, it seems here:

from line 327 to 329, they check, but I’m not sure what else (flags, etc) is set when appropriate function is called.

FYI: I use AMD 7840U (780M) with 64GB RAM, no dgpu.
Do you guys have idea what to do next? Where to dig into?

65

for llama.cpp on rocm the magic is here:

=> no need for gtt upgrade only define GGML_CUDA_ENABLE_UNIFIED_MEMORY=ON to get all RAM useable . And do not une mmap to. (–no-mmap)

66

I have the same issues with the MES errors/ GPU resets and was able to resolve them to get Pytorch + ROCm working as expected on a similar system (Ryzen 5 8640HS, 760M iGPU)

I am not 100% sure of which steps directly helped it work out in the end, but here is what I tried:

  1. Add the following to the kernel arguments: amdgpu.mcbp=0 amdgpu.vm_update_mode=0 amdgpu.cwsr_enable=0

All tests prior to adding amdgpu.cwsr_enable=0 had the same crashes. This leads me to believe that the main variable that solved the issue is amdgpu.cwsr_enable=0

  1. Add the following environment variables:
export PYTORCH_ROCM_ARCH=gfx1100
export HSA_OVERRIDE_GFX_VERSION=11.0.0

Let me know how this goes, and hopefully all works as expected! Cheers

Note that I am using rocm-6.4.3 from the fedora repos. Kernel is 6.17-rc7, but I suspect it would work on 6.16.x as well. Note that I checked the kernel git source manually and I do not believe 6.17 has the lr_compute_wa. Anyhow rpm -q linux-firmware says that I have the August 8th version so it wouldn’t work anyway.

1 Like
67

something like the “amdgpu.cswr_enable” kind of makes some sense as a possible cause.

Say, if one was doing a long calculation and waiting around for a “sync” to finish, the GUI might wish to update the desktop when the user select something, so it might need to interrupt the long calculation in order to update the display as a result of user actions.

If that context switching has a bug, it could cause what we are seeing.

I did see far fewer (but not zero. I had one crash across many tests) crashes when I disabled the GUI and ran my test program on the command line, at runlevel 3 (the one without a GUI desktop).

Another odd observation I saw, is sometimes the crashed are a “once” crash.

I.e. you power on, it crashes some app, but then that app restarts and runs ok without crashes after that.

1 Like
68

I only applied this; it did not hang at the beginning, but as soon as I clicked on the already opened system monitor, the reset came in. But I think it is a nice option to reduce. i will test this further.

I have the same experience, but not always.

69

Yup, I was going through the kernel module arguments and thought it might be reasonable especially given that mid command buffer preemption amdgpu.mcbp=0 was a fix for a similar issue in the past

Seeing that you have updated to linux-firmware-20250917, I would try downgrading to linux-firmware-20250808 as I suspect there are issues introduced with it. This is because when I was originally testing with linux-firmware-20250917 I noticed that it crashed immediately with a batch size of 128 whereas the prior version allowed it to work without issue for far longer. (When testing without the cwsr workaround).

I am not on a new enough kernel to see whether the patch that sets the lr_compute_wa bit interferes somehow with this so ruling this out may be helpful.

Do note that when testing with the cwsr workaround, there are no observable differences with either version in PyTorch, at least for a short 5-10 minute run.

Are you using patched kernel and linux-firmware-20250917? Try downgrading linux-firmware or using an older kernel so that lr_compute_wa is not used. (to replicate my setup where I do not use a patched kernel with lr_compute_wa)

1 Like
70

Yes, I’ve used linux-firmware 20250901.git993ff19b-0ubuntu1 to update my firmware files from Ubuntu 25.0X builds because for 24.04 there is no updated package (latest is for 2024).
BTW, as James also mentioned, after adding amdgpu.cwsr_enable=0, I only have one crash with my C++ or James’s Rust rocblas matrix multiplication program with 50k-element matrices. Once I had the crash, I ran my and Jame’s test program multiple times, and it did not hang or reset at all.
Right after this, I’ve tested ComfyUI, loading multiple models, unloading them, and running long workflows, which did not trigger any GPU reset.

So, Thanks to you, @gliu20 ,I guess this workaround is working with a high success rate, but it is not a fix for the problem.

I should repeat my tests, and if something happens or the results change, I’ll post here.

1 Like
71

Hi.

I can confirm that “amdgpu.cwsr_enable=0" helps.
After a new reboot.
my “rocm-rust 50000” test program crashes amdgpu the first time one runs it after rebooting, but then, after that, it succeeds each time.
Without the “amdgpu.cwsr_enable=0" it fails every time.

1 Like
72

Do you guys have problem with rocm loading quantized models in comfyui ? I see models weights are converted to bfloat16 from Q4_K_M or Q8 when I check their memory footprints, and the log:

Warning: This gguf model file is loaded in compatibility mode 'sd.cpp' [arch:flux]
gguf qtypes: F32 (476), Q8_0 (304)
model weight dtype torch.bfloat16, manual cast: None

To recall, I have gfx1103 but I override it as gfx1100 for pytorch.

Since RDNA3 is supporting int8, I don’t know if it is the problem of sdcpp library or it is the rocm, or it is gfx1103? In addition, I don’t get how Vulkan backend deal with this when using by LMStudio for text models.

I would like to here your comments on this.

1 Like
73

I was curious if ollama can use the GPU at the same time as a game…Well, it can, but it’s not perfect. This happens when I ask the LLM a question. Kinda funny:

74

Some details on CWSR from AMD:

CWSR is a fundamental feature in our command-submission/execution model and memory management model with user mode queues. When you contrast that with kernel mode command submission it’s like comparing preemptive scheduling and cooperative scheduling on CPUs.

Without CWSR we can only preempt dispatches at wavefront granularity. I.e. an endless loop in a shader kernel would hang the GPU and require a GPU reset to recover. CWSR allows preempting wavefronts at instruction granularity, which enables arbitrarily long-running or persistent shader kernels on the GPU without introducing stalls or hangs affecting the GPU scheduler and memory manager. We also use the saved GPU state from CWSR in the ROCm debugger. Therefore disabling CWSR is not recommended. We don’t test our SW with CWSR disabled. This mode only exists for hardware bring-up, triage and debugging purposes.

CWSR is implemented by a combination of CP firmware and a trap handler, which is like an interrupt handler running on the GPU compute units. The driver is only responsible for allocating the memory needed to save the GPU state (registers, LDS, and metadata needed to restarte waves later). When a request comes to the CP to preempt a queue that’s currently executing a dispatch, it first tries to drain running waves from the compute units. If a grace period expires it sends a trap to all the remaining waves and the trap handler does the heavy lifting of saving register and LDS state to memory. The trap handler code can be found in our kernel mode driver. The CP writes restart information about the wave-fronts to the control stack.

AFAIK resuming a queue is handled by the CP firmware and supporting fixed-function HW. After restoring the wave, register and LDS state it resumes execution of each wave at the instruction that was previously interrupted by the trap handler.

In this context CP is another name for MEC. CP stands for “command processor”. MEC stands for “Micro Engine Compute”.

1 Like
75

From that CWSR description above, and the kernel log saying the MES is not completing all the items in the queue. My first guess is a race condition.
If an item in the queue can be preempted, it makes queue management a little more difficult. Instead of simple managing 3 states (waiting in queue, running, completed), one needs to handle 4 states (waiting in queue, preempted-partially complete, running, completed).
There is then the saving of the state, and associating that state with the correct item in the queue. Then the restoring of that state to continue on.
On top of this, there is kernel code saving state, but also amd firmware (binary blob) that is involved during restore, making it very difficult to track down the cause.

My guess is that there maybe be a race conditions where the gpu is preempted but the queue state mismatches. Or the saving state is missing a register it needs.
Note: this is a APU 7840HS (no dGPU) problem, so might not affect other GPUs.

I guess one way to test this is by simply putting a count on the amount of preemptions. Do the preemptions happen while rocm is not running?

As “If a grace period expires it sends a trap”. So, the trap handler might only be being called during “rocm-rust”, thus hitting a piece of code that is rarely exercised otherwise. The system seems to function OK with “CWSR=0”, this further supports the idea that the trap is maybe not often used. So, again, having a trap counter would prove this.

I don’t think this problem is related to GPU power states, because the GPU is pretty much maxed out during my rocm-rust 50000 test.

76

Hello, I have uploaded Pytorch build, rocm, and flash attention artifacts for gfx1103 here:

With this, you don’t have to do any overrides.
Flash attention works too; the Triton version is very simple to install and use. My observation is that it does not expedite much.

77

The build Doc and how long it needs to build and tool ninja, swap space (24GB+).

So many unknown requirement to build it.

78

This topic was automatically closed 180 days after the last reply. New replies are no longer allowed.