fix(clr): let the fat-binary readable bound span a split mapping

[magaonka-amd]

Motivation

hipModuleLoadData rejects valid code objects with hipErrorInvalidImage. The
readable-size bound is measured to the end of one /proc/self/maps entry, but one
allocation can be split across several, so the bound comes back short.

How it shows up in JAX

Random kernels fail to load part-way through a run, with no pattern to which ones.
A 16-process JAX test suite hits it within about 30 minutes. The same run also
fails inside MIOpen and in jaxlib's own FFI kernels, because all three load code
objects the same way.

How to reproduce

Build a code object for your GPU:

echo '__global__ void k(float* p) { *p = 1.0f; }' > kernel.hip
hipcc --genco -x hip kernel.hip --offload-arch=$(rocminfo | grep -om1 'gfx[0-9a-z]*') -o kernel.co

repro.cpp:

#include <hip/hip_runtime.h>
#include <malloc.h>
#include <numaif.h>
#include <unistd.h>

#include <cstdio>
#include <cstring>
#include <fstream>
#include <vector>

int main(int argc, char** argv) {
  std::ifstream f(argv[1], std::ios::binary | std::ios::ate);
  std::vector<char> co(f.tellg());
  f.seekg(0);
  f.read(co.data(), co.size());

  const size_t page = sysconf(_SC_PAGESIZE);
  mallopt(M_MMAP_THRESHOLD, 1 << 30);  // keep the buffer on the brk heap
  char* buf = static_cast<char*>(malloc(64 * page));
  memset(buf, 0, 64 * page);
  hipFree(nullptr);

  // Put the code object across a page boundary inside the allocation.
  const uintptr_t seam = (reinterpret_cast<uintptr_t>(buf) + 16 * page) & ~(page - 1);
  char* image = reinterpret_cast<char*>(seam - 2048);
  memcpy(image, co.data(), co.size());

  // Split [heap] the way libhsakmt does when it places memory on a NUMA node.
  unsigned long mask = 1;
  if (mbind(reinterpret_cast<void*>(seam), 8 * page, MPOL_PREFERRED, &mask, sizeof(mask) * 8, 0)) {
    printf("mbind blocked (seccomp?); cannot reproduce here\n");
    return 2;
  }

  // The image never moved and every byte is readable.
  volatile char sink = 0;
  for (size_t i = 0; i < co.size(); ++i) sink = image[i];
  printf("image %p, %zu bytes, all readable, memcmp %s\n", image, co.size(),
         memcmp(image, co.data(), co.size()) ? "FAILED" : "ok");

  hipModule_t m = nullptr;
  const hipError_t e = hipModuleLoadData(&m, image);
  printf("hipModuleLoadData -> %d (%s)\n", static_cast<int>(e), hipGetErrorString(e));
  return e != hipSuccess;
}

hipcc -O2 repro.cpp -o repro -lnuma
AMD_LOG_LEVEL=1 ./repro kernel.co

On develop today:

image 0x55c080759800, 9696 bytes, all readable, memcmp ok
:1:hip_fatbin.cpp :417 : Rejecting fat binary: code object for isa
  'amdgcn-amd-amdhsa--gfx942:sramecc+:xnack-' is out of bounds
  (offset=4096 size=5600 image bound=2048)
hipModuleLoadData -> 200 (device kernel image is invalid)

The image is 9696 bytes and every one of them was just read. The bound says 2048.

This needs mbind to be permitted. Docker's default seccomp blocks it, which
libhsakmt already works around at fmm.c:2037.

Why the split happens

The kernel splits a mapping when part of it gets an attribute the rest does not
have, and the pieces do not merge back afterwards. A NUMA memory policy is one such
attribute, and it shows up in neither the /proc/self/maps line nor smaps
VmFlags, so the pieces look identical:

before:  562668960000-562669671000 rw-p 00000000 00:00 0   [heap]

after:   562668960000-562668a05000 rw-p 00000000 00:00 0   [heap]
         562668a05000-562668a0d000 rw-p 00000000 00:00 0   [heap]
         562668a0d000-562669671000 rw-p 00000000 00:00 0   [heap]

This is not something the application has to do. libhsakmt does it itself, in
bind_mem_to_numa() (rocr-runtime/libhsakmt/src/fmm.c:2032).

An strace of a twenty-line JAX program shows 1290 mbind(..., MPOL_PREFERRED, ...) calls, plus
52039 mprotect and 124 madvise(MADV_DONTFORK | MADV_HUGEPAGE), all of which
split mappings too.

What the fix does

FindFileNameFromAddress walks forward from the entry holding the image while the
next one is directly adjacent, readable, and backed by the same object (same
device, inode, path), and reports the end of that run.

Requiring the same backing object keeps the bound inside the allocation instead of
letting it walk into whatever is mapped next. The walk stops after 32 entries so
its cost does not grow with the number of mappings in the process.

Issue Tracking

JIRA ID: ROCM-29159

Test Plan

New test, OOB_hipModuleLoadData_Positive_ImageSpansSplitMapping in
catch/unit/oob/oob_module.cc. It compiles a code object with hiprtc so the arch
is right, puts it across a split anonymous mapping, checks every byte is readable,
and requires the load to succeed. It fails on develop today.

Test Result

Suite	unmodified	patched
oob_module, 6 existing cases, 200 assertions	pass	pass
Unit_hipModuleLoad_Negative_MalformedFatBinaryBounds	pass	pass
ModuleTest, full run, 5,850,405 assertions	9 pre-existing failures	same 9, byte-identical
OOB_hipModuleLoadData_Positive_ImageSpansSplitMapping (new)	fail	pass

Latency, 2000 hipModuleLoadData iterations:

mappings in the process	unmodified	patched
219	2466 us	2448 us
10218	2004 us	2017 us
44218	2455 us	2450 us

End to end: the JAX UT passes with no flakes with this fix.

Not tested: Windows. I'm not sure about windows usecase, I'm open to suggestions.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/clr/rocclr/os/os.hpp, projects/clr/rocclr/os/os_posix.cpp, projects/hip-tests/catch/unit/oob/oob_module.cc; no test file found
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?