Skip to content

Add SSA compiler pipeline and unified backends#164

Merged
voltjia merged 18 commits into
masterfrom
ssa-compiler-with-multibackend
Jul 16, 2026
Merged

Add SSA compiler pipeline and unified backends#164
voltjia merged 18 commits into
masterfrom
ssa-compiler-with-multibackend

Conversation

@whjthu

@whjthu whjthu commented Jun 28, 2026

Copy link
Copy Markdown
Contributor

Summary

This PR introduces an SSA-first compiler pipeline and a unified multi-backend architecture for NineToothed. The new path is integrated with the existing jit, make, aot, and build entry points, while keeping frontend lowering, compiler passes, backend emission, and runtime materialization as separate layers.

Highlights

  • Add structured kernel, layout, launch, and SSA IRs with a configurable pass pipeline.
  • Add Triton, CUDA, and TileLang backend emission and materialization paths, together with caching, specialization, autotuning, and reloadable artifacts.
  • Add Triton AOT multi-context support and fresh-context reload coverage.
  • Expand tests across lowering, passes, backend boundaries, runtime behavior, and AOT materialization.

pytest output:

$ python3 -m pytest --doctest-modules --junitxml=junit/test-results.xml --cov=ninetoothed --cov-report=xml --cov-report=html
collected 353 items
156 passed, 197 skipped in 22.61s

All current GitHub Actions checks pass, including pytest, Ruff, distribution build, metadata, and documentation.

@whjthu
whjthu force-pushed the ssa-compiler-with-multibackend branch 4 times, most recently from ad873c2 to a8249a9 Compare June 28, 2026 08:21
@whjthu
whjthu force-pushed the ssa-compiler-with-multibackend branch from a8249a9 to 9a4ba7a Compare June 28, 2026 09:12
@voltjia
voltjia force-pushed the ssa-compiler-with-multibackend branch from 437f718 to 9a4ba7a Compare June 29, 2026 05:18
@voltjia

This comment was marked as resolved.

@voltjia

This comment was marked as resolved.

Comment thread src/ninetoothed/backends/base.py Outdated
Comment thread docs/design/ir_design_zh.md Outdated
Comment thread docs/design/multi_backend_ir.md Outdated
Comment thread docs/design/README.md Outdated
Comment thread docs/design/ssa_ir_lowering_pipeline_zh.md Outdated
Comment thread tests/test_ssa_pass_pipeline.py
Comment thread src/ninetoothed/__init__.py
Comment thread src/ninetoothed/aot.py
Comment thread src/ninetoothed/generation.py Outdated
Comment thread src/ninetoothed/compiler/passes.py

@voltjia voltjia left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Submitting pending inline review comments.

Comment thread src/ninetoothed/backends/triton.py Outdated
Comment thread src/ninetoothed/ssa_passes.py Outdated
Comment thread src/ninetoothed/ssa_passes.py Outdated
Comment thread src/ninetoothed/ssa_passes.py Outdated
Comment thread src/ninetoothed/lowering.py Outdated
Comment thread src/ninetoothed/backends/emitters/common.py Outdated
Comment thread tests/test_ssa_first_backend_lowering.py Outdated
Comment thread tests/test_backend_registry.py Outdated
Comment thread tests/test_ssa_first_backend_lowering.py Outdated
Comment thread scripts/check_contributing_style.py Outdated

voltjia commented Jul 2, 2026

Copy link
Copy Markdown
Collaborator

关于模块和命名,建议把这个 PR 收敛成下面这套确定方案。核心原则是:模块负责承载上下文,public-ish 名称去掉重复的 IR / SSA / Backend 前后缀;因为 PR 尚未合入,不需要保留旧名 alias;已知为 legacy 的 ProgramIR 也不要进入这个 PR。

模块布局

src/ninetoothed/
  ir/
    __init__.py
    kernel.py
    ssa.py
  frontend/
    __init__.py
    python.py
  compiler/
    __init__.py
    passes.py
  backends/
    __init__.py
    core.py
    registry.py
    triton.py
    cuda.py
    tilelang.py
    tvm.py
    emitters/
      __init__.py
      ssa.py
  lowering.py

