[DO NOT MERGE] H100 CAGRA hang: JIT cache, module loading, and library pin experiments

ci/b3_cache_info.py

@@ -0,0 +1,45 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# DO NOT MERGE: summarize the CUDA driver JIT cache and the cupy kernel cache.
+
+import collections
+import os
+import re
+import sys
+
+for root in sys.argv[1:]:
+    root = os.path.expanduser(root)
+    files = []
+    for d, _, fs in os.walk(root):
+        for f in fs:
+            p = os.path.join(d, f)
+            try:
+                files.append((p, os.path.getsize(p)))
+            except OSError:
+                pass
+    print(
+        f"== {root}: {len(files)} files, {sum(s for _, s in files) / 1e6:.1f} MB"
+    )
+    names = collections.Counter()
+    for p, _ in files:
+        try:
+            data = open(p, "rb").read()
+        except OSError:
+            continue
+        found = set(
+            m.decode(errors="replace")
+            for m in re.findall(rb"_Z[0-9A-Za-z_]{12,}", data)
+        )
+        for n in found:
+            names[n[:160]] += 1
+    print(f"   distinct mangled symbols: {len(names)}")
+    # Group by a short prefix (namespace/function) to see which libraries the modules come from.
+    groups = collections.Counter()
+    for n in names:
+        m = re.findall(r"(?:N|_Z)\d+([A-Za-z_]\w*)", n)
+        groups["::".join(m[:3]) if m else n[:40]] += 1
+    for g, c in groups.most_common(40):
+        print(f"   {c:5d} {g}")
+    for n in sorted(names)[:15]:
+        print(f"   e.g. {n}")

ci/b3_ltolink.py

@@ -0,0 +1,98 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# DO NOT MERGE: NVRTC -> LTO-IR -> nvJitLink(-lto -arch=sm_XX) -> GetLinkedCubin,
+# then load the image into the driver (no kernel launch) and watch CUDA_CACHE_PATH.
+import ctypes
+import os
+import sys
+
+
+def list_cache(tag):
+    p = os.environ.get("CUDA_CACHE_PATH", "")
+    files = []
+    for d, _, fs in os.walk(p):
+        files += [os.path.join(d, f) for f in fs]
+    total = sum(os.path.getsize(f) for f in files)
+    print(f"[{tag}] CUDA_CACHE_PATH={p}: {len(files)} files, {total} bytes")
+
+
+env = sys.argv[1]
+arch = sys.argv[2]  # e.g. 90
+out = sys.argv[3]
+
+nvrtc = ctypes.CDLL(f"{env}/lib/libnvrtc.so.12")
+jl = ctypes.CDLL(f"{env}/lib/libnvJitLink.so.12")
+
+src = b"""
+extern "C" __global__ void spin_or_add(int* flag, float* x, int n) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n) x[i] += 1.0f;
+  if (i == 0) { while (atomicAdd(flag, 0) == 0) {} }
+}
+"""
+prog = ctypes.c_void_p()
+assert (
+    nvrtc.nvrtcCreateProgram(ctypes.byref(prog), src, b"k.cu", 0, None, None)
+    == 0
+)
+opts = [f"-arch=compute_{arch}".encode(), b"-dlto"]
+arr = (ctypes.c_char_p * len(opts))(*opts)
+rc = nvrtc.nvrtcCompileProgram(prog, len(opts), arr)
+print("nvrtc compile rc", rc)
+sz = ctypes.c_size_t()
+assert nvrtc.nvrtcGetLTOIRSize(prog, ctypes.byref(sz)) == 0
+ltoir = ctypes.create_string_buffer(sz.value)
+assert nvrtc.nvrtcGetLTOIR(prog, ltoir) == 0
+print("ltoir bytes", sz.value)
+
+h = ctypes.c_void_p()
+lopts = [b"-lto", f"-arch=sm_{arch}".encode()]
+larr = (ctypes.c_char_p * len(lopts))(*lopts)
+rc = jl.nvJitLinkCreate(ctypes.byref(h), len(lopts), larr)
+print("create rc", rc)
+NVJITLINK_INPUT_LTOIR = (
+    3  # nvJitLinkInputType: NONE=0, CUBIN=1, PTX=2, LTOIR=3
+)
+rc = jl.nvJitLinkAddData(h, NVJITLINK_INPUT_LTOIR, ltoir, sz, b"frag")
+print("add rc", rc)
+rc = jl.nvJitLinkComplete(h)
+print("complete rc", rc)
+csz = ctypes.c_size_t()
+rc = jl.nvJitLinkGetLinkedCubinSize(h, ctypes.byref(csz))
+print("cubin size rc", rc, csz.value)
+cub = ctypes.create_string_buffer(csz.value)
+rc = jl.nvJitLinkGetLinkedCubin(h, cub)
+open(out, "wb").write(cub.raw)
+psz = ctypes.c_size_t()
+rc = jl.nvJitLinkGetLinkedPtxSize(h, ctypes.byref(psz))
+print("ptx size rc", rc, psz.value)
+lsz = ctypes.c_size_t()
+jl.nvJitLinkGetInfoLogSize(h, ctypes.byref(lsz))
+log = ctypes.create_string_buffer(lsz.value + 1)
+jl.nvJitLinkGetInfoLog(h, log)
+print("info log:", log.value.decode()[:500])
+data = cub.raw
+print(
+    "ELF magic",
+    data[:4],
+    "contains .target:",
+    b".target" in data,
+    "contains NVVM header:",
+    b"NVVM" in data,
+)
+
+list_cache("after nvJitLinkComplete")
+cu = ctypes.CDLL("libcuda.so.1")
+print("cuInit", cu.cuInit(0))
+dev = ctypes.c_int()
+ctx = ctypes.c_void_p()
+cu.cuDeviceGet(ctypes.byref(dev), 0)
+cu.cuDevicePrimaryCtxRetain(ctypes.byref(ctx), dev)
+cu.cuCtxSetCurrent(ctx)
+lib = ctypes.c_void_p()
+rc = cu.cuLibraryLoadData(ctypes.byref(lib), cub, None, None, 0, None, None, 0)
+print("cuLibraryLoadData rc", rc)
+mod = ctypes.c_void_p()
+print("cuLibraryGetModule rc", cu.cuLibraryGetModule(ctypes.byref(mod), lib))
+list_cache("after cuLibraryLoadData + cuLibraryGetModule")

