diff --git a/src/srdatalog/ir/codegen/cuda/lower_ctx.py b/src/srdatalog/ir/codegen/cuda/lower_ctx.py new file mode 100644 index 0000000..fcbac70 --- /dev/null +++ b/src/srdatalog/ir/codegen/cuda/lower_ctx.py @@ -0,0 +1,115 @@ +'''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 dc3414c..6e4e9c2 100644 --- a/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py +++ b/src/srdatalog/ir/dialects/relation/sorted_array/lowerings/__init__.py @@ -19,9 +19,12 @@ 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, @@ -67,105 +70,215 @@ 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. - 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. + 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. ''' - 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) + # 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) ---------------------------- 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 new file mode 100644 index 0000000..f1a4ea3 --- /dev/null +++ b/tests/test_cuda_render_ctx_extraction.py @@ -0,0 +1,222 @@ +'''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