模块单复数建议也固定下来:表示一个领域或命名空间的 package 用单数/概念名,比如 irfrontendcompiler;表示一组可扩展实现或集合的 package 用复数,比如 backendspassesemitters。具体文件默认用单数名,比如 kernel.pyssa.pyregistry.pytriton.py

IR 命名

TensorTypeIR        -> TensorSpec
LaunchIR            -> Launch
KernelIR            -> Kernel
SSATypeIR           -> ssa.Type
SSAValueIR          -> ssa.Value
SSAOperationIR      -> ssa.Operation
SSABlockIR          -> ssa.Block
SSAProgramIR        -> ssa.Program

TensorTypeIR 建议改成 TensorSpec,而不是 TensorTensorType:它现在承载的不只是 element type / shape / layout,还包含参数名、source/view、constexpr、jagged 等边界信息,更像 kernel 参数规格。ssa.Operationssa.Op 更适合作为公共结构名,避免缩写进入上层 API。

下面这些 legacy / transitional kernel-level op IR 不建议保留在这个 PR 里,统一删掉或并入 SSA 入口:

ProgramIR
ExprIR
ElementwiseBinaryOpIR
ElementwiseAssignOpIR
AxisReductionAssignOpIR
RowwiseAssignOpIR
FillOpIR
CopyOpIR
ReductionOpIR
MatmulOpIR
FlashAttentionOpIR
TransposeOpIR
program_to_ssa

Frontend / render 命名

application_to_ssa   -> frontend.python.from_application
source_to_ssa        -> frontend.python.from_source
render_ssa_program   -> ir.ssa.render
SSALoweringError     -> frontend.python.LoweringError

这里的重点是避免全局函数名重复携带 ssa / python / lowering 上下文;调用点通过 module 已经能读出来。

Pass 命名

SSAPipelineSpec                    -> PipelineSpec
SSAPassContext                     -> Context
SSAPass                            -> Pass
SSAPassDescriptor                  -> Descriptor
SSAPassRegistry                    -> Registry
SSAPassPipeline                    -> Pipeline
CanonicalizeSSAPass                -> Canonicalize
AnalyzeSSAEffectsPass              -> AnalyzeEffects
DecomposeLinalgPass                -> DecomposeLinalg
SelectSchedulePass                 -> SelectSchedule
BackendScheduleOptimizationPass    -> OptimizeSchedule
BackendMemoryScopesLoweringPass    -> LowerMemoryScopes
BackendIntrinsicsLoweringPass      -> LowerIntrinsics

对应 helper 也建议收短:

create_default_ssa_pass_registry   -> create_default_registry
registered_ssa_passes              -> registered
default_ssa_pipeline_spec          -> default_spec
default_ssa_pipeline               -> default_pipeline
build_ssa_pipeline                 -> build
lower_ssa_for_backend              -> lower_for_target
annotate_ssa_operations            -> annotate_operations

这些名字在 compiler.passes 里已经有足够上下文,继续重复 SSA / Pass / Backend 会让未来贡献者写调用点时很噪。

Backend 命名

BackendName        -> Target
BackendOptions     -> Options
BackendCapability  -> Capability
BackendArtifact    -> Artifact
BackendRegistry    -> Registry
Backend            -> Backend

Backend 本身可以保留,因为这是对外抽象的核心名;其它类型放在 backends namespace 后不需要再重复 Backend

后端入口建议从 lower 改成 emit,以免和高层 lowering 混淆:

backends.lower                    -> backends.emit
backends/ssa_unified.py           -> backends/emitters/ssa.py
lower_unified_ssa_artifact        -> emit

这样外层仍然可以是 lowering.lower(...) 或 compiler pipeline 的 lower,而 backend 阶段语义更明确:把已经 lowered/selected 的 SSA program emit 成目标 artifact。

@whjthu
whjthu requested a review from voltjia July 3, 2026 07:18
@voltjia

voltjia commented Jul 7, 2026

Copy link
Copy Markdown
Collaborator

建议把新增 IR/backend/compiler 数据结构的构造方式统一收紧一下,避免这个 PR 还没合入时就固化一批容易错位的位置参数 API。

