srdatalog.ir.codegen.cuda.materialized

Materialized binary-join code generation.

Port of src/srdatalog/codegen/target_jit/jit_materialized.nim.

Materialized joins take a different shape from fused joins (kernel functors with nested loops). They emit host-side C++ that:

  1. Materializes row-id pairs in buffers between join stages

  2. Uses binary-search probes against already-indexed relations

  3. Applies merge-path load balancing for unbalanced outputs

The generated C++ uses Thrust primitives (thrust::lower_bound / upper_bound, exclusive_scan, gather) plus three CUDA kernels declared in gen_materialized_join_helpers().

Public API:

is_materialized_pipeline(ops) -> bool True iff any op is a ProbeJoin. jit_complete_runner uses this as an early-dispatch check; when True it emits a materialized runner instead of a kernel functor.

gen_materialized_runner(node, db_type_name) -> str The main entry. Given an ExecutePipeline MIR node, emits a host- side Thrust executor: struct JitRunner_ { using DB = ; using FirstSchema = …; … static void execute(DB& db, uint32_t iteration = 0) { … } };

gen_materialized_join_helpers() -> str Returns the three global CUDA kernels (probe_count_matches_kernel, probe_materialize_pairs_kernel, gather_column_kernel) that the runner’s execute() body calls.

gen_materialized_join_kernel(ops, rule_name, ctx) -> str Legacy in-kernel variant. Not used by Nim’s live code and not exercised by our fixtures — ported for completeness.

Module Contents

Functions

gen_materialized_join_helpers

Three global CUDA helper kernels + comments. Embedded verbatim into the generated batch file when a materialized pipeline is present.

gen_materialized_join_kernel

Emit in-kernel materialized-join code. Unused by Nim’s live codegen (superseded by the host-side runner below) but ported for completeness.

gen_materialized_runner

Emit a host-side Thrust executor for a materialized-join pipeline.

is_materialized_pipeline

True iff the pipeline contains a ProbeJoin — which switches dispatch to the materialized runner.

API

srdatalog.ir.codegen.cuda.materialized.gen_materialized_join_helpers() → str[source]

Three global CUDA helper kernels + comments. Embedded verbatim into the generated batch file when a materialized pipeline is present.

srdatalog.ir.codegen.cuda.materialized.gen_materialized_join_kernel(ops: list[srdatalog.ir.mir.types.MirNode], rule_name: str, ctx: srdatalog.ir.codegen.cuda.context.CodeGenContext) → str[source]

Emit in-kernel materialized-join code. Unused by Nim’s live codegen (superseded by the host-side runner below) but ported for completeness.

srdatalog.ir.codegen.cuda.materialized.gen_materialized_runner(node: srdatalog.ir.mir.types.ExecutePipeline, db_type_name: str) → str[source]

Emit a host-side Thrust executor for a materialized-join pipeline.

Produces struct JitRunner_<rule_name> with an execute(DB&, iter) static method that walks Scan → ProbeJoin+ → GatherColumn* → InsertInto using Thrust primitives.

Matches Nim’s genMaterializedRunner byte-for-byte (modulo the usual clang-format whitespace).

srdatalog.ir.codegen.cuda.materialized.is_materialized_pipeline(ops: list[srdatalog.ir.mir.types.MirNode]) → bool[source]

True iff the pipeline contains a ProbeJoin — which switches dispatch to the materialized runner.