DE

Issues / #3358

#3358 [Engine] HIP virtual memory (0.2.0-beta): vmm_test memory fault on gfx1011 (BC-160) with ROCm 7.14

open · @andyelka-creator · 1 Kommentare · Auf GitHub

needs-format

Beschreibung

**Area:** [Engine]
**Desktop app, Terminal or source?** Source (a build)
**Reproduction status:** Reproduced locally
**Strata version:** Strata 0.2.0-beta · engine 0.1.45.3 (tag `v0.2.0-beta`, commit `3d1fb480`)
**Graphics card, memory and system:** 3× AMD BC-160 (Navi 12, gfx1011, 8 GB HBM2 each) · 16 GB RAM · Ubuntu, kernel 7.0.0-38-generic (in-tree amdgpu), HIP 7.14.60850 from TheRock wheels `rocm[libraries,devel]==7.14.0a20260608` (`gfx101X-dgpu` index)
**Model, size and settings:** none — this is `ctest`, no model was loaded

### What happened

Re-test of #2374 on 0.2.0-beta. The original problem is fixed: `hip_q2_zero` gives 1024 values, 0 failures on all three cards.

`vmm_test` now fails. On 0.1.43.1 it reported Passed, but there `vmm.cpp` had no HIP path, so the test took its "skipped" branch. The HIP ranges from 3980a5ce are the first time it really runs on this card.

The failure is at the `cudaMemset` after `a.map_range(3, 7, ...)`, which covers the chunks just remapped into the hole:

```
:0:rocdevice.cpp :3608: ... Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x743b89049000, kernel: __amd_rocclr_fillBufferAligned]
FAIL tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == cudaSuccess
FAIL tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G
FAIL tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess
vmm_test: FAILED
```

Kernel log at the same moment:

```
amdgpu 0000:09:00.0: [gfxhub] page fault (src_id:0 ring:40 vmid:8 pasid:2450)
amdgpu 0000:09:00.0: GCVM_L2_PROTECTION_FAULT_STATUS:0x00841051
amdgpu 0000:09:00.0:          Faulty UTCL2 client ID: TCP (0x8)
amdgpu 0000:09:00.0:          MORE_FAULTS: 0x1
amdgpu 0000:09:00.0:          PERMISSION_FAULTS: 0x5
```

There was no GPU reset, and the cards kept working afterwards.

Three runs per card, with nothing else on the GPUs:

| card (`HIP_VISIBLE_DEVICES`) | PCI | passed |
|---|---|---|
| 0 | 0000:03:00.0 | 1 of 3 |
| 1 | 0000:09:00.0 | 0 of 3 |
| 2 | 0000:0f:00.0 | 0 of 3 |

Page faults appear in the kernel log for all three cards. With `STRATA_HIP_VMM=0` the test prints `no CUDA device with virtual memory - skipped`. I have not run `--kv-grow` with a model, so I do not know how this shows up in real use.

I have not found whether the cause is in Strata's mapping/access sequence or in ROCm 7.14 on gfx1011. The permission-fault status suggests the remapped chunks are mapped without read/write access, but that is only my reading of the status bits.

### Smallest steps or command that reproduce it

```
cmake -G Ninja -S . -B build-hip -DCMAKE_BUILD_TYPE=Release -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF \
  -DSTRATA_BUILD_TESTS=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1011
cmake --build build-hip
for d in 0 1 2; do for r in 1 2 3; do HIP_VISIBLE_DEVICES=$d build-hip/vmm_test; done; done
```

### Other ctest results on this build (same machine, for completeness)

`ctest` per card, excluding `ple_parity` and `platform_memory_test`: 133 tests. Card 0: 5 failed; cards 1 and 2: 6 failed (the extra one is `vmm_test`). The others are the same on all three cards and look like test or build issues, not the card:

- `native_hadamard_test`, `tq_cpu_parity`, `tq_cpu_parity_avx2` print `this ggml has no trellis types (build with STRATA_GGML_DIR = the agention fork)` and exit as failed rather than skipped.
- `hip_prefill_mmq_parity`: the Q2_0 products pass (`rel_l2` ≈ 0.0005–0.0013), then `MMQ parity screen failed: basic_string: construction from null is not valid`. From the source, the trellis loop sees `supported(145)` as true, while this ggml has no name for type 145, so `std::string tn = ggml_type_name(t)` gets null. This is from reading the code; I did not debug it further.
- `isa_guard_check`: `VIOLATION strata::kernels::cpu::tq_detail::tq_rows_avx512(...) (530 wide instructions, first vmovdqa32): wide-ISA code outside the wide-ISA translation unit`.

Skipped as expected for RDNA1 / this host: the WMMA, hipBLASLt, rocBLAS and fused prefill tests, `qsa_scores_gfx1100_parity`, `expert_cache_segmented_test`, `expert_multi_test`.

### Logs

<details><summary>vmm_test, 3 runs per card (vmm.out)</summary>

