Source code for nki.compiler

# AUTO GENERATED FILE - DO NOT MANUALLY MODIFY.

from __future__ import annotations
from dataclasses import dataclass
from enum import Enum
from typing import *

[docs]class CompileConfig(): r"""Per-kernel compile configuration passed to ``@nki.jit(compile_config=...)``. Options that impact downstream compiler behavior.""" address_rotation: AddressRotation = ... r"""Backend address-rotation policy (an :class:`AddressRotation`).""" non_ssa_legalization: bool = ... r"""Enable the backend's non-SSA legalization pass on this kernel's memlocs.""" dead_code_elim: bool = ... r"""Leave this kernel's memlocs eligible for backend dead-code elimination.""" barrier_check: bool = ... r"""Compile with the LNC core barrier checker enabled.""" prefetch_auto: bool = ... r"""Enable AUTO DMA-prefetch allocation: the compiler classifies load destinations automatically, reserves a fixed prefetch ring + slack band in SBUF, and lets compute/prefetch spill into the shared slack. Mutually exclusive with a manual ``NKI_PREFETCH_RING_SIZE``. Off by default.""" nki_alloc_fallback: bool = ... r"""Fall back to the default neuronx-cc allocation flow when NKI's own LinearScanAllocation fails during compilation. On by default; set False to surface the failure (e.g. an SBUF/PSUM out-of-memory) instead of retrying.""" ...
[docs]class SpecializationConfig(): r"""Per-kernel specialization-time config passed to ``@nki.jit(specialization_config=...)``. Options that tune the specialization phase (turning the kernel into IR).""" allow_nested_config_override: bool = ... r"""When this kernel is the top-level compiled kernel, warn (instead of raising) if a nested @nki.jit sub-kernel's config disagrees, and let it inline under this kernel's config.""" ...
[docs]class AddressRotation(Enum): r"""Whether the backend may rotate a memloc's address for better layout. - ``AUTO`` (default): decide per memloc by allocation origin — a memloc with a fixed address (``bind_memloc``, whether user-supplied or chosen by NKI's LinearScanAllocation) is opted OUT of rotation so the backend does not reshuffle an address the kernel already fixed; a memloc the backend allocates keeps rotation ON so it can still pick a good layout. - ``ENABLE``: keep rotation on for every memloc. - ``DISABLE``: opt every memloc out of rotation.""" AUTO = 'auto' ENABLE = 'enable' DISABLE = 'disable' ...