Check if it works on AMD

[Qubitol]

This implies:

• MadMatrix (standalone and not)
• MadSpace (including the automatic installation)

[theoheimel]

I ran some tests on the lemaitre4 cluster with MI300A APUs. Everything compiled without issues but crashed with illegal memory access errors. There appear to be issues both in madspace and madmatrix. I let claude work on it for a few hours. I'd say the findings are somewhat inconclusive and there was no simple fix. Before we continue, we should first check on other machines with AMD GPUs if we have similar issues or if this is specific to that system. At the LUMI hackathon, the full madspace+madmatrix pipeline was working, but I didn't test it since then. Here is what the clanker has to say:

AMD/HIP GPU illegal-memory-access bugs on MI300A (ROCm 7.2.1, gfx942)

Environment

AMD Instinct MI300A APU (gfx942), ROCm 7.2.1 (HIP version: 7.2.53211), unified
memory, XNACK disabled. All three bugs below are HIP-only — CPU and CUDA builds of
the same code are unaffected and are the trusted reference throughout. All fixes are
gated behind __HIPCC__ / the HIP branch of the abstraction headers, so CUDA is
untouched.

Summary

Running GPU event generation on this hardware surfaced three independent bugs, all
presenting as the same symptom (hipErrorIllegalAddress / "an illegal memory access
was encountered", sometimes HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION):

#	Where	Status	Cause
1	generated calculate_jamps kernel (matrix element)	Fixed	HIP/clang codegen bug, optimization-level dependent
2	hipMallocAsync/hipFreeAsync (stream-ordered pool allocator)	Fixed	Known ROCm 7.x regression
3	madspace::gpu::GpuTensorView::operator[] (kernel_cut_all)	Open	Suspected AMDGPU kernarg address-space codegen bug

Fixing only one of Bugs 1/2 was not enough for even the simplest test process — both
were required simultaneously (bisection table below). Bug 3 was found while stress
testing a more complex process and is still open; this issue is to get more eyes on it.

---

Bug 1 — compiler codegen bug in generated calculate_jamps kernel (FIXED)

Repro: g g > t t~, any optimization level (-O1, -O2, or -O3 — not just
-O3) applied to the calculate_jamps __global__ kernel makes hipcc/clang emit a
momenta-buffer load with a corrupted address. -O0 avoids it.

Same class of bug as the long-standing, never-root-caused
madgraph4gpu#806 (HIP-only
"Memory access fault", -O3 broke it, -O2 was the historical workaround, root cause
never found, still open) — but worse here, since even -O1 triggers it.

Fix: mark only the calculate_jamps kernel __attribute__((optnone)), HIP-only, in
the template that generates CPPProcess.cc (madmatrix/model_handling.py, right before
the kernel signature is emitted):

#ifdef __HIPCC__
  __global__ void __attribute__((optnone)) /* clang-format off */
#else
  __global__ void /* clang-format off */