```
dev 0:
  run 1: vmm_test: ok|
  run 2: :0:rocdevice.cpp            :3608: 50677180454 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x79cb65240000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
  run 3: :0:rocdevice.cpp            :3608: 50679348459 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x7c9e36e01000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
dev 1:
  run 1: :0:rocdevice.cpp            :3608: 50681642604 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x7f93d5017000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
  run 2: :0:rocdevice.cpp            :3608: 50683826731 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x77dfff027000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
  run 3: :0:rocdevice.cpp            :3608: 50686011166 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x7cfb1360b000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
dev 2:
  run 1: :0:rocdevice.cpp            :3608: 50688301723 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x7a58c9215000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
  run 2: :0:rocdevice.cpp            :3608: 50690483507 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x708a01e05000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
  run 3: :0:rocdevice.cpp            :3608: 50692662099 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x741bbe80e000, kernel: __|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == c|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G|FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess|vmm_test: FAILED|
PCI bus of HIP_VISIBLE_DEVICES=0,1,2 (hipDeviceGetPCIBusId): 0000:03:00.0 0000:09:00.0 0000:0f:00.0 
"page fault" lines in the kernel log during these runs:      10 0000:03:00.0      10 0000:09:00.0      10 0000:0f:00.0 
```
</details>

<details><summary>ctest output of the failed tests, card 2 (excerpt of LastTest.log)</summary>

```
19/133 Testing: vmm_test
19/133 Test: vmm_test
Output:
VGPU=0x5e8c6f370910 SWq=0x743b991ce000, HWq=0x743b88200000, id=1
	Dispatch Header =0x1502 (type=2, barrier=1, acquire=2, release=2), setup=1
	grid=[4608, 1, 1], workgroup=[256, 1, 1]
	private_seg_size=0, group_seg_size=0
	kernel_obj=0x743b991be4c0, kernarg_address=0x0x743b80c00100
	completion_signal=0x0, correlation_id=0
	rptr=2, wptr=4
:0:rocdevice.cpp            :3608: 26887096612 us:  Memory Fault Error [host: bc160, GPU index: 0, faulting addr: 0x743b89049000, kernel: __amd_rocclr_fillBufferAligned]
FAIL /home/codex/strata020/tests/core/vmm_test.cpp:57: cudaMemset(a.base() + 4 * G, 7, 3 * G) == cudaSuccess && cudaDeviceSynchronize() == cudaSuccess
FAIL /home/codex/strata020/tests/core/vmm_test.cpp:67: f1 >= f0 + 8 * G
FAIL /home/codex/strata020/tests/core/vmm_test.cpp:68: cudaGetLastError() == cudaSuccess
vmm_test: FAILED
Test time =   2.18 sec
Test Failed.

119/133 Testing: hip_prefill_mmq_parity
119/133 Test: hip_prefill_mmq_parity
Output:
synthetic-Q2_0-GU-pass0: type=42 rows=7 experts=3 ref_rms=1.98041 rel_l2=0.00106754 max_abs/ref_rms=0.00380476 zero_row_max=0
synthetic-Q2_0-down-pass0: type=42 rows=7 experts=3 ref_rms=2.10697 rel_l2=0.000617685 max_abs/ref_rms=0.00253579 zero_row_max=0
synthetic-Q2_0-GU-pass1: type=42 rows=7 experts=3 ref_rms=1.98512 rel_l2=0.0011973 max_abs/ref_rms=0.00416649 zero_row_max=0
synthetic-Q2_0-down-pass1: type=42 rows=7 experts=3 ref_rms=2.12958 rel_l2=0.000538054 max_abs/ref_rms=0.00216068 zero_row_max=0
synthetic-Q2_0-GU-pass2: type=42 rows=5 experts=2 ref_rms=1.90121 rel_l2=0.00125928 max_abs/ref_rms=0.00443003 zero_row_max=0
synthetic-Q2_0-down-pass2: type=42 rows=5 experts=2 ref_rms=2.0192 rel_l2=0.000491595 max_abs/ref_rms=0.00196968 zero_row_max=0
synthetic-Q4_K skipped: not built (STRATA_MMQ_KQUANTS off?)
synthetic-Q5_K skipped: not built (STRATA_MMQ_KQUANTS off?)
synthetic-Q6_K skipped: not built (STRATA_MMQ_KQUANTS off?)
synthetic-Q5_1 skipped: not built (STRATA_MMQ_KQUANTS off?)
MMQ parity screen failed: basic_string: construction from null is not valid
Test time =   0.98 sec
Test Failed.

124/133 Testing: native_hadamard_test
124/133 Test: native_hadamard_test
Output:
this ggml has no trellis types (build with STRATA_GGML_DIR = the agention fork)
FWHT vs dense normalized Sylvester matrix: max abs diff 3.35e-07  ok
Test time =   0.00 sec
Test Failed.

125/133 Testing: tq_cpu_parity
125/133 Test: tq_cpu_parity
Output:
this ggml has no trellis types (build with STRATA_GGML_DIR = the agention llama.cpp fork)
Test time =   0.00 sec
Test Failed.

126/133 Testing: tq_cpu_parity_avx2
126/133 Test: tq_cpu_parity_avx2
Output:
this ggml has no trellis types (build with STRATA_GGML_DIR = the agention llama.cpp fork)
Test time =   0.01 sec
Test Failed.

131/133 Testing: isa_guard_check
131/133 Test: isa_guard_check
Output:
check_isa_guard: /home/codex/strata020/build-hip/libstrata_kernels_cpu.a: 477 functions, 69 hold wide-ISA instructions, 1 not allowed
  VIOLATION strata::kernels::cpu::tq_detail::tq_rows_avx512(int, unsigned char const*, unsigned long, int, strata::kernels::cpu::tq_detail::TqAct const* const*, int, float* (530 wide instructions, first vmovdqa32): wide-ISA code outside the wide-ISA translation units (object tq_avx512.cpp.o) [tq_avx512.cpp.o]
Test time =   1.37 sec
Test Failed.

```
</details>

Mehr auf der Site

Links zu Install, Modellen, Releases.