Bug: make_constant_array over-aligns constant arrays (uses element size instead of ABI alignment)

[shaunc]

Describe the bug

CUDATargetContext.make_constant_array sets the constant global's
alignment from the element ABI size (get_abi_sizeof) instead of its
ABI alignment (get_abi_alignment), then rounds up to a power of two
(target.py:311-312). For scalar dtypes size ≈ alignment, so the result
is accidentally correct and the bug is masked. But for a Record
(structured) dtype the value becomes next_pow2(itemsize) — e.g. a
49,824-byte record is emitted with .const .align 65536, where the
correct alignment is the record's max field alignment (4 bytes for an
all-float32 record).

A 64 KiB alignment on a __constant__ object is meaningless to the
hardware (constant loads have no such alignment requirement) and is
actively harmful: a 64 KiB-aligned object must sit at offset 0 of the
64 KiB constant bank, so it alone consumes the bank base and crowds out
other constants — turning a modest constant footprint into a spurious
File uses too much global constant data link failure.

Steps/Code to reproduce bug

import numpy as np
from numba import cuda, float32

# A Record dtype: itemsize = 12456 * 4 = 49,824 B, but every field is
# float32 (natural alignment 4).
_REC = np.dtype([("cfg", np.float32, (12456,))])
_CONST = np.zeros(1, dtype=_REC)

@cuda.jit
def k(out):
    rec = cuda.const.array_like(_CONST)
    i = cuda.grid(1)
    if i < out.size:
        out[i] = rec[0]["cfg"][i % 12456]

k.compile((float32[:],))
_, ptx = next(iter(k.inspect_asm().items()))
print([ln.strip() for ln in ptx.splitlines() if "_cudapy_cmem" in ln][0])

Actual output:

.const .align 65536 .b8 _cudapy_cmem[49824];

record itemsize = 49824 B, record field alignment = 4 B, yet the
emitted alignment is 65536. A plain float32 constant array (scalar
dtype, same code path) correctly emits .align 4, which is why the bug is
only visible for large structured dtypes.

Expected behavior

A bounded alignment that does not scale with the array/record size.
The minimal correct value is the element's natural ABI alignment (4 B
here):

.const .align 4 .b8 _cudapy_cmem[49824];

Root cause — numba_cuda/numba/cuda/target.py,
CUDATargetContext.make_constant_array (defined at line 278), alignment
block at lines 309-312; the defect is line 311:

# target.py:309
        # Preserve the underlying alignment
        lldtype = self.get_data_type(aryty.dtype)
        align = self.get_abi_sizeof(lldtype)       # line 311: ABI SIZE, not alignment
        gv.align = 2 ** (align - 1).bit_length()   # line 312: next pow2 >= SIZE

get_abi_sizeof(ty) returns ty.get_abi_size(...) (the ABI size),
whereas get_abi_alignment(ty) returns ty.get_abi_alignment(...) (the
ABI alignment). The variable is named align and the comment says
"Preserve the underlying alignment," but the value read is the size.

Proposed fix (one line, target.py:311):

         lldtype = self.get_data_type(aryty.dtype)
-        align = self.get_abi_sizeof(lldtype)       # line 311
+        align = self.get_abi_alignment(lldtype)    # line 311
         gv.align = 2 ** (align - 1).bit_length()   # line 312

get_abi_alignment already returns a power-of-two alignment for scalar
and struct types, so the subsequent next_pow2 round becomes a harmless
safety net and existing scalar-dtype behaviour is unchanged.

Note on policy (not required for the fix): the essential defect is that
alignment is derived from size (unbounded → 64 KiB here), not that it
must be exactly the natural ABI alignment. A deliberate cache-line
minimum (e.g. the 128 B L1/L2 line, or the constant-cache line) would be a
defensible enhancement for access locality — but as a bounded named
constant, max(get_abi_alignment(lldtype), CACHE_LINE), never as a
function of the array size.

Environment details (please complete the following information):

• Environment location: Bare-metal (Ubuntu 24.04.4 LTS, x86_64)
• Method of numba-cuda install: pip (from PyPI)
• numba-cuda 0.24.0, numba 0.62.1, Python 3.13.1
• CUDA runtime 13.1, driver 13.3 (NVIDIA driver 610.43.02)
• GPU: NVIDIA RTX 5090 (sm_120); also reproduced on RTX 3070 (sm_86)

Additional context

Reproduces deterministically at compile time (no launch / no specific GPU
needed — the PTX is emitted for the compiled kernel). Downstream, this
inflates constant-bank pressure for any kernel that freezes a large
structured array into constant memory via cuda.const.array_like, and can
cause spurious 64 KiB constant-bank overflows at device-link time. It is
also a portability hazard, since the emitted alignment tracks struct size
rather than the ABI.

[arcusbuilds]

/assign I would like to work on it.