Bug: `ggml_set_rows` faults on HIP/ROCm after many reuses of a cached single-token graph (Strix Halo / gfx1151)

[howard0su]

Environment

• GPU: AMD Strix Halo, gfx1151, wavefront size 32
• Backend: ggml-hip (ROCm), built with -DGGML_HIP=ON, AMDGPU_TARGETS=gfx1151
• llama.cpp / ggml pinned at commit 53859c2
• Kernel: ggml/src/ggml-cuda/set-rows.cu (k_set_rows / k_set_rows_quant), compiled through HIP
• CUDA (same graph, same code path) does not fault.

Symptom

• We build a cached single-token attention/decode graph once and re-execute it every
  decode step via ggml_backend_graph_compute, updating only input tensors (row
  indices, positions, embeddings) between runs.
• The graph contains several ggml_set_rows nodes that write one row into persistent
  state / KV ring buffers each step:
  • KV compressor state: set_rows(state_kv, kv_cur, state_rows) and
    set_rows(state_score, sc_cur, state_rows)
  • Raw SWA KV ring: set_rows(raw_kv, kv_f32, raw_kv_rows)
  • Compressed cache flush: set_rows(comp_cache, pooled, comp_rows)
• On HIP, after enough tokens (the destination row index advances / wraps the ring
  buffer), the compute faults (GPU memory access fault / incorrect writes). The same
  path is stable on CUDA and stable on HIP when we rebuild the graph dynamically each
  token.

What changes across reuses (the only moving parts)

• The src1 index tensor of set_rows (state_rows, raw_kv_rows, comp_rows) is
  updated each step via ggml_backend_tensor_set with a new int64 destination row
  (e.g. kv_start % n_swa, token_pos / ratio, (ratio + pos_mod)), so
  dst_row = *(src1 + ...) reads a different offset in the destination buffer each run.
• Destination tensors (state_kv, state_score, raw_kv, comp_cache) are persistent
  buffers that are written in-place and then read by later ops in the same graph.

Suspected root cause (for AMD to confirm)

The cached graph is allocated once and re-run; only leaf input tensors change. Two
candidate issues on the HIP backend:

1. set_rows writing in-place into a persistent destination that is also consumed
   downstream in the same reused graph may expose a missing dependency / synchronization
   or aliasing assumption that only manifests on the HIP scheduler (RDNA3.5, wave32),
   not on CUDA.
2. The set-rows.cu kernel uses int64 index math and
   blockDim.x * blockIdx.x + threadIdx.x; on gfx1151 (wave size 32, which we must
   force via -DDFLASH_WAVE_SIZE=32) there may be a launch-config / index-overflow /
   write-out-of-bounds edge when dst_row points into a high offset of the ring buffer.

Reproduction sketch

1. Build a small graph: a persistent 2D buffer dst (e.g. [width, n_rows]), one
   set_rows(dst, src_row, idx) node, then an op that reads dst.
2. Allocate once; run graph_compute in a loop, changing only idx (the int64
   destination row) each iteration to sweep across all rows / wrap a ring.
3. On HIP/gfx1151 this should reproduce the fault / incorrect writes after N iterations;
   on CUDA it stays correct.

Current workaround (in our code)

• We disabled reuse of the cached single-token decode graph on HIP only and fall back to
  rebuilding the graph dynamically each token. Non-HIP backends keep graph reuse.
• Detection: ggml_backend_name(backend) contains "HIP" / "ROCm".
• Reference:
  • server/src/deepseek4/deepseek4_graph.cpp:3250
    (reuse_decode_attn = n_tokens == 1 && !ds4_backend_is_hip(backend))
  • set_rows sites: deepseek4_graph.cpp:655-656, :744, :1009

Questions

1. Is re-executing a pre-allocated ggml graph containing in-place ggml_set_rows into a
   persistent buffer (that is also read later in the same graph) expected to be safe on
   the HIP backend? Any required synchronization / allocator flags?
2. Is there a known set-rows.cu issue on RDNA3.5 / wave32 (gfx1151) with int64 index
   writes into high offsets of a large destination?
3. Does forcing DFLASH_WAVE_SIZE=32 interact with the CUDA-derived kernel launch
   geometry in a way that could cause out-of-bounds writes?