Backport two patches from rocm-systems for gfx1151 (Strix Halo):
hsakmt: bump vgpr count for gfx1151
We apply only the change that adds a GFX_VERSION_GFX1151 define but
not the default vgpr change which causes crashes on old kernels.
The vgpr change was later partially reverted upstream.
hsakmt: Expose and use CWSR and Control stack sizes
Fixes potential mismatches between stack sizes configured in the
amdgpu kernel module and userspace by relying on kernel value when
available. Falls back to hardcoded size based on ISA if kernel is
too old. For gfx1151, also prints a warning that KFD ABI 1.20+ is
recommended. This allows new kernels to work with this runtime
detection mechanism. Only gfx1151 warns loudly because only it has
had the size updated in kernel.
Hoping to resolve MES hang errors on gfx1151.
Found by @06kellyjac, possibly triggered by bug in linux-6.18.4
https://lore.kernel.org/linux-iommu/870872aa-28e9-412a-bac6-8020bf560e4f@amd.com/t/
Resolving this *will not fix the underlying issue*, but does transmute a
segfault into an actual error result.
ROCm error: out of memory
current device: 0, in function stream at /build/source/ggml/src/ggml-hip/../ggml-cuda/common.cuh:1345
hipStreamCreateWithFlags(&streams[device][stream], 0x01)
Eman Resu