ci/b3_nsys_query.py

@@ -0,0 +1,174 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# DO NOT MERGE: summarize an nsys sqlite export around the end of the capture
+# (the stall) for the H100 hang reproducer.
+
+import sqlite3
+import sys
+
+db = sqlite3.connect(sys.argv[1])
+db.row_factory = sqlite3.Row
+
+
+def q(sql, *args):
+    try:
+        return db.execute(sql, args).fetchall()
+    except sqlite3.Error as e:
+        print(f"  (query failed: {e})")
+        return []
+
+
+def tables():
+    return [
+        r[0] for r in q("SELECT name FROM sqlite_master WHERE type='table'")
+    ]
+
+
+def cols(t):
+    return [r[1] for r in q(f"PRAGMA table_info({t})")]
+
+
+def s(sid):
+    r = q("SELECT value FROM StringIds WHERE id=?", sid)
+    return r[0][0] if r else str(sid)
+
+
+T = tables()
+print("== tables with rows")
+t_end = 0
+for t in T:
+    n = q(f"SELECT count(*) FROM {t}")[0][0]
+    c = cols(t)
+    extra = ""
+    if n and "start" in c:
+        mn, mx = q(f"SELECT min(start), max(start) FROM {t}")[0]
+        extra = f" start {mn / 1e9:.3f}..{mx / 1e9:.3f} s"
+        if t.startswith("CUPTI_ACTIVITY_KIND_KERNEL") or t.startswith(
+            "CUPTI_ACTIVITY_KIND_RUNTIME"
+        ):
+            t_end = max(t_end, mx)
+    if n:
+        print(f"  {t}: {n}{extra}")
+
+print(f"== last CUDA activity at {t_end / 1e9:.3f} s")
+
+if "CUPTI_ACTIVITY_KIND_KERNEL" in T:
+    print("== last 25 kernels (start s, dur ms, stream, name)")
+    for r in q(
+        "SELECT start, end, streamId, demangledName, gridX, gridY, gridZ FROM CUPTI_ACTIVITY_KIND_KERNEL ORDER BY start DESC LIMIT 25"
+    )[::-1]:
+        print(
+            f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} s{r['streamId']} grid=({r['gridX']},{r['gridY']},{r['gridZ']}) {s(r['demangledName'])[:150]}"
+        )
+    print("== longest 10 kernels")
+    for r in q(
+        "SELECT start, end, streamId, demangledName FROM CUPTI_ACTIVITY_KIND_KERNEL ORDER BY end-start DESC LIMIT 10"
+    ):
+        print(
+            f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} ms s{r['streamId']} {s(r['demangledName'])[:150]}"
+        )
+
+if "CUPTI_ACTIVITY_KIND_MEMCPY" in T:
+    # CUPTI memory kinds: 0 unknown, 1 pageable, 2 pinned, 3 device, 4 array, 5 managed
+    print("== memcpy histogram by (copyKind, srcKind, dstKind)")
+    for r in q(
+        "SELECT copyKind, srcKind, dstKind, count(*) n, sum(bytes) b, sum(end-start) d, max(end-start) mx FROM CUPTI_ACTIVITY_KIND_MEMCPY GROUP BY copyKind, srcKind, dstKind ORDER BY n DESC"
+    ):
+        print(
+            f"  copyKind={r['copyKind']} src={r['srcKind']} dst={r['dstKind']}: n={r['n']} bytes={r['b']} total={r['d'] / 1e6:.1f} ms max={r['mx'] / 1e6:.2f} ms"
+        )
+    print("== pageable-source/destination memcpys per 10 s bucket")
+    for r in q(
+        "SELECT CAST(start/10e9 AS INT) b, count(*) n, sum(end-start) d FROM CUPTI_ACTIVITY_KIND_MEMCPY WHERE srcKind=1 OR dstKind=1 GROUP BY b ORDER BY b"
+    ):
+        print(
+            f"    t={r['b'] * 10:5d}s n={r['n']} total={r['d'] / 1e6:.1f} ms"
+        )
+    print("== API calls by name for memcpy-like calls")
+    if "CUPTI_ACTIVITY_KIND_RUNTIME" in T:
+        for r in q(
+            "SELECT nameId, count(*) n, sum(end-start) d, max(end-start) mx FROM CUPTI_ACTIVITY_KIND_RUNTIME GROUP BY nameId ORDER BY d DESC LIMIT 25"
+        ):
+            print(
+                f"  {s(r['nameId'])}: n={r['n']} total={r['d'] / 1e6:.1f} ms max={r['mx'] / 1e6:.2f} ms"
+            )
+
+for t in ("CUPTI_ACTIVITY_KIND_MEMCPY", "CUPTI_ACTIVITY_KIND_MEMSET"):
+    if t in T:
+        print(f"== last 15 {t} (start s, dur ms, bytes, kinds)")
+        c = cols(t)
+        kinds = [
+            k for k in ("copyKind", "srcKind", "dstKind", "memKind") if k in c
+        ]
+        for r in q(f"SELECT * FROM {t} ORDER BY start DESC LIMIT 15")[::-1]:
+            print(
+                f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} {r['bytes']} s{r['streamId']} "
+                + " ".join(f"{k}={r[k]}" for k in kinds)
+            )
+
+if "CUPTI_ACTIVITY_KIND_RUNTIME" in T:
+    print(
+        "== last 40 CUDA runtime/driver API calls (start s, dur ms, tid, name)"
+    )
+    for r in q(
+        "SELECT start, end, globalTid, nameId FROM CUPTI_ACTIVITY_KIND_RUNTIME ORDER BY start DESC LIMIT 40"
+    )[::-1]:
+        print(
+            f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} {r['globalTid'] % 0x1000000} {s(r['nameId'])}"
+        )
+    print("== slowest 10 API calls")
+    for r in q(
+        "SELECT start, end, globalTid, nameId FROM CUPTI_ACTIVITY_KIND_RUNTIME ORDER BY end-start DESC LIMIT 10"
+    ):
+        print(
+            f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} ms {s(r['nameId'])}"
+        )
+
+for t in T:
+    if not t.startswith("CUDA_UM_") and "UNIFIED_MEMORY" not in t:
+        continue
+    c = cols(t)
+    print(f"== {t} columns: {c}")
+    if "start" not in c:
+        continue
+    cnt = (
+        "sum(numberOfPageFaults)" if "numberOfPageFaults" in c else "count(*)"
+    )
+    print(f"  per 10 s bucket ({cnt}):")
+    for r in q(
+        f"SELECT CAST(start/10e9 AS INT) AS b, count(*) AS n, {cnt} AS f FROM {t} GROUP BY b ORDER BY b"
+    ):
+        print(f"    t={r['b'] * 10:5d}s events={r['n']} faults={r['f']}")
+    if "address" in c:
+        print("  top 15 2MiB regions by events:")
+        for r in q(
+            f"SELECT (address >> 21) AS reg, count(*) AS n, min(start) AS t0, max(start) AS t1 FROM {t} GROUP BY reg ORDER BY n DESC LIMIT 15"
+        ):
+            print(
+                f"    0x{r['reg'] << 21:x} n={r['n']} {r['t0'] / 1e9:.3f}..{r['t1'] / 1e9:.3f}s"
+            )
+    print("  last 15 events:")
+    for r in q(f"SELECT * FROM {t} ORDER BY start DESC LIMIT 15")[::-1]:
+        print(
+            "    "
+            + " ".join(f"{k}={r[k]}" for k in c if k not in ("globalPid",))
+        )
+
+if "OSRT_API" in T:
+    print("== OS runtime calls in the last 120 s of the capture, by name")
+    mx = q("SELECT max(start) FROM OSRT_API")[0][0] or 0
+    for r in q(
+        "SELECT nameId, count(*) n, sum(end-start) d, max(end) e FROM OSRT_API WHERE start > ? GROUP BY nameId ORDER BY d DESC LIMIT 15",
+        mx - 120e9,
+    ):
+        print(
+            f"  {s(r['nameId'])}: n={r['n']} total={r['d'] / 1e6:.1f} ms last_end={r['e'] / 1e9:.3f}s"
+        )
+    print("== last 20 OS runtime calls")
+    for r in q(
+        "SELECT start, end, globalTid, nameId FROM OSRT_API ORDER BY start DESC LIMIT 20"
+    )[::-1]:
+        print(
+            f"  {r['start'] / 1e9:.4f} {(r['end'] - r['start']) / 1e6:9.3f} {r['globalTid'] % 0x1000000} {s(r['nameId'])}"
+        )

ci/b3_ptrattr.py

@@ -0,0 +1,78 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# DO NOT MERGE: cudaPointerGetAttributes on several kinds of pageable host
+# memory, before and after the GPU touches host memory (host-input PQ build).
+
+import ctypes
+
+import cupy as cp
+import numpy as np
+
+rt = cp.cuda.runtime
+libc = ctypes.CDLL("libc.so.6")
+libc.malloc.restype = ctypes.c_void_p
+libc.malloc.argtypes = [ctypes.c_size_t]
+libc.memset.argtypes = [ctypes.c_void_p, ctypes.c_int, ctypes.c_size_t]
+
+TYPES = {0: "unregistered", 1: "host", 2: "device", 3: "managed"}
+
+
+def probe(tag, ptr):
+    try:
+        a = rt.pointerGetAttributes(ptr)
+        print(
+            f"  {tag:<44} ptr=0x{ptr:x} type={TYPES.get(a.type, a.type)} "
+            f"devicePointer=0x{a.devicePointer:x} hostPointer=0x{a.hostPointer:x}",
+            flush=True,
+        )
+    except Exception as e:  # noqa: BLE001
+        print(f"  {tag:<44} ptr=0x{ptr:x} ERROR {e!r}", flush=True)
+
+
+def round_(label):
+    print(f"== {label}")
+    big = np.zeros(1 << 20, dtype=np.float32)
+    probe("(a) fresh np.zeros(1<<20)", big.ctypes.data)
+    small = np.empty(100, dtype=np.int64)
+    small[:] = 7
+    probe("(b) np.empty(100) after write", small.ctypes.data)
+    m = libc.malloc(1024)
+    libc.memset(m, 1, 1024)
+    probe("(c) malloc(1024) after memset", m)
+    m2 = libc.malloc(64 * 8)
+    libc.memset(m2, 0, 64 * 8)
+    probe("(c2) malloc(64*8) after memset", m2)
+    return big, small
+
+
+def main():
+    print(
+        "pageableMemoryAccess",
+        rt.deviceGetAttribute(rt.cudaDevAttrPageableMemoryAccess, 0),
+    )
+    keep = [round_("before any GPU access to host memory")]
+
+    # (d) the GPU copies from pageable memory, then probe the same array again.
+    src = np.arange(1 << 16, dtype=np.float32)
+    keep.append(cp.asarray(src))
+    probe("(d) np array after cp.asarray (H2D copy)", src.ctypes.data)
+
+    # (e) host-input PQ build/transform, as in test_product_quantizer.
+    try:
+        from cuvs.preprocessing.quantize import pq
+
+        x = np.random.random_sample((1000, 64)).astype(np.float32)
+        params = pq.QuantizerParams(pq_bits=8, pq_dim=32)
+        q = pq.build(params, x)
+        pq.transform(q, x)
+        cp.cuda.runtime.deviceSynchronize()
+        probe("(e) PQ host input after pq.build/transform", x.ctypes.data)
+    except Exception as e:  # noqa: BLE001
+        print("PQ step failed:", repr(e))
+
+    keep.append(round_("after host-input PQ build"))
+
+
+if __name__ == "__main__":
+    main()