总体原则:

  • 新增 dataclass 建议全部使用 kw_only=True
  • 涉及张量属性的字段顺序统一对齐现有 Tensorndim, shape, dtype, jagged_dim, ...
  • name 类字段不要放在位置参数最前面;保留其语义,但全部 keyword-only。
  • 不需要兼容旧写法,因为 PR 尚未合入。

具体建议:

@dataclass(frozen=True, kw_only=True)
class TensorSpec:
    ndim: int
    shape: tuple[str, ...] = ()
    dtype: str | None = None
    jagged_dim: int | None = None
    constexpr: bool = False
    name: str
    attrs: Mapping[str, Any] = field(default_factory=dict)

TensorSpec.name 不是普通 info,它现在用于 frontend 参数绑定、SSA value 生成和 backend 参数/输出匹配,应继续必填;但它不应该作为第一个位置参数。attrs["source_name"] 才更像 provenance/debug 信息。

ssa.Type 也涉及张量属性,字段顺序建议改成:

@dataclass(frozen=True, kw_only=True)
class Type:
    kind: str
    shape: tuple[str, ...] = ()
    dtype: str | None = None
    attrs: Mapping[str, Any] = field(default_factory=dict)

Type 描述 SSA value 的类型,不需要 name,也不建议加 ndim;维度可由 shape 推出。名字应在 ssa.Value.name 上。

其它 SSA IR 也建议 keyword-only:

@dataclass(frozen=True, kw_only=True)
class Value:
    name: str
    type: Type

@dataclass(frozen=True, kw_only=True)
class Operation:
    opcode: str
    operands: tuple[str, ...] = ()
    results: tuple[Value, ...] = ()
    attrs: Mapping[str, Any] = field(default_factory=dict)
    regions: tuple["Block", ...] = ()

@dataclass(frozen=True, kw_only=True)
class Block:
    name: str = "entry"
    args: tuple[Value, ...] = ()
    operations: tuple[Operation, ...] = ()

@dataclass(frozen=True, kw_only=True)
class Program:
    kind: str
    inputs: tuple[Value, ...] = ()
    outputs: tuple[Value, ...] = ()
    blocks: tuple[Block, ...] = ()
    metadata: Mapping[str, Any] = field(default_factory=dict)

Launch / Kernel 也建议全 keyword-only。Kernel.kernel_name 是功能字段,会参与文件名、函数名和 entrypoint 生成;Launch.name 目前更偏 launch metadata,但两者都不适合继续作为位置参数暴露。

@dataclass(frozen=True, kw_only=True)
class Launch:
    name: str
    args: tuple[str, ...] = ()
    grid: str | None = None

@dataclass(frozen=True, kw_only=True)
class Kernel:
    kernel_name: str
    source: str
    source_path: str | None = None
    source_language: str = "triton"
    entrypoint: str | None = None
    launch: Launch | None = None
    tensors: tuple[TensorSpec, ...] = ()
    compiler_options: Mapping[str, Any] = field(default_factory=dict)
    metadata: Mapping[str, Any] = field(default_factory=dict)
    ssa: Program | None = None

Backend/compiler 侧的 dataclass 也建议同样收紧:

  • Options
  • Capability
  • Artifact
  • PipelineSpec
  • Context
  • Descriptor
  • _TensorInfo
  • _Target
  • _EmitContext

其中 _TensorInfo 同样按张量属性顺序整理:

@dataclass(frozen=True, kw_only=True)
class _TensorInfo:
    ndim: int = 1
    shape: tuple[str, ...] = ()
    dtype: str = "float32"
    name: str
    source_name: str | None = None
    source_shape: tuple[str, ...] = ()
    source_strides: tuple[str, ...] = ()
    view_linear_offset: str | None = None
    view_mask: str | None = None
    attrs: Mapping[str, Any] | None = None

调用点统一改成 keyword 形式,例如:

