[Issue]: hipHccModuleLaunchKernel in amdhip64_7.dll access violation

[kdrapel]

Problem Description

hipHccModuleLaunchKernel in amdhip64_7.dll (ROCm 10.0.0) dereferences a null pointer at a fixed struct offset during kernel launch, causing a reliable STATUS_ACCESS_VIOLATION (0xC0000005) crash on Windows. The fault reproduces identically — same Windows crash-classification hash — across two separate amdhip64_7.dll builds (10.0.3661.0 and 10.0.3679.0), including after a full, verified-clean ROCm 10 reinstall with no mixed-version libraries present.

The faulting instruction reads a field at a fixed offset (0x388 in build 10.0.3661.0, 0x390 in build 10.0.3679.0 — an 8-byte shift consistent with a struct layout change between builds, e.g. one added or removed pointer field) from a base pointer (rcx) that is null. This appears to be a kernel-launch descriptor or module handle that was never populated before being dereferenced.

The crash occurs during the client application's startup device-capability/VRAM-fitting probe — i.e. on essentially the first real kernel-launch call HIP makes for the session, before any weight loading or actual inference work begins.

This has been reproduced with two structurally distinct models (a REAP-pruned MoE model and an unrelated dense 9B model), indicating the fault is general to gfx1200 under ROCm 10.0.0 at this kernel-launch stage, rather than specific to any particular model.

This build worked previously with ROCM 7.14, no crash.

Exception:      0xC0000005 (Access violation - Read)
Faulting module: amdhip64_7.dll
Faulting func:  hipHccModuleLaunchKernel
Offset:         +0xf4220 (build 10.0.3679.0) / +0xdc090 (build 10.0.3661.0)
Faulting instr: mov rcx, qword ptr [rcx+390h]   ; (+388h in the 3661.0 build)
Fault address:  0x0000000000000390 (0x388 in earlier build)
rcx at fault:   0x0 (null)
Failure.Hash:   {b8d1dc6f-3fdd-cc00-5915-402226ee5221}   (identical across both builds)

Operating System

Windows 11 (build 10.0.26100.1)

CPU

Ryzen 7 7800X3D

GPU

AMD Radeon RX 9060

ROCm Version

ROCM 10.0.3661.0

ROCm Component

No response

Steps to Reproduce

Environment:

• ROCm 10.0.0 (Windows)
• GPU: AMD Radeon RX 9060 (gfx1200)
• OS: Windows 11 (build 10.0.26100.1)
• Client: llama.cpp (llama-server.exe), built from source against ROCm 10 via CMake, HIP backend (ggml-hip.dll)

Steps:

1. Install ROCm 10.0.0 on Windows 11 for a gfx1200 GPU (RX 9060), following the standard tarball/installer flow.
2. Build llama.cpp from source with the HIP backend, configured against this ROCm 10 install (confirm hipblas.dll, rocblas.dll, rocsolver.dll, libhipblaslt.dll all resolve from the same ROCm 10 install directory — no mixed versions).
3. Run llama-server.exe with the following flags against a GGUF model (a REAP-pruned/expert-pruned MoE model, Q3_K_M quantization, was used when this was found):

   llama-server.exe -m <model>.gguf -ngl 99 -sm none -fa 1 -np 1 --temp 0.4 --no-mmap \
     --top-k 20 --top-p 0.8 --min-p 0.05 -t 12 -tb 8 --prio 2 --prio-batch 2 --poll 100 \
     --poll-batch 1 --presence-penalty 1.1 --ctx-size 131072 -n -1 --reasoning off \
     --host 127.0.0.1 --port 9090 --jinja --cache-type-k q4_0 --cache-type-v q4_0 \
     --alias Qwen3.6-35B-A3B
   

4. Process crashes during model load, before inference begins, with the access violation described above.

Reproduction rate: 100% — every attempt so far, across two ROCm 10 runtime builds and a clean reinstall.

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

rocminfo --support output

Paste output here

Additional Information

• Call stack (client → HIP runtime), captured via WinDbg:

  llama_server_impl!llama_server
   → llama_common!common_init_from_params
    → llama_common!common_init_result::common_init_result
     → llama_common!common_fit_params
      → llama_common!common_get_device_memory_data   ← device VRAM/kernel-launch probe
       → llama!llama_init_from_model
        → ggml_hip!ggml_backend_cuda_host_buffer_type
         → amdhip64_7!hipRegisterTracerCallback+0x...  (internal)
          → amdhip64_7!hipHccModuleLaunchKernel        ← FAULT HERE
  

• Confirmed not model-specific: reproduced the identical crash (same fault location, same call path through common_get_device_memory_data → hipHccModuleLaunchKernel) with a completely different model — a dense, non-MoE 9B model (not REAP-pruned, standard architecture) — ruling out any connection to the original model's pruned-MoE expert count or quantization scheme. This indicates the bug is general to gfx1200 + ROCm 10.0.0 at this kernel-launch stage, not specific to any one model or model family.
• Ruled out: mixed/stale ROCm library versions — initial reproduction happened to have a leftover ROCm 7.14 hipblas/rocblas/rocsolver/libhipblaslt install shadowing the ROCm 10 runtime; after a full clean reinstall with all libraries confirmed loading from a single ROCm 10 tree, the identical crash still occurs, ruling this out as the cause.
• amdhip64_7.dll ships without public symbols, so the exact struct/field involved couldn't be identified beyond the offset shown above. Happy to provide the full .dmp file if useful for triage.
• This looks like it belongs in the closed-source HIP runtime rather than in llama.cpp/ggml (open-source client code) — but since it will affect any HIP-based application hitting this launch path on gfx1200 under ROCm 10, cross-linking a corresponding issue in the llama.cpp tracker may also be worthwhile for visibility.

[JM00NJ]

Related: same crash pattern on ROCm 7.14.0 / gfx1201 (RX 9070 XT) during
PyTorch training. Confirmed via WinDbg — corrupted/garbage pointer in the
profiler callback linked list (not null, but invalid usermode address:
7b2c450fc892ac84).

Workaround: NOP out call r10 at hipLaunchKernel+0x85 in amdhip64_7.dll
(file offset 0x4549B5: 41 FF D2 → 90 90 90). Full analysis in issue #7732 .

[JM00NJ]

Fix script : https://github.com/JM00NJ/amdhip64-hipLaunchKernel-crash-fix

[kdrapel]

Thanks @JM00NJ, my amdhip64_7.dll is the updated version I guess as it did not match your patch sequence but definitely this looks like a similar crash.