diff --git a/src/srdatalog/ir/codegen/cuda/lower_ctx.py b/src/srdatalog/ir/codegen/cuda/lower_ctx.py deleted file mode 100644 index fcbac70..0000000 --- a/src/srdatalog/ir/codegen/cuda/lower_ctx.py +++ /dev/null @@ -1,115 +0,0 @@ -'''CudaRenderCtx — CUDA-target-private render scratch for MIR -> IIR lowering. - -Per `docs/phase_decomposition_redesign.md` § 3.2.1. - -The dialect-level `LoweringCtx` -(`src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py`) -historically carried both: - - - structural / lowering-time state (e.g. `name_counter`, `bound_vars`, - `handle_vars`, `inside_cartesian`) — target-agnostic; - - CUDA C++ identifier scratch (e.g. `view_var_names`, `output_var`, - `tiled_cartesian_valid_var`) — CUDA render-private. - -PR-1d extracts the second bucket into `CudaRenderCtx`, the -target-private bag. `LoweringCtx` retains the structural fields and -exposes the render-private fields via forwarding properties for back- -compat. Per the spec, pragma scratch booleans -(`is_counting`, `dedup_hash`, `tiled_cartesian`, `ws_enabled`, -`bg_enabled`) are NOT moved here — they remain on `LoweringCtx` for -now and are scheduled for removal in PR-1e. - -Design notes: - - Plain `@dataclass` (NOT `frozen=True`, NOT `slots=True`). The - spec shows `frozen=True, slots=True` but block-group test paths - rely on dynamic attribute writes; the dataclass must remain - mutable + non-slotted. - - The companion framework-level `LowerCtx` - (`src/srdatalog/ir/core/lower_ctx.py`) is OFF LIMITS per D10. - `CudaRenderCtx` lives entirely under `codegen/cuda/`; it never - appears in the framework-level type signatures. - -This module also hosts `NegPreNarrowInfo`. It used to live in -`lowerings/__init__.py`; PR-1d moves it here because it is purely -CUDA-render shape (C++ var names, view-var names, etc.) and the -`neg_pre_narrow` field type-references it. `lowerings/__init__.py` -re-exports the name for any existing imports. -''' - -from __future__ import annotations - -from dataclasses import dataclass, field - - -@dataclass -class NegPreNarrowInfo: - '''Pre-narrowed handle info for a Negation that follows a Cartesian. - - When a Negation's prefix vars are all (or partly) bound *before* - the Cartesian, those vars don't change inside the Cartesian loop — - so we can apply them once cooperatively before the loop, then - cheaply check `valid()` per iteration. The remaining (in-Cartesian) - vars are applied per-thread inside the loop via `prefix_seq`. - - Every field carries CUDA C++ identifier shape (var names emitted - into rendered kernel source), so this lives under `codegen/cuda/`. - - Mirrors the legacy `NegPreNarrowInfo` in - ir/codegen/cuda/context.py. - ''' - - var_name: str - pre_vars: list[str] - in_cartesian_vars: list[str] - pre_consts: list[tuple[int, int]] - view_var: str - rel_name: str - - -@dataclass -class CudaRenderCtx: - '''CUDA-target-private render scratch threaded through MIR -> IIR. - - All fields hold CUDA C++ identifier scratch or render-time scratch - read by CUDA-shaped IIR-op emission. The renderer (and the IIR-op - lowerings that emit CUDA-shaped ops) read these via the forwarding - properties on `LoweringCtx`; the structural lowering logic does - not touch them. - - Field tally: - 1. view_var_names — handle_idx -> C++ view var name - 2. output_var — current C++ OutputContext var - 3. output_var_overrides — rel_name -> override OutputContext - 4. view_slot_bases — handle_idx -> base slot in views[] - 5. rel_index_types — rel_name -> C++ index template - 6. tiled_cartesian_valid_var — tiled-Cart ballot validity var - 7. ws_cartesian_valid_var — WS Cart ballot validity var - 8. neg_pre_narrow — handle_idx -> pre-narrowed handle info - 9. debug — emit `printf` debug breadcrumbs - 10. tile_var — CUB cooperative_groups tile var name - - Total: 10 fields. The spec table (§ 2.2.2) listed `dedup_hash_vars` - and `ws_cartesian_bound_vars` too — but those live on the legacy - `CodeGenContext` (`ir/codegen/cuda/context.py`), not on the current - `LoweringCtx`. Per PR-1d's discipline (trust the monolith), they - are NOT extracted here. They will land on `CudaRenderCtx` once the - dialect-level `LoweringCtx` grows them (planned alongside the - pragma-scratch-flag removal in PR-1e+). - ''' - - view_var_names: dict[str, str] = field(default_factory=dict) - output_var: str = 'output' - output_var_overrides: dict[str, str] = field(default_factory=dict) - view_slot_bases: dict[str, int] = field(default_factory=dict) - rel_index_types: dict[str, str] = field(default_factory=dict) - tiled_cartesian_valid_var: str = '' - ws_cartesian_valid_var: str = '' - neg_pre_narrow: dict[int, NegPreNarrowInfo] = field(default_factory=dict) - debug: bool = True - tile_var: str = 'tile' - - -__all__ = [ - 'CudaRenderCtx', - 'NegPreNarrowInfo', -] diff --git a/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py b/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py index 6e4e9c2..dc3414c 100644 --- a/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py +++ b/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py @@ -19,12 +19,9 @@ from __future__ import annotations from collections.abc import Callable +from dataclasses import dataclass, field import srdatalog.ir.mir.types as mir -from srdatalog.ir.codegen.cuda.lower_ctx import ( - CudaRenderCtx, - NegPreNarrowInfo, -) from srdatalog.ir.core import Op from srdatalog.ir.dialects.iir.cf import ( AddCount, @@ -70,215 +67,105 @@ from srdatalog.ir.hir.types import Version +@dataclass +class NegPreNarrowInfo: + '''Pre-narrowed handle info for a Negation that follows a Cartesian. + + When a Negation's prefix vars are all (or partly) bound *before* + the Cartesian, those vars don't change inside the Cartesian loop — + so we can apply them once cooperatively before the loop, then + cheaply check `valid()` per iteration. The remaining (in-Cartesian) + vars are applied per-thread inside the loop via `prefix_seq`. + + Mirrors the legacy `NegPreNarrowInfo` in + ir/dialects/target/cuda/context.py. + ''' + + var_name: str + pre_vars: list[str] + in_cartesian_vars: list[str] + pre_consts: list[tuple[int, int]] + view_var: str + rel_name: str + + +@dataclass class LoweringCtx: '''Mutable state during MIR -> IIR walk. - Carries structural lowering state (counters, lexical-scope flags, - bound vars, handle aliasing) directly; carries CUDA-target-private - render scratch via the embedded `render_ctx: CudaRenderCtx`. The - render-private fields are exposed as forwarding properties so the - 100+ existing call sites that say e.g. `ctx.view_var_names` keep - working without an edit (back-compat seam per D20). - - PR-1d split — per `docs/phase_decomposition_redesign.md` § 3.2.1. - Pragma scratch booleans (`is_counting`, `dedup_hash`, - `tiled_cartesian`, `ws_enabled`, `bg_enabled`) remain on the - dialect-level `LoweringCtx` for now; they will move OFF both - contexts in PR-1e (per the spec's pragma-lowering plan). - - Constructor accepts both the legacy flat keyword names - (`view_var_names=...`, `output_var=...`, etc.) and an explicit - `render_ctx=` — the back-compat path routes the flat kwargs into a - fresh `CudaRenderCtx`. Calls that pass `render_ctx=` (the new - preferred shape) must not also pass flat render-side kwargs. + Mirrors the legacy `CodeGenContext` for the fields that matter to + the dialect's emission decisions today. Other legacy fields + (tiled_cartesian state, ws state, etc.) aren't needed yet — + milestones add them as they cover those paths. ''' - # Structural fields (defaults match the legacy dataclass). - name_counter: int - inside_cartesian: bool - bound_vars: list[str] - is_counting: bool - dedup_hash: bool - tiled_cartesian: bool - ws_enabled: bool - bg_enabled: bool - handle_vars: dict[str, str] - cartesian_bound_vars: list[str] - # CUDA-target-private render scratch. - render_ctx: CudaRenderCtx - - # ---- Constructor ------------------------------------------------- - - def __init__( - self, - *, - # Structural fields. - name_counter: int = 0, - inside_cartesian: bool = False, - bound_vars: list[str] | None = None, - handle_vars: dict[str, str] | None = None, - cartesian_bound_vars: list[str] | None = None, - # Pragma scratch (kept here pending PR-1e removal). - is_counting: bool = False, - dedup_hash: bool = False, - tiled_cartesian: bool = False, - ws_enabled: bool = False, - bg_enabled: bool = False, - # New: explicit render-ctx (preferred). Mutually exclusive with - # the flat render-side kwargs below. - render_ctx: CudaRenderCtx | None = None, - # Back-compat: flat render-side kwargs. Routed into a fresh - # CudaRenderCtx when `render_ctx` is not supplied. - view_var_names: dict[str, str] | None = None, - output_var: str | None = None, - output_var_overrides: dict[str, str] | None = None, - view_slot_bases: dict[str, int] | None = None, - rel_index_types: dict[str, str] | None = None, - tiled_cartesian_valid_var: str | None = None, - ws_cartesian_valid_var: str | None = None, - neg_pre_narrow: dict[int, NegPreNarrowInfo] | None = None, - debug: bool | None = None, - tile_var: str | None = None, - ) -> None: - flat_render_kwargs: dict[str, object] = { - 'view_var_names': view_var_names, - 'output_var': output_var, - 'output_var_overrides': output_var_overrides, - 'view_slot_bases': view_slot_bases, - 'rel_index_types': rel_index_types, - 'tiled_cartesian_valid_var': tiled_cartesian_valid_var, - 'ws_cartesian_valid_var': ws_cartesian_valid_var, - 'neg_pre_narrow': neg_pre_narrow, - 'debug': debug, - 'tile_var': tile_var, - } - if render_ctx is not None: - conflicting = [k for k, v in flat_render_kwargs.items() if v is not None] - if conflicting: - raise TypeError( - f'LoweringCtx: cannot pass `render_ctx=` together with flat render kwargs: {conflicting}' - ) - self.render_ctx = render_ctx - else: - # Build a CudaRenderCtx, then assign any explicitly-supplied - # flat kwargs onto it. Explicit per-field assignment (rather - # than `**supplied`) keeps mypy happy on the field types. - self.render_ctx = CudaRenderCtx() - if view_var_names is not None: - self.render_ctx.view_var_names = view_var_names - if output_var is not None: - self.render_ctx.output_var = output_var - if output_var_overrides is not None: - self.render_ctx.output_var_overrides = output_var_overrides - if view_slot_bases is not None: - self.render_ctx.view_slot_bases = view_slot_bases - if rel_index_types is not None: - self.render_ctx.rel_index_types = rel_index_types - if tiled_cartesian_valid_var is not None: - self.render_ctx.tiled_cartesian_valid_var = tiled_cartesian_valid_var - if ws_cartesian_valid_var is not None: - self.render_ctx.ws_cartesian_valid_var = ws_cartesian_valid_var - if neg_pre_narrow is not None: - self.render_ctx.neg_pre_narrow = neg_pre_narrow - if debug is not None: - self.render_ctx.debug = debug - if tile_var is not None: - self.render_ctx.tile_var = tile_var - - self.name_counter = name_counter - self.inside_cartesian = inside_cartesian - self.bound_vars = bound_vars if bound_vars is not None else [] - self.handle_vars = handle_vars if handle_vars is not None else {} - self.cartesian_bound_vars = cartesian_bound_vars if cartesian_bound_vars is not None else [] - self.is_counting = is_counting - self.dedup_hash = dedup_hash - self.tiled_cartesian = tiled_cartesian - self.ws_enabled = ws_enabled - self.bg_enabled = bg_enabled - - # ---- Forwarding properties (back-compat seam, D20) --------------- - - @property - def view_var_names(self) -> dict[str, str]: - return self.render_ctx.view_var_names - - @view_var_names.setter - def view_var_names(self, value: dict[str, str]) -> None: - self.render_ctx.view_var_names = value - - @property - def output_var(self) -> str: - return self.render_ctx.output_var - - @output_var.setter - def output_var(self, value: str) -> None: - self.render_ctx.output_var = value - - @property - def output_var_overrides(self) -> dict[str, str]: - return self.render_ctx.output_var_overrides - - @output_var_overrides.setter - def output_var_overrides(self, value: dict[str, str]) -> None: - self.render_ctx.output_var_overrides = value - - @property - def view_slot_bases(self) -> dict[str, int]: - return self.render_ctx.view_slot_bases - - @view_slot_bases.setter - def view_slot_bases(self, value: dict[str, int]) -> None: - self.render_ctx.view_slot_bases = value - - @property - def rel_index_types(self) -> dict[str, str]: - return self.render_ctx.rel_index_types - - @rel_index_types.setter - def rel_index_types(self, value: dict[str, str]) -> None: - self.render_ctx.rel_index_types = value - - @property - def tiled_cartesian_valid_var(self) -> str: - return self.render_ctx.tiled_cartesian_valid_var - - @tiled_cartesian_valid_var.setter - def tiled_cartesian_valid_var(self, value: str) -> None: - self.render_ctx.tiled_cartesian_valid_var = value - - @property - def ws_cartesian_valid_var(self) -> str: - return self.render_ctx.ws_cartesian_valid_var - - @ws_cartesian_valid_var.setter - def ws_cartesian_valid_var(self, value: str) -> None: - self.render_ctx.ws_cartesian_valid_var = value - - @property - def neg_pre_narrow(self) -> dict[int, NegPreNarrowInfo]: - return self.render_ctx.neg_pre_narrow - - @neg_pre_narrow.setter - def neg_pre_narrow(self, value: dict[int, NegPreNarrowInfo]) -> None: - self.render_ctx.neg_pre_narrow = value - - @property - def debug(self) -> bool: - return self.render_ctx.debug - - @debug.setter - def debug(self, value: bool) -> None: - self.render_ctx.debug = value - - @property - def tile_var(self) -> str: - return self.render_ctx.tile_var - - @tile_var.setter - def tile_var(self, value: str) -> None: - self.render_ctx.tile_var = value - - # ---- Fresh-name generator (unchanged) ---------------------------- + name_counter: int = 0 + view_var_names: dict[str, str] = field(default_factory=dict) + is_counting: bool = False + inside_cartesian: bool = False + output_var: str = 'output' + tile_var: str = 'tile' + debug: bool = True + output_var_overrides: dict[str, str] = field(default_factory=dict) + bound_vars: list[str] = field(default_factory=list) + # rel_name -> custom index type code (e.g. 'Device2LevelIndex'). + # Empty string / missing entry = plain DSAI single-segment. + rel_index_types: dict[str, str] = field(default_factory=dict) + # handle_idx (str) -> base slot in views[]. Populated by + # `compile_kernel_body` from `emit_view_declarations` so D2L + # segment-loop emission can reference the right HEAD/FULL pair. + view_slot_bases: dict[str, int] = field(default_factory=dict) + # When True, InsertInto emit wraps the output write in a + # `{ bool _p = dedup_table.try_insert(thread_id, ...); if (_p) { + # ... } }` gate, and the materialize-phase write goes through + # `atomicAdd(atomic_write_pos, 1u)` + `out_data_0[...]` instead of + # `output.emit_direct(...)`. Threaded from `ep.dedup_hash`. + dedup_hash: bool = False + # When True, eligible 2-source / 1-var-per-source nested Cartesians + # in materialize phase emit the `if (total > 32) { tiled smem path } + # else { fallback path }` dispatch. Bodies inside both branches use + # the `tiled_cartesian_valid_var` ballot-write variant of InsertInto. + # Threaded from `_dialect_safe_kernel`'s `tiled_cartesian_eligible` + # gate via `compile_kernel_body`. + tiled_cartesian: bool = False + # Empty (default) — InsertInto emits the standard + # `output.emit_direct(...)` write. Non-empty — emits the tiled- + # Cartesian ballot-write variant guarded by this var. Set by + # `_lower_nested_cart` when rendering the tiled-mode body. + tiled_cartesian_valid_var: str = '' + # Work-stealing flag (mirrors legacy `ctx.ws_enabled`). In count + # phase, InsertInto emits `++` instead of + # `.emit_direct()` — the WS count uses a per-thread + # local counter that the runner aggregates. The legacy emitter has + # only this kernel-functor-level WS support; the runner-side WS + # scaffolding (WCOJTask queue) was never finished. + ws_enabled: bool = False + # Work-stealing batched-Cartesian valid var. When set, Filter / + # Negation fold their guard into ` = && ();` and + # InsertInto materialize emits `.emit_warp_coalesced( + # tile, , )` — a cooperative warp write instead of a + # lane-zero-guarded emit_direct. Mirrors legacy + # `ctx.ws_cartesian_valid_var`. + ws_cartesian_valid_var: str = '' + # Block-group flag (mirrors legacy `ctx.bg_enabled`). When True, + # the root multi-source ColumnJoin emits via `BgRootCjMulti` + # (block-group work-balanced partition + binary-search key loop) + # instead of the standard grid-stride root_unique_values loop. + # Cleared when descending into the body (the BG root's narrowed + # handle already restricts work to this warp's slice). + bg_enabled: bool = False + # State-key -> handle var name. Lets nested CJ find the parent + # handle to alias by the same (rel, cols, prefix_vars, ver) key + # that the outer CJ used to register it. + handle_vars: dict[str, str] = field(default_factory=dict) + # Cartesian-bound var names. Used to decide which Negation prefix + # vars are pre-Cartesian (= bound by an outer scope) vs in-Cartesian + # (= bound by the current Cart and per-thread). + cartesian_bound_vars: list[str] = field(default_factory=list) + # handle_idx -> NegPreNarrowInfo. Populated by `_lower_nested_cart` + # before its body renders so that the body's Negation handler can + # pick up the pre-allocated handle. + neg_pre_narrow: dict[int, NegPreNarrowInfo] = field(default_factory=dict) def fresh(self, prefix: str) -> str: self.name_counter += 1 diff --git a/tests/test_cuda_render_ctx_extraction.py b/tests/test_cuda_render_ctx_extraction.py deleted file mode 100644 index f1a4ea3..0000000 --- a/tests/test_cuda_render_ctx_extraction.py +++ /dev/null @@ -1,222 +0,0 @@ -'''PR-1d: CudaRenderCtx extraction tests. - -Verifies: - - 1. `CudaRenderCtx` is importable from - `srdatalog.ir.codegen.cuda.lower_ctx` and constructs with - defaults. - 2. `LoweringCtx.render_ctx` is a real `CudaRenderCtx` instance. - 3. The forwarding properties round-trip read + write through - `ctx.` <-> `ctx.render_ctx.` for every field - extracted into `CudaRenderCtx`. - 4. Passing flat render kwargs (back-compat path) and passing - `render_ctx=` (preferred path) both produce a ctx whose - observable state agrees on every render-side field. - 5. A sample MIR pipeline (`Scan -> InsertInto`) lowers to the - identical rendered IIR text under both ctx-construction shapes - — i.e. extraction is byte-equivalent. - -Per `docs/phase_decomposition_redesign.md` § 3.2.1 and the PR-1d -spec brief. -''' - -from __future__ import annotations - -from typing import Any - -import pytest - -import srdatalog.ir.mir.types as mir -from srdatalog.ir.codegen.cuda.emit import EmitCtx, emit -from srdatalog.ir.codegen.cuda.lower_ctx import CudaRenderCtx, NegPreNarrowInfo -from srdatalog.ir.dialects.relation.sorted_array.lowerings import ( - LoweringCtx, - lower_scan_pipeline, -) -from srdatalog.ir.hir.types import Version - -# ----------------------------------------------------------------------------- -# 1. CudaRenderCtx surface -# ----------------------------------------------------------------------------- - - -def test_cuda_render_ctx_default_construct() -> None: - '''CudaRenderCtx() with no args yields a fresh dataclass with - every field at its declared default. Mutable dataclass — NOT - frozen, NOT slotted (per the spec brief: dynamic setattr must - work).''' - rc = CudaRenderCtx() - assert rc.view_var_names == {} - assert rc.output_var == 'output' - assert rc.output_var_overrides == {} - assert rc.view_slot_bases == {} - assert rc.rel_index_types == {} - assert rc.tiled_cartesian_valid_var == '' - assert rc.ws_cartesian_valid_var == '' - assert rc.neg_pre_narrow == {} - assert rc.debug is True - assert rc.tile_var == 'tile' - - # Mutable — direct field assignment works (no FrozenInstanceError). - rc.output_var = 'ctx_other' - assert rc.output_var == 'ctx_other' - - -def test_cuda_render_ctx_not_frozen() -> None: - '''Dynamic attribute write must succeed — the BG pragma test path - relies on `object.__setattr__`-style mutation. If we ever switch - to `frozen=True` or `slots=True`, this test catches the - regression.''' - rc = CudaRenderCtx() - # Plain setattr (would FrozenInstanceError if frozen=True). - rc.tiled_cartesian_valid_var = 'tc_valid_1' - assert rc.tiled_cartesian_valid_var == 'tc_valid_1' - # Dynamic setattr — would AttributeError on a slotted dataclass - # if `__dict__` is missing. - object.__setattr__(rc, 'output_var', 'overridden') - assert rc.output_var == 'overridden' - - -# ----------------------------------------------------------------------------- -# 2. LoweringCtx embeds CudaRenderCtx + forwards through properties -# ----------------------------------------------------------------------------- - - -def test_lowering_ctx_default_render_ctx_is_fresh() -> None: - '''A bare `LoweringCtx()` constructs its own CudaRenderCtx via the - default-factory shape. The render_ctx is a real instance — not a - shared mutable.''' - ctx_a = LoweringCtx() - ctx_b = LoweringCtx() - assert isinstance(ctx_a.render_ctx, CudaRenderCtx) - assert isinstance(ctx_b.render_ctx, CudaRenderCtx) - assert ctx_a.render_ctx is not ctx_b.render_ctx - # Mutating one's dict must NOT bleed into the other. - ctx_a.render_ctx.view_var_names['x'] = 'view_x' - assert ctx_b.render_ctx.view_var_names == {} - - -FORWARD_FIELD_FIXTURES: list[tuple[str, Any, Any]] = [ - ('view_var_names', {'0': 'view_one'}, {'1': 'view_two'}), - ('output_var', 'ctx0', 'ctx_other'), - ('output_var_overrides', {'Rel': 'ctx_rel'}, {'Rel2': 'ctx_rel2'}), - ('view_slot_bases', {'0': 0, '1': 2}, {'2': 4}), - ('rel_index_types', {'R': 'Device2LevelIndex'}, {'R': ''}), - ('tiled_cartesian_valid_var', 'tc_valid_1', 'tc_valid_2'), - ('ws_cartesian_valid_var', 'ws_valid_3', 'ws_valid_4'), - ( - 'neg_pre_narrow', - { - 0: NegPreNarrowInfo( - var_name='h_pre_1', - pre_vars=['v0'], - in_cartesian_vars=['v1'], - pre_consts=[(0, 5)], - view_var='view_neg_0', - rel_name='Neg', - ), - }, - {}, - ), - ('debug', False, True), - ('tile_var', 'tile', 'tile_other'), -] - - -@pytest.mark.parametrize(('field', 'initial', 'replacement'), FORWARD_FIELD_FIXTURES) -def test_forwarding_property_read_roundtrip(field: str, initial: Any, replacement: Any) -> None: - '''`ctx.` reads must alias `ctx.render_ctx.` — - no shadowing dataclass field on LoweringCtx.''' - rc = CudaRenderCtx() - setattr(rc, field, initial) - ctx = LoweringCtx(render_ctx=rc) - assert getattr(ctx, field) == initial - # Read again after a direct mutation on the inner ctx — the - # property must NOT have cached a snapshot. - setattr(rc, field, replacement) - assert getattr(ctx, field) == replacement - - -@pytest.mark.parametrize(('field', 'initial', 'replacement'), FORWARD_FIELD_FIXTURES) -def test_forwarding_property_write_roundtrip(field: str, initial: Any, replacement: Any) -> None: - '''`ctx. = X` must route to `ctx.render_ctx.` — no - shadow attribute on the outer LoweringCtx.''' - ctx = LoweringCtx() - setattr(ctx, field, initial) - assert getattr(ctx.render_ctx, field) == initial - setattr(ctx, field, replacement) - assert getattr(ctx.render_ctx, field) == replacement - - -# ----------------------------------------------------------------------------- -# 3. Flat-kwargs construction equivalence -# ----------------------------------------------------------------------------- - - -def test_flat_kwargs_and_render_ctx_kwargs_agree() -> None: - '''Constructing with flat render kwargs vs. constructing with a - pre-built render_ctx yields the same observable render-side state.''' - view_var_names = {'0': 'view_src'} - ctx_flat = LoweringCtx(view_var_names=view_var_names, output_var='ctx0') - ctx_built = LoweringCtx( - render_ctx=CudaRenderCtx(view_var_names=dict(view_var_names), output_var='ctx0') - ) - assert ctx_flat.view_var_names == ctx_built.view_var_names - assert ctx_flat.output_var == ctx_built.output_var - # Identity check on render_ctx itself. - assert ctx_flat.render_ctx.view_var_names == ctx_built.render_ctx.view_var_names - - -def test_render_ctx_kwarg_rejects_flat_render_kwargs() -> None: - '''Passing both `render_ctx=` AND flat render kwargs must fail - loudly — silently merging would be a footgun.''' - rc = CudaRenderCtx(output_var='ctx0') - with pytest.raises(TypeError, match='render_ctx'): - LoweringCtx(render_ctx=rc, view_var_names={'0': 'view_x'}) - - -# ----------------------------------------------------------------------------- -# 4. Sample MIR lowers to byte-equivalent output through extracted ctx -# ----------------------------------------------------------------------------- - - -def _scan_insert_pipeline() -> tuple[mir.Scan, mir.InsertInto]: - scan = mir.Scan( - vars=['v0', 'v1'], - rel_name='Src', - version=Version.FULL, - index=[0, 1], - handle_start=0, - ) - ins = mir.InsertInto( - rel_name='Dst', - version=Version.NEW, - vars=['v0', 'v1'], - index=[0, 1], - ) - return scan, ins - - -def _render_iir(op: Any) -> str: - return emit(op, EmitCtx(indent_level=2)) - - -def test_byte_equivalent_under_flat_vs_render_ctx_kwargs() -> None: - '''A sample MIR program lowered with a flat-kwargs LoweringCtx vs. - a render_ctx-kwarg LoweringCtx must yield identical rendered IIR - text. Extraction MUST be byte-preserving.''' - scan, ins = _scan_insert_pipeline() - view_var_names = {'0': 'view_src'} - - ctx_flat = LoweringCtx(view_var_names=view_var_names, output_var='ctx0') - text_flat = _render_iir(lower_scan_pipeline([scan, ins], ctx_flat)) - - ctx_built = LoweringCtx( - render_ctx=CudaRenderCtx(view_var_names=dict(view_var_names), output_var='ctx0') - ) - text_built = _render_iir(lower_scan_pipeline([scan, ins], ctx_built)) - - assert text_flat == text_built - # Sanity: the render reached actual content. - assert 'view_src' in text_flat - assert 'ctx0' in text_flat