TensorSpec(ndim=1, shape=("n",), dtype="float32", name="x")
ssa.Type(kind="tensor", shape=("n",), dtype="float32")
ssa.Value(name="x", type=type_)
ssa.Operation(opcode="mem.store", operands=(value, "out"))
Kernel(kernel_name="add", source=source, tensors=tensors, ssa=program)
Launch(name="launch_add", args=("x", "out"), grid="lambda meta: (1,)")

最后建议补几条测试,防止后续回退:

  • 旧式位置参数构造 dataclass 会抛 TypeError
  • TensorSpec.name 仍参与 frontend/backend binding。
  • attrs["source_name"] 只作为 provenance/debug 信息,不参与 binding。
  • Kernel.kernel_name 仍影响 artifact 文件名、entrypoint 和生成源码函数名。

@voltjia

voltjia commented Jul 10, 2026

Copy link
Copy Markdown
Collaborator

这个 PR 目前已经引入了新的 SSA lowering / multi-backend 路径,但从现有主线入口来看,它还没有和原来的用户路径形成良好的集成。make / jit / build / AOT 这些入口现在大体仍然各自沿用旧路径;新路径更多是通过 lowering.lower 和 backend emitter 直接被测试。因此当前 CI 通过主要能说明两件事:旧路径没有明显回归,以及新 SSA/backend path 作为旁路 API 可以生成一些 artifact/source;但还不能说明用户通过原有 public API 就能可靠地使用这条新路径。

我做了一些临时验证来确认这个集成状态。在 ssh nvidia 上用 accelerator-dev/nvidia:latest,临时把 jit 接到 SSA Triton artifact/import/launch 路径后,旧的最小 JIT add 用例还不能直接跑通:一开始 generated launch 里缺 runtime shape 参数;临时补 shape/meta 参数绑定后,tests/test_add.py 仍然数值不等价,输出只写了少量位置;再测 tests/test_generation.py::test_non_int_constexpr,也遇到了 launch 参数数量不匹配。同时,开着这个临时 SSA JIT 开关跑 AOT 标量用例仍然通过,这也说明默认 AOT/build 类路径并没有因为 jit 的临时切换而被真正覆盖到。

所以这里的主要问题不是“应该在某个函数里改一行切换过去”,而是这套新路径目前还缺少和主线 API 的系统性集成与验证。建议在这个 PR 或后续 PR 中明确补上:

  1. 明确新 SSA path 在 public API 中的入口和返回语义,例如是 opt-in 的 source/artifact API,还是要兼容旧的 callable kernel handle。
  2. 如果要接入 jit,需要补齐旧 JIT 的 runtime 行为,包括 runtime tensor shape、meta parameter、constexpr/scalar 参数、launch 参数顺序,以及 arrangement/tile/view 的 index lowering 语义。
  3. 如果要接入 build / AOT,需要单独设计接入点,因为默认 build 现在走的是 caller=cuda 的 AOT 路径,不会自然被 jit 的改动覆盖。
  4. 增加面向主线入口的 integration tests,而不只是直接测 lowering.lower 或 backend emitter。至少需要覆盖 legacy 默认路径不变、opt-in SSA JIT 的最小 add/constexpr scalar、以及 opt-in build/AOT 的 artifact 或可执行路径。

换句话说,这条新路径本身可以继续演进,但目前还不宜认为它已经和原有 make / jit / build 用户路径连通并可用;需要把 integration contract 和验证矩阵补清楚。

Comment thread src/ninetoothed/backends/materializers/__init__.py
Comment thread src/ninetoothed/__init__.py
Comment thread tests/test_triton_runtime_auto_tuning.py
Comment thread src/ninetoothed/compiler/driver.py Outdated
voltjia added 2 commits July 16, 2026 17:10
* Add Triton AOT multi-context design

* Fix Triton AOT multi-context launches

* Cover fresh Triton AOT context guard

* Remove development process documents
* Address remaining SSA review comments

* Sort test imports for Ruff

* Remove redundant entrypoint assertions
@voltjia
voltjia self-requested a review July 16, 2026 09:13
@voltjia
voltjia merged commit b29a7e1 into master Jul 16, 2026
10 checks passed
@voltjia
voltjia deleted the ssa-compiler-with-multibackend branch July 16, 2026 09:26
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants