Summary
A numba.cuda device function that does arr.view(record_dtype)[0]
where record_dtype.itemsize is very large (order ~133 MB — a single
structured-dtype record spanning a whole memory region) silently
hangs a cooperative kernel instead of failing fast.
Indexing [0] materializes the record by value — i.e. a
per-thread local-memory copy of the entire region. At tens-to-hundreds
of MB per thread this never completes on device; all blocks of a
cooperative-grid kernel wedge uniformly. There is no error, no
diagnostic, no out-of-local-memory failure — just a hang.
The .view(dtype) itself is cheap. The by-value [0] copy is the
killer.
Reproducer shape
A device function that:
- takes a plain
uint32 (or uint8) device buffer,
.view(record_dtype) where record_dtype is ONE structured dtype
whose itemsize is tens-to-hundreds of MB,- indexes
[0] to get the record, and
- passes it somewhere (even to a no-op device function).
The kernel hangs on launch (cooperative grid: every block wedges at
the same point).
Small records — per-sub-region dtypes at KB scale — work fine with the
exact same idiom, because the by-value copy is bounded.
Expected
Diagnose / error / bound on oversized single-record local-memory
materialization, rather than silently emitting an unbounded per-thread
copy that wedges the kernel. Even a compile-time or launch-time
warning when a record-by-value copy exceeds a sane local-memory
threshold would turn a silent multi-minute hang into an actionable
failure.
Notes
This may well be expected numba record-by-value semantics at
extreme sizes and could be WONTFIX — filing to document the mechanism
and to ask whether a fast-fail / bound is feasible.
The consumer (xanadu) has already mitigated on its side: adapt
transfers now pass the raw region slice rather than performing a
by-value record copy, so the extreme-size [0] materialization no
longer occurs. This report is about numba-cuda's silent-hang behavior,
not about the downstream fix.
(Mirrored from an internal tracker bead, xanadu-lzx6.)
Summary
A
numba.cudadevice function that doesarr.view(record_dtype)[0]where
record_dtype.itemsizeis very large (order ~133 MB — a singlestructured-dtype record spanning a whole memory region) silently
hangs a cooperative kernel instead of failing fast.
Indexing
[0]materializes the record by value — i.e. aper-thread local-memory copy of the entire region. At tens-to-hundreds
of MB per thread this never completes on device; all blocks of a
cooperative-grid kernel wedge uniformly. There is no error, no
diagnostic, no out-of-local-memory failure — just a hang.
The
.view(dtype)itself is cheap. The by-value[0]copy is thekiller.
Reproducer shape
A device function that:
uint32(oruint8) device buffer,.view(record_dtype)whererecord_dtypeis ONE structured dtypewhose
itemsizeis tens-to-hundreds of MB,[0]to get the record, andThe kernel hangs on launch (cooperative grid: every block wedges at
the same point).
Small records — per-sub-region dtypes at KB scale — work fine with the
exact same idiom, because the by-value copy is bounded.
Expected
Diagnose / error / bound on oversized single-record local-memory
materialization, rather than silently emitting an unbounded per-thread
copy that wedges the kernel. Even a compile-time or launch-time
warning when a record-by-value copy exceeds a sane local-memory
threshold would turn a silent multi-minute hang into an actionable
failure.
Notes
This may well be expected numba record-by-value semantics at
extreme sizes and could be WONTFIX — filing to document the mechanism
and to ask whether a fast-fail / bound is feasible.
The consumer (xanadu) has already mitigated on its side: adapt
transfers now pass the raw region slice rather than performing a
by-value record copy, so the extreme-size
[0]materialization nolonger occurs. This report is about numba-cuda's silent-hang behavior,
not about the downstream fix.
(Mirrored from an internal tracker bead, xanadu-lzx6.)