[sbates130272]

What this does

Adds --data_validation support for ROCm device memory, which today is refused for --use_rocm.

One commit, 10 files, +1837/−97. Four new files (rocm_validation.{c,h}, rocm_kernels.{h,hip}); the rest is build glue plus the changes described under Behaviour changes below.

Depends on nothing unmerged. Builds clean with --enable-rocm and hipcc, with --enable-rocm without hipcc, and with neither.

Why

perftest validates host memory and CUDA memory. ROCm has neither a validation backend nor working copy ops, so --data_validation is rejected at parameter-parse time for --use_rocm. That gap matters most exactly where a bandwidth number is least trustworthy: a wrong address translation moves the wrong bytes and the throughput figure looks fine.

The ROCm memory backend also set all three copy ops to plain memcpy into and out of device memory. Those are now hipMemcpy with the right direction.

How it works

Two paths, mirroring the CUDA backend's shape.

Live (default). src/rocm_kernels.hip is a port of src/cuda_kernels.cu: a resident kernel where block 0 monitors the tail markers and blocks 1..N drain a lock-free queue, so each cycle is checked as it lands.

Host-staged (fallback). Where the live kernel cannot run — a build without hipcc, a device the kernel will not launch on, or PERFTEST_ROCM_VALIDATION=staged — the buffer is copied back and compared on the host. This is strictly weaker: it sees only each chunk's final state, so transient corruption that a later cycle overwrote is missed and bytes_validated is correspondingly lower. Which path ran is printed at stop, and a live validator that observed no completed cycle falls back rather than reporting an empty result.

Both reuse upstream's classify_validation_mismatch(), so a run stopping mid-cycle is scored as an overwrite race rather than as corruption, consistently with the other backends.

The port, and where it diverges from CUDA

Deliberately close to the original: same monitor/validator split, same queue, same structures from cuda_validation.h, same exported entry points. Divergences are marked PORT: at the site. The ones worth a reviewer's time:

Warp width — a correctness fix, not a tuning choice. The CUDA reduction hardcodes a 32-lane warp (tid & 31, shuffle offset from 16). An AMD wavefront is 64 lanes on CDNA and 32 on RDNA, so the literal reduces only within half a wave and lets two lanes observe lane == 0, double-counting the atomics that follow. The port uses warpSize, which is a built-in rather than a constant, so the #pragma unrolls go with it.

Scoped atomics instead of volatile. The CUDA code polls through *(volatile T *)p and orders with __threadfence_system(). volatile is not a synchronisation primitive — it constrains the compiler for one access and says nothing about which agents it must be coherent with; it works on NVIDIA because the emitted loads happen to bypass cache. On AMD the memory scope selects which cache levels are bypassed, so this uses __hip_atomic_load/store with it stated: system scope for the markers and slots the NIC DMAs into and for flags shared with the host, agent scope for the queue only this kernel touches. A uniform __threadfence_system() over-synchronised the queue while under-specifying the markers.

Two __ldg reads replaced, both on data the NIC is actively writing. __ldg is a read-only-cache hint asserting the data will not change for the kernel's life, which is false for a DMA target. The tail retry re-read through __ldg was the load least likely to observe the write it was waiting for; and the forward-byte snapshot feeds classify_validation_mismatch(), where a stale read biases toward reporting corruption that did not occur. The bulk compare loop keeps __ldg, ordered by the acquire on the marker.

Poll loops back off with s_sleep. A device that runs a wave until it blocks never retires a tight poll loop, so the kernel cannot observe should_stop. validation_start() also honours num_blocks/threads_per_block, which the CUDA original accepts and (void)s — 33 blocks × 1024 spinning threads is not a geometry every device retires.

rocm_kernels_available() is the one addition with no CUDA counterpart: the CUDA plugin answers that question by whether libperftest_kernels.so dlopen'd, while these kernels link into libperftest.a and need a runtime probe.

Behaviour changes outside the ROCm backend — please look here

data_validation_stop_and_report() is reordered: errors, then nothing-validated, then pass. This affects every backend.

Testing "nothing validated" first reports a run whose chunks all mismatched (errors > 0 with bytes_validated == 0) as a warning about missing data, hiding a real failure. And backends set passed = (errors == 0 && chunks > 0), so a run that observed no completed cycle also has passed == 0 and would otherwise print a failure with every error field still at its initialised zero:

VALIDATION: FAILED - 0 error(s), first at QP 0, Chunk 0, Offset 0 (expected 0x00, got 0x00)

bytes_validated == 0 is that same condition counted differently, since a chunk mismatching at its first byte contributes no compared bytes but still increments the chunk count.

The ROCm dma-buf path no longer greps the kernel config before running. The old check required CONFIG_DMABUF_MOVE_NOTIFY, which no longer exists upstream — dma-buf always builds with it now, and .move_notify became .invalidate_mappings — so the probe began refusing to run because support had become universal. What it gated was whether amdgpu_dma_buf_pin() kept AMDGPU_GEM_DOMAIN_VRAM or fell back to GTT: driver policy, versioned with the driver rather than the running kernel, which on a DKMS install moves independently of it. CONFIG_PCI_P2PDMA fared no better, being blind to whether the NIC's driver implements .reg_user_mr_dmabuf. Both questions are already answered accurately by hsa_amd_portable_export_dmabuf() and ibv_reg_dmabuf_mr(); their diagnostics are improved instead of being predicted. This also removes PERFTEST_ASSUME_DMABUF, which existed only to defeat that check.

VALIDATION_INLINE now also annotates __host__ __device__ for __HIPCC__. hipcc compiles .hip as nvcc compiles .cu; both need it for the shared helpers to be callable from device code. No change to the CUDA path.

configure learns --with-rocm-arch, defaulting to rocm_agent_enumerator. hipcc otherwise targets whatever GPU the build host reports, which is wrong whenever perftest is cross-built or built in a VM — and the symptom is "device kernel image is invalid" at launch rather than anything at build time.

The --data_validation rejection message now mentions --use_rocm.

Environment variables

All default off; none changes behaviour unless set.

variable	effect
PERFTEST_ROCM_VALIDATION=staged	force the host-staged compare
PERFTEST_ROCM_VALIDATION_STAGED=1	same
PERFTEST_ROCM_VALIDATION_BLOCKS/_THREADS	override the live kernel's launch geometry
PERFTEST_ROCM_VALIDATION_FAULT=n	corrupt n bytes of the staged copy before comparing — a validator that only ever passes is indistinguishable from one that does not look
PERFTEST_ROCM_HOST_ALLOC=1	allocate host-mapped rather than VRAM, for devices whose MR registers cleanly and then fails to gather. Measures host memory, so it diagnoses reachability rather than benchmarking peer-to-peer

Testing, and its limits

Tested on an emulated gfx1250 behind an emulated RDMA NIC, two guests with real RDMA traffic between them. The live kernel retires, the host-staged fallback passes with real byte counts, and injected faults are caught on both paths — so a PASSED here means the checker was looking, not merely silent.

What that does not cover: no discrete AMD GPU, and no performance figures (the device has no real compute, so throughput numbers from it would be meaningless). Reviewers with Instinct hardware would be very welcome to run it.

Build-tested in containers across --enable-rocm with hipcc, --enable-rocm without hipcc, and a plain build — all warning-free.