// Minimal HIP smoke test: is the host ROCm usable from a plain process?
#include <cstdio>
#include <cstring>
#include <hip/hip_runtime.h>

int main(void) {
    int n = -1;
    hipError_t e = hipGetDeviceCount(&n);
    printf("hipGetDeviceCount -> %d (%s), count=%d\n",
           (int)e, hipGetErrorString(e), n);
    if (e != hipSuccess || n <= 0) return 1;
    setvbuf(stdout, NULL, _IONBF, 0);   /* an earlier crash lost the buffer, hiding the crash point */
    // Device properties are deliberately NOT queried here: hipGetDeviceProperties segfaults on the
    // ROCm build this probe was written for, even though hipGetDeviceCount above works fine. Card
    // identity comes from sysfs anyway — scripts/pcie-probe.sh prints the link topology.

    const size_t bytes = 2459136;   /* 2.344 MiB: the engine's expert slab */
    void *dev = NULL, *host = NULL;
    if (hipMalloc(&dev, bytes) != hipSuccess) { printf("hipMalloc failed\n"); return 1; }
    if (hipHostMalloc(&host, bytes, hipHostMallocDefault) != hipSuccess) {
        printf("hipHostMalloc failed\n"); return 1;
    }
    memset(host, 0x5a, bytes);

    hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b);
    hipDeviceSynchronize();
    hipEventRecord(a, 0);
    const int iters = 200;
    for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice);
    hipEventRecord(b, 0); hipEventSynchronize(b);
    float ms = 0; hipEventElapsedTime(&ms, a, b);
    printf("\n  %.3f MiB x %d iters in %.2f ms -> %.2f GB/s H2D (pinned)\n",
           bytes/1048576.0, iters, ms, (double)bytes*iters/(ms/1000.0)/1e9);
    printf("  Gen5 x8 ceiling = 31.5 GB/s\n");
    return 0;
}