diff --git a/Makefile b/Makefile index be56103..e86c541 100644 --- a/Makefile +++ b/Makefile @@ -2,11 +2,14 @@ PYTHON ?= python3 REPO_ROOT := $(CURDIR) export PYTHONPATH := $(REPO_ROOT)/python:$(REPO_ROOT)/compiler$(if $(PYTHONPATH),:$(PYTHONPATH)) -.PHONY: check examples native-sdk-check package package-test python-check repository-check shell-check test validate +.PHONY: check examples native-sdk-check package package-test python-check repository-check shell-check static-check test validate python-check: $(PYTHON) -m compileall -q compiler/netra_compiler python/netra_kernel tools/compiler tools/ci tests/compiler tests/package +static-check: + $(PYTHON) -m ruff check compiler/netra_compiler python/netra_kernel tools tests --select F401,F811,F821,F841 + test: $(PYTHON) -m unittest discover -s tests/compiler -v @@ -29,7 +32,7 @@ shell-check: git ls-files -z '*.sh' | xargs -0 -r -n1 bash -n git diff --check -check: python-check test package-test repository-check native-sdk-check validate examples shell-check +check: python-check static-check test package-test repository-check native-sdk-check validate examples shell-check package: PYTHON=$(PYTHON) bash tools/ci/build_package.sh diff --git a/README.md b/README.md index 7389432..62a17ef 100644 --- a/README.md +++ b/README.md @@ -110,6 +110,44 @@ stream. Direct fixed launches remain available when graph capture is disabled. The installed SDK exports `libnetra_engine.so.1`, ``, CMake and pkg-config metadata, plus a framework-neutral `netra_kernel.Engine` binding. +## Repository layout + +The public integration surface is separated from target-specific kernels and +inference-engine adapters: + +```text +netra-kernel/ +├── compiler/netra_compiler/ # AOT compiler, IR, planners, frontends, backends +│ ├── backends/gfx950/ # Target-specific compiler implementation +│ └── library_data/ # Packaged kernels, manifests, and schemas +├── python/netra_kernel/ # Bundle, runtime, CLI, and adapter APIs +│ ├── runtime/ # Framework-neutral Python FFI +│ ├── sglang/ # SGLang adapter and startup plugin only +│ └── assets/ # Deployment profiles and packaged overlays +├── include/netra/engine.h # Stable framework-neutral C ABI +├── runtime/ # Native runtime implementations +├── examples/reference_engine/ # Minimal C consumer with no framework dependency +├── kernels/ # Architecture-specific kernel sources +├── manifests/ # Model, deployment, and tactic contracts +├── schemas/ # Public JSON schemas +├── docker/sglang/ # Thin-plugin and accepted-source images +├── tools/ # Build, compiler, benchmark, and CI tooling +├── tests/ # Compiler, package, and GPU validation +└── docs/ # Integration guides and acceptance evidence +``` + +`netra-compiler` and `netra-kernel` are separate installable distributions. +The compiler emits a deterministic engine directory; `libnetra_engine.so.1` +loads and launches it; an optional adapter translates only the surrounding +inference-engine boundary. New engines integrate through the +`netra.inference_engines` entry-point group without importing SGLang code. + +The top-level `kernels/gfx950`, `manifests`, and `schemas` paths point to the +same canonical data packaged in the wheels. This preserves convenient +source-tree commands without maintaining duplicate catalogs that can drift. +Architecture-specific harnesses, profiling scripts, and historical results +remain outside the public SDK surface. + ## Supported today | Capability | Current support | diff --git a/compiler/netra_compiler/engine.py b/compiler/netra_compiler/engine.py index 9d84783..6047d78 100644 --- a/compiler/netra_compiler/engine.py +++ b/compiler/netra_compiler/engine.py @@ -1,7 +1,6 @@ from __future__ import annotations import hashlib -import json import shutil from collections.abc import Mapping from pathlib import Path @@ -16,7 +15,7 @@ from .memory import BufferRequest, checked_mul, plan_buffers from .planner import Plan, plan_graph from .profiles import ShapeProfile -from .types import canonical_json, stable_hash, write_json +from .types import stable_hash, write_json FORMAT_VERSION = "netra-engine-1" diff --git a/compiler/netra_compiler/tactics.py b/compiler/netra_compiler/tactics.py index d668aef..4559759 100644 --- a/compiler/netra_compiler/tactics.py +++ b/compiler/netra_compiler/tactics.py @@ -2,7 +2,6 @@ from collections.abc import Mapping from dataclasses import dataclass -from pathlib import Path from typing import Any, Iterable from .contracts import GoldenKernelContract, KernelArgument, KernelContract, Launch, Workspace diff --git a/docs/notes/gfx1151-causal-conv-stream64-2026-07-29.md b/docs/notes/gfx1151-causal-conv-stream64-2026-07-29.md index 274368b..01f843b 100644 --- a/docs/notes/gfx1151-causal-conv-stream64-2026-07-29.md +++ b/docs/notes/gfx1151-causal-conv-stream64-2026-07-29.md @@ -37,7 +37,7 @@ The raw path delivered the expected ~1.07 s host benefit, but pair-b changed its The retained raw ASM and harness are: - `kernels/gfx1151/gdn/experiments/causal_conv1d_stream64_gfx1151.s` -- `kernels/gfx1151/gdn/experiments/causal_conv1d_state_update_gfx1151.s` +- `kernels/gfx1151/gdn/causal_conv1d_state_update_gfx1151.s` - `harness/gfx1151/gdn/causal_conv1d_stream64_launcher.hip` - `tools/benchmark/benchmark_causal_conv1d_stream64.py` diff --git a/kernels/gfx1151/gdn/experiments/causal_conv1d_state_update_gfx1151.s b/kernels/gfx1151/gdn/experiments/causal_conv1d_state_update_gfx1151.s deleted file mode 100644 index 0c6af3b..0000000 --- a/kernels/gfx1151/gdn/experiments/causal_conv1d_state_update_gfx1151.s +++ /dev/null @@ -1,92 +0,0 @@ -// SPDX-License-Identifier: MIT -// Raw gfx1151 state writeback for T=8192,D=8192,W=4 BF16 causal convolution. -.amdgcn_target "amdgcn-amd-amdhsa--gfx1151" -.amdhsa_code_object_version 6 -.text -.protected causal_conv1d_state_update_gfx1151 -.globl causal_conv1d_state_update_gfx1151 -.p2align 8 -.type causal_conv1d_state_update_gfx1151,@function -causal_conv1d_state_update_gfx1151: -s_load_b128 s[4:7], s[0:1], 0 -s_load_b64 s[8:9], s[0:1], 16 -s_waitcnt lgkmcnt(0) -v_mov_b32_e32 v8, 0 -global_load_dword v9, v8, s[8:9] -s_lshl_b32 s10, s2, 8 -v_lshlrev_b32_e32 v1, 1, v0 -v_add_nc_u32_e32 v1, s10, v1 -v_add_nc_u32_e32 v2, 0x07ff4000, v1 -global_load_ushort v3, v2, s[4:5] -v_add_nc_u32_e32 v2, 16384, v2 -global_load_ushort v4, v2, s[4:5] -v_add_nc_u32_e32 v2, 16384, v2 -global_load_ushort v5, v2, s[4:5] -s_waitcnt vmcnt(0) -v_readfirstlane_b32 s11, v9 -s_mul_i32 s11, s11, 49152 -s_add_u32 s6, s6, s11 -s_addc_u32 s7, s7, 0 -s_mul_i32 s12, s2, 768 -s_waitcnt_depctr 0 -v_mul_lo_u32 v6, v0, 6 -v_add_nc_u32_e32 v6, s12, v6 -global_store_short v6, v3, s[6:7] -global_store_short v6, v4, s[6:7] offset:2 -global_store_short v6, v5, s[6:7] offset:4 -s_endpgm -.section .rodata,"a",@progbits -.p2align 6,0 -.amdhsa_kernel causal_conv1d_state_update_gfx1151 -.amdhsa_group_segment_fixed_size 0 -.amdhsa_private_segment_fixed_size 0 -.amdhsa_kernarg_size 24 -.amdhsa_user_sgpr_count 2 -.amdhsa_user_sgpr_kernarg_segment_ptr 1 -.amdhsa_wavefront_size32 1 -.amdhsa_enable_private_segment 0 -.amdhsa_system_sgpr_workgroup_id_x 1 -.amdhsa_system_sgpr_workgroup_id_y 0 -.amdhsa_system_sgpr_workgroup_id_z 0 -.amdhsa_system_vgpr_workitem_id 0 -.amdhsa_next_free_vgpr 10 -.amdhsa_next_free_sgpr 13 -.amdhsa_reserve_vcc 1 -.amdhsa_float_denorm_mode_32 3 -.amdhsa_float_denorm_mode_16_64 3 -.amdhsa_dx10_clamp 1 -.amdhsa_ieee_mode 1 -.amdhsa_workgroup_processor_mode 1 -.amdhsa_memory_ordered 1 -.amdhsa_forward_progress 1 -.end_amdhsa_kernel -.text -.Lend: .size causal_conv1d_state_update_gfx1151,.Lend-causal_conv1d_state_update_gfx1151 -.amdgpu_metadata ---- -amdhsa.kernels: - - .args: - - { .name: x, .offset: 0, .size: 8, .value_kind: global_buffer, .address_space: global, .actual_access: read_only } - - { .name: state, .offset: 8, .size: 8, .value_kind: global_buffer, .address_space: global, .actual_access: write_only } - - { .name: cache_index, .offset: 16, .size: 8, .value_kind: global_buffer, .address_space: global, .actual_access: read_only } - .group_segment_fixed_size: 0 - .kernarg_segment_align: 8 - .kernarg_segment_size: 24 - .language: OpenCL C - .language_version: [2,0] - .max_flat_workgroup_size: 128 - .name: causal_conv1d_state_update_gfx1151 - .private_segment_fixed_size: 0 - .sgpr_count: 15 - .sgpr_spill_count: 0 - .symbol: causal_conv1d_state_update_gfx1151.kd - .uniform_work_group_size: 1 - .uses_dynamic_stack: false - .vgpr_count: 10 - .vgpr_spill_count: 0 - .wavefront_size: 32 - .workgroup_processor_mode: 1 -amdhsa.target: amdgcn-amd-amdhsa--gfx1151 -amdhsa.version: [1,2] -... -.end_amdgpu_metadata diff --git a/netra_compiler/__init__.py b/netra_compiler/__init__.py deleted file mode 100644 index e25120b..0000000 --- a/netra_compiler/__init__.py +++ /dev/null @@ -1,14 +0,0 @@ -"""Source-tree import shim for ``python -m netra_compiler.cli``. - -Installed packages use ``compiler/pyproject.toml`` directly. -""" - -from pathlib import Path - -__path__.append(str(Path(__file__).resolve().parents[1] / "compiler" / "netra_compiler")) - -from .contracts import KernelContract -from .engine import compile_engine - -__all__ = ["KernelContract", "compile_engine"] -__version__ = "0.2.0" diff --git a/pyproject.toml b/pyproject.toml index d346721..df3eee7 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -18,7 +18,7 @@ classifiers = [ [project.optional-dependencies] sglang = ["torch>=2.8"] -test = ["pytest>=8", "build>=1.2"] +test = ["pytest>=8", "build>=1.2", "ruff>=0.16"] [project.scripts] netra-kernel = "netra_kernel.kernel_cli:main" diff --git a/python/netra_kernel/cli.py b/python/netra_kernel/cli.py index f14046c..1f36dbb 100644 --- a/python/netra_kernel/cli.py +++ b/python/netra_kernel/cli.py @@ -77,7 +77,6 @@ def _doctor(bundle_path: Path | None) -> dict: expected_source_snapshot = None bundle_result = None if bundle_path is not None: - payload = load_bundle(bundle_path) integration = integration_configuration(bundle_path, "sglang") required = tuple(integration.get("required_modules", ())) expected_source_snapshot = integration.get("source_snapshot") diff --git a/python/netra_kernel/native.py b/python/netra_kernel/native.py index cf6adbb..fd95967 100644 --- a/python/netra_kernel/native.py +++ b/python/netra_kernel/native.py @@ -10,7 +10,7 @@ import ctypes import ctypes.util import os -from importlib.resources import as_file, files +from importlib.resources import files from pathlib import Path from typing import Iterator diff --git a/tests/compiler/test_engine_manifest.py b/tests/compiler/test_engine_manifest.py index ebcfe7b..5206ec1 100644 --- a/tests/compiler/test_engine_manifest.py +++ b/tests/compiler/test_engine_manifest.py @@ -1,6 +1,5 @@ from __future__ import annotations -import hashlib import json import tempfile import unittest diff --git a/tests/gfx1151/test_netra_mxfp4_sgl_decode.py b/tests/gfx1151/test_netra_mxfp4_sgl_decode.py index 72c68cf..d73b015 100644 --- a/tests/gfx1151/test_netra_mxfp4_sgl_decode.py +++ b/tests/gfx1151/test_netra_mxfp4_sgl_decode.py @@ -3,7 +3,6 @@ import ctypes import json -import math import os from pathlib import Path diff --git a/tests/gfx1151/test_netra_mxfp4_sgl_linear.py b/tests/gfx1151/test_netra_mxfp4_sgl_linear.py index b4829b4..102e7db 100644 --- a/tests/gfx1151/test_netra_mxfp4_sgl_linear.py +++ b/tests/gfx1151/test_netra_mxfp4_sgl_linear.py @@ -93,12 +93,12 @@ def main() -> None: if status: raise RuntimeError(f"launch failed: {status}") torch.cuda.synchronize() + output = output_groups.view(-1, n)[:m].to(torch.bfloat16) if os.environ.get("NETRA_ARGS_DEBUG") == "1": print(output.view(torch.int32).cpu()[0, :8].tolist()) return weight64 = dequant(packed_cpu, scales_cpu) - output = output_groups.view(-1, n)[:m].to(torch.bfloat16) reference = activation.cpu().to(torch.float64) @ weight64.t() ref_bf16 = reference.to(torch.bfloat16) got = output.cpu() diff --git a/tools/benchmark/benchmark_extend_attention_tiles.py b/tools/benchmark/benchmark_extend_attention_tiles.py index 400ef57..4cfb585 100755 --- a/tools/benchmark/benchmark_extend_attention_tiles.py +++ b/tools/benchmark/benchmark_extend_attention_tiles.py @@ -96,8 +96,11 @@ def call() -> None: "max_abs_vs_64x64": max_abs, "normalized_l2_vs_64x64": normalized_l2, }) - del q, k, v, k_buffer, v_buffer, reference - + # Drop the closure before its captured tensors so the next prefix does + # not temporarily retain two full attention fixtures. + call = output = None + q = k = v = k_buffer = v_buffer = reference = None + torch.cuda.empty_cache() module._get_block_sizes_for_extend_attention = original_blocks result = { "target": "gfx1151", diff --git a/tools/benchmark/benchmark_qkvzba_split_copy.py b/tools/benchmark/benchmark_qkvzba_split_copy.py index 4b678f9..a56f264 100755 --- a/tools/benchmark/benchmark_qkvzba_split_copy.py +++ b/tools/benchmark/benchmark_qkvzba_split_copy.py @@ -154,7 +154,10 @@ def launch_triton() -> None: "bit_mismatches_qkv_z_b_a": mismatches, } ) - del raw_tensors, triton_tensors + # The launch closures capture the tensors; release them first so + # empty_cache can return storage before the next token shape. + launch_raw = launch_triton = None + raw_tensors = triton_tensors = None torch.cuda.empty_cache() report = { diff --git a/tools/ci/check.sh b/tools/ci/check.sh deleted file mode 100755 index 316b5d5..0000000 --- a/tools/ci/check.sh +++ /dev/null @@ -1,21 +0,0 @@ -#!/usr/bin/env bash -set -euo pipefail - -script_dir=$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd) -repo_root=$(cd -- "$script_dir/../.." && pwd) -python_bin=${PYTHON:-python3} - -export PYTHONPATH="$repo_root/compiler${PYTHONPATH:+:$PYTHONPATH}" -cd "$repo_root" - -"$python_bin" -m compileall -q compiler/netra_compiler tools/compiler tools/ci tests/compiler -"$python_bin" -m unittest discover -s tests/compiler -v -"$python_bin" tools/ci/check_repository.py -"$python_bin" tools/compiler/validate_gfx950_tactic_catalog.py -"$python_bin" tools/ci/compile_examples.py - -while IFS= read -r -d '' shell_file; do - bash -n "$shell_file" -done < <(git ls-files -z '*.sh') - -git diff --check diff --git a/tools/ci/smoke_package.py b/tools/ci/smoke_package.py index 0bed006..43944a6 100644 --- a/tools/ci/smoke_package.py +++ b/tools/ci/smoke_package.py @@ -6,7 +6,6 @@ import argparse import os import subprocess -import sys import tempfile import venv from pathlib import Path diff --git a/tools/compiler/netra_compile.py b/tools/compiler/netra_compile.py deleted file mode 100755 index 11746d3..0000000 --- a/tools/compiler/netra_compile.py +++ /dev/null @@ -1,13 +0,0 @@ -#!/usr/bin/env python3 -from __future__ import annotations - -import sys -from pathlib import Path - -ROOT = Path(__file__).resolve().parents[2] -sys.path.insert(0, str(ROOT / "compiler")) - -from netra_compiler.cli import main - -if __name__ == "__main__": - raise SystemExit(main()) diff --git a/tools/compiler/validate_engine.py b/tools/compiler/validate_engine.py deleted file mode 100755 index 68da3e6..0000000 --- a/tools/compiler/validate_engine.py +++ /dev/null @@ -1,11 +0,0 @@ -#!/usr/bin/env python3 -from __future__ import annotations - -import sys -from pathlib import Path - -ROOT = Path(__file__).resolve().parents[2] -sys.path.insert(0, str(ROOT / "compiler")) -from netra_compiler.cli import main - -if __name__ == "__main__": raise SystemExit(main(["validate", "--engine", *sys.argv[1:], "--static"])) diff --git a/tools/profiling/rocprof_attach.py b/tools/profiling/rocprof_attach.py index 8efa56d..719d2e2 100755 --- a/tools/profiling/rocprof_attach.py +++ b/tools/profiling/rocprof_attach.py @@ -10,7 +10,6 @@ import ctypes import os import signal -import sys import time DEFAULT_LIBRARY = (