#endif
  calculate_jamps( int ihel, ...

Scoping the fix to just this kernel (rather than -O0 for the whole translation unit)
measured ~2x faster for the gg > ttx test process, since coupling computation, color
sum, helicity selection, etc. in the same file stay fully optimized.

---

Bug 2 — ROCm 7.x stream-ordered pool allocator regression (FIXED)

Repro: hipMallocAsync/hipFreeAsync under allocation churn (one alloc+free per
matrix-element batch call) independently produced the same illegal-memory-access
symptom, confirmed via rocgdb + AMD_LOG_LEVEL=3 to be unrelated to Bug 1 (fault
address falls within the nominally allocated buffer when Bug 1 is fixed and this one
isn't).

External corroboration this is a real, known ROCm 7.x issue:

• AMD's own rocBLAS stopped using stream-order allocation by default from ROCm 7.1.1
  onwards (needs ROCBLAS_STREAM_ORDER_ALLOC=1 to re-enable).
• ROCm/ROCm#6178: hipMallocAsync /
  hipMallocFromPoolAsync fail on gfx1100/gfx1101 with ROCm 7.x (6.4.x fine on the same
  hardware); open, unresolved.
• HIP docs still mark the Stream Ordered Memory Allocator API as Beta.

Fix: at the GPU abstraction layer, HIP-only, fall back to the plain non-pooled
allocator (CUDA keeps the real stream-ordered API). Since hipFreeAsync's contract
allows freeing while the stream still has in-flight work referencing the pointer, the
synchronous hipFree replacement needs an explicit stream drain first:

#define gpuMallocAsync( ptr, size, stream ) checkGpu( hipMalloc( ptr, size ) )
#define gpuFreeAsync( ptr, stream ) do { checkGpu( hipStreamSynchronize( stream ) ); checkGpu( hipFree( ptr ) ); } while( 0 )

Applied in both madgraph/iolibs/template_files/madmatrix/GpuAbstraction.h (HIP branch)
and madspace/src/gpu/gpu_abstraction.cuh (as hipMallocAsyncWorkaround/
hipFreeAsyncWorkaround wrappers).

Bisection: why both fixes are needed

g g > t t~, HIP_LAUNCH_BLOCKING=1 bin/generate_events -f, repeated several times per
configuration (a single pass is not reliable — see caveat below):

optnone fix	allocator fix	result
no	no	crash, ~100%
no	yes	crash — fault confirmed (via rocgdb/AMD_LOG_LEVEL) to be a codegen problem, not an allocator problem
yes	no	crash, 8/8
yes	yes	pass, 10/10, with and without HIP_LAUNCH_BLOCKING=1

Caveat: an early "confirmed fix" (allocator fix alone) turned out to be a false
negative — that run happened to use a smaller batch size and got lucky. Both bugs are
intermittent/data-size-dependent in whether they visibly crash; always verify with
several repeated runs, not one.

---

Bug 3 — suspected AMDGPU kernarg address-space codegen bug in GpuTensorView (OPEN)

Repro: g g > t t~ g g (2→4, devices = ["hip"]), crashes during the very first
survey batch — before MadNIS training even starts — always inside
madspace::gpu::kernels::kernel_cut_all (madspace/src/kernels/cuts.hpp:34-40), the
phase-space cuts kernel. Confirmed not a logic/shape bug and not
optimization-level dependent (details below) — this looks like a genuine AMDGPU backend
bug and would benefit from more eyes / an upstream report.

template <typename T>
KERNELSPEC void
kernel_cut_all(FIn<T, 1> obs, FIn<T, 1> min, FIn<T, 1> max, FOut<T, 0> w) {
    FVal<T> cut = 1.;
    for (std::size_t i = 0; i < obs.size(); ++i) {
        cut = where((obs[i] < min[i]) | (obs[i] > max[i]), 0., cut);
    }
    w = cut;
}

Root cause, pinned down via rocgdb

rocgdb -x with set amdgpu precise-memory on gives the exact fault location and the
live argument values at the crash:

Thread 8 received signal SIGBUS / HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION
#0 operator double (this=0x20000000001e8) at .../gpu/tensor.cuh:87   [return *_data;]
#1 kernel_cut_all<GpuTypes>(obs=..., min=..., max=..., w=...) at .../kernels/cuts.hpp:37
#2 run_kernel<kernel_cut_all,...>(batch_size=1000, ...) at .../gpu/tensor.cuh:195

(gdb) print obs
$ = {_data = 0x15475c544800, _stride = 0x2000000000028, _shape = 0x2000000000038}
(gdb) print min
$ = {_data = 0x155501474000, _stride = 0x2000000000050, _shape = 0x2000000000060}
(gdb) print max
$ = {_data = 0x155501475000, _stride = 0x2000000000078, _shape = 0x2000000000088}

Every _data pointer is a legitimate device address. Every _stride/_shape pointer,
on all three arguments, shares the exact same bogus high bits 0x200000000, with only a
small low offset varying. That's not a real pointer — it looks like a kernel-argument
segment offset with a missing/garbage base, i.e. a botched kernarg → flat/generic
address-space cast.

Mechanism: GpuTensorView::operator[] (madspace/src/gpu/tensor.cuh:52-62) reduces
a tensor view by one rank via

return {&_data[index * _stride[0]], &_stride[1], &_shape[1]};

i.e. it returns a view whose _stride/_shape are addresses of elements inside the
original view's own _stride/_shape arrays. For the first slice performed by
run_kernel (views[index], reducing the packed, by-value kernel argument down by the
batch dimension), *this is the raw kernel argument itself — for a packed=true
GpuTensorView, SizeType = std::size_t[_dim], so the shape/stride arrays are stored
inline in the by-value argument, which lives in the kernarg (constant, address-space-4)
segment. Taking &_stride[1] there requires correctly folding the kernarg base into the
resulting pointer; on this toolchain that conversion produces garbage high bits instead.
Any later dereference (min[i], max[i], or a second-level obs[i]) either reads
garbage silently or, when the address happens to fall outside any mapped aperture,
faults.

This also explains why two back-to-back, byte-identical kernel_cut_all launches were
seen where the first "succeeded" (hipSuccess) and only the second faulted — the bad
pointer is present on every call, but whether it actually faults depends on whether that
particular garbage address happens to be unmapped. The "successful" call was likely
silently reading wrong min/max values, not genuinely correct — treat any
HIP-backend output from this code path as untrustworthy until this is fixed, even from
runs that didn't crash.

What's been ruled out

• Not a logic/shape bug: running the identical process/cuts on the CPU backend
  (devices = ["cppnone"]) completes the survey cleanly. The cut_all opcode,
  Cuts::build_function_impl, and kernel_cut_all itself are all correct.
• Not optimization-level dependent: madspace's CMake Debug build actually compiles
  HIP at -O1 (bare -O in the ninja flags), not -O0 as the name suggests. Forcing a
  real -O0 via CMAKE_HIP_FLAGS_DEBUG still crashed.
• __attribute__((optnone)) on operator[] (mirroring the Bug 1 fix): made it
  worse — went from intermittent to 100% reproducible, faulting on the very first call.
• __attribute__((always_inline)) on operator[] (the opposite bet): also 100%
  reproducible failure. So this isn't simple inlining-heuristic sensitivity either — both
  directions make it deterministic rather than fixing it; only the compiler's own default
  inlining choice was ever observed to succeed at all (intermittently).
• Standalone minimal reproduction attempts: two hand-written HIP programs replicating
  the exact GpuTensorView/run_kernel<auto func, dims, V...> template pattern (one
  simplified, one a near-verbatim copy) both ran correctly for 50+ iterations outside the
  full madspace build. So the trigger depends on something in the fuller call chain not
  yet isolated — plausibly the std::apply/tuple/lambda machinery tensor_foreach goes
  through before reaching run_kernel, the real MemPool-backed device memory, or a
  non-default stream. A clean, fully standalone reproducer is still needed — this is
  probably the most valuable next contribution.

Suggested fix direction (not yet implemented)

Stop ever forming a pointer into the by-value kernel argument's inline _stride/_shape
arrays. When GpuTensorView::operator[] slices a packed=true view (only the first
reduction, directly off a raw kernel argument), it should copy the remaining
_stride/_shape values into the returned object's own inline storage, instead of
returning pointers into the parent — so every subsequent slice addresses normal
stack/register-resident memory, not the kernarg segment. Deeper slices (already off a
local, non-kernarg object) can keep the current pointer-aliasing scheme.

Not yet implemented because: it touches core indexing machinery used by essentially
every GPU kernel in madspace, not just cut_all/cut_any; it introduces a
self-referential struct (an object holding a pointer to its own member) which is only
safe if C++17 guaranteed copy elision reliably applies on this toolchain — something we
have direct evidence not to assume here; and it needs a full regression pass across
existing GPU processes to land safely.

How to help

1. Isolate a fully standalone reproducer (no madspace/madgraph7 dependency) —
   probably needs to replicate more of tensor_foreach's std::apply/lambda launch path,
   not just run_kernel/GpuTensorView in isolation.
2. Implement and test the copy-on-first-slice fix above; validate with repeated runs (not
   a single pass) plus a regression pass on other GPU processes.
3. Once confirmed, consider filing upstream with AMD/ROCm — this looks like a genuine
   backend bug in kernarg address-space legalization, independent of madspace.
4. Worth checking whether a different ROCm version (older or newer than 7.2.1) is
   affected, to help scope/report it upstream.

---

Debugging notes

• rocgdb -x <cmdfile> needs set amdgpu precise-memory on for an accurate fault
  PC/address — without it, the reported location is just wherever the wave next hit
  s_waitcnt, not the real faulting instruction.
• AMD_LOG_LEVEL=3 gives a full HIP API trace including ShaderName per kernel launch
  and allocation return addresses — useful for cross-referencing a fault address against
  real allocations (out-of-bounds vs. compiler-generated-garbage-address).
• dmesg shows the underlying amdgpu VM_L2_PROTECTION_FAULT entries, confirming a
  real hardware page fault independent of any HIP-level software check.
• nm -D lib*.so | grep hipMallocAsync is a quick way to check whether a built library
  still references the pooled allocator or picked up the Bug 2 workaround.
• Rebuild madspace via python3 madspace/install.py (not raw cmake) so build products
  land under madspace/install/, which the run scripts add to sys.path automatically.

[stloufra]

I tried on the NGT cluster with W7900.
Standalone was compiling with a wrong arch, fixed that nad madmatrix seems to work (same for Theo's try).
With MadSpace right at the survey step it fails with Memory access fault. This is prob in madspace's own GPU sampling kernels.