diff --git a/.agents/riscv_operator_workflow/README.md b/.agents/riscv_operator_workflow/README.md new file mode 100644 index 0000000000..544824abd9 --- /dev/null +++ b/.agents/riscv_operator_workflow/README.md @@ -0,0 +1,33 @@ +# TileLang-RISC-V Operator Workflow + +This directory contains the Codex-compatible operator workflow for TileLang-RISCV. + +## Entry Points + +- `operators.json`: source of truth for initial operator task context +- `validation_report.md`: validation results, SG2044 baseline status, and PR summary draft +- `.agents/tasks/riscv/*.md`: generated Codex task prompts +- `.agents/skills/riscv-operator-workflow/SKILL.md`: Codex workflow instructions +- `docs/get_started/RiscvOperatorAgentWorkflow.md`: user-facing workflow guide + +## Generate Tasks + +```bash +python maint/scripts/generate_riscv_operator_task.py --list +python maint/scripts/generate_riscv_operator_task.py --all +python maint/scripts/generate_riscv_operator_task.py --operator vector_add +``` + +## Validate + +```bash +bash maint/scripts/run_riscv_operator_workflow_validation.sh +``` + +The initial workflow validates `vector_add`, `reduce_sum`, and `matmul` through host checks, RISC-V artifact export, and QEMU functional smoke checks. + +## SG2044 Boundary + +No new SG2044 native run is claimed unless the workflow is actually rerun on SG2044 hardware. + +Existing SG2044 baseline data from `docs/get_started/BuildOnSG2044.md` is kept as historical baseline data. QEMU functional smoke checks do not replace SG2044 native RVV validation. diff --git a/.agents/riscv_operator_workflow/operators.json b/.agents/riscv_operator_workflow/operators.json new file mode 100644 index 0000000000..4724cbe6b7 --- /dev/null +++ b/.agents/riscv_operator_workflow/operators.json @@ -0,0 +1,177 @@ +{ + "workflow_version": 1, + "target": "riscv", + "target_alias": "linalg_riscv", + "validation_note": "The Codex operator workflow can be pre-validated without SG2044 hardware by using x86/arm host execution, RISC-V artifact export, and QEMU. Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available.", + "common_context": { + "riscv_flow": "TileLang -> MLIR Linalg -> MLIR vector/RVV lowering -> LLVM/RISC-V artifact -> host check -> QEMU smoke check -> optional SG2044 native replay", + "sg2044_baseline": [ + "Existing documented SG2044 native baseline from docs/get_started/BuildOnSG2044.md:", + "testing/python/riscv/test_riscv_target_parse.py: 2 passed", + "testing/python/riscv/test_riscv_mlir_codegen.py: 33 passed", + "testing/python/riscv/test_riscv_toolchain.py: 3 passed", + "testing/python/riscv/test_riscv_jit_runtime.py: 16 passed", + "This baseline was not rerun during the no-hardware workflow validation pass." + ], + "primary_docs": [ + ".agents/riscv_operator_workflow/README.md", + "docs/get_started/BuildOnSG2044.md", + "docs/get_started/targets.md" + ], + "reference_tests": [ + "testing/python/riscv/test_riscv_target_parse.py", + "testing/python/riscv/test_riscv_mlir_codegen.py", + "testing/python/riscv/test_riscv_toolchain.py", + "testing/python/riscv/test_riscv_jit_runtime.py", + "testing/python/riscv/test_riscv_examples.py", + "testing/python/riscv/test_riscv_operator_workflow_examples.py", + "testing/python/riscv/test_riscv_qemu_smoke.py" + ], + "common_failure_buckets": [ + "Generated TileLang/TIR source error", + "Unsupported scheduling or TileLang intrinsic for linalg_riscv", + "MLIR lowering or structured codegen limitation", + "LLVM/RISC-V toolchain setup issue", + "Runtime wrapper, host adapter, QEMU, or optional SG2044 environment issue", + "Test configuration or tolerance mismatch" + ] + }, + "operators": [ + { + "name": "vector_add", + "category": "elementwise", + "task_title": "Implement and validate vector_add for TileLang-RISCV", + "semantics": "For each i in [0, 8), compute C[i] = A[i] + B[i].", + "inputs": [ + "A: float32 tensor with shape (8,)", + "B: float32 tensor with shape (8,)" + ], + "outputs": [ + "C: float32 tensor with shape (8,)" + ], + "baseline_reference": [ + "PyTorch: lhs + rhs", + "NumPy: lhs + rhs" + ], + "baseline_implementation": [ + "PyTorch: expected = lhs + rhs", + "NumPy: expected = lhs + rhs" + ], + "current_example": "examples/riscv/example_vector_add.py", + "suggested_files": [ + "examples/riscv/example_vector_add.py", + "testing/python/riscv/test_riscv_examples.py", + "testing/python/riscv/test_riscv_operator_workflow_examples.py", + "testing/python/riscv/test_riscv_qemu_smoke.py" + ], + "schedule_strategy": [ + "Use a single serial loop over N.", + "Use a T.block with one spatial axis.", + "Avoid GPU-specific T.Kernel threading, shared memory, warp, or async-copy constructs for this first RISC-V path." + ], + "compile_commands": [ + "python examples/riscv/example_vector_add.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add", + "python examples/riscv/example_vector_add.py --run-host --output-dir /tmp/vector_add", + "python examples/riscv/example_vector_add.py --run-qemu --output-dir /tmp/vector_add" + ], + "test_commands": [ + "python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k vector_add", + "python -m pytest testing/python/riscv/test_riscv_qemu_smoke.py -q" + ], + "success_criteria": [ + "Host adapter output matches the PyTorch reference.", + "Artifact export produces one non-empty .s and one non-empty .o file.", + "QEMU smoke emits vector_add.qemu.elf and output matches NumPy reference when a runner is available." + ] + }, + { + "name": "reduce_sum", + "category": "reduction", + "task_title": "Implement and validate row-wise reduce_sum for TileLang-RISCV", + "semantics": "For each row i in [0, 4), compute B[i] = sum(A[i, k] for k in [0, 8)).", + "inputs": [ + "A: float32 tensor with shape (4, 8)" + ], + "outputs": [ + "B: float32 tensor with shape (4,)" + ], + "baseline_reference": [ + "PyTorch: data.sum(dim=1)", + "NumPy: data.sum(axis=1)" + ], + "baseline_implementation": [ + "PyTorch: expected = data.sum(dim=1)", + "NumPy: expected = data.sum(axis=1)" + ], + "current_example": "examples/riscv/example_reduce_sum.py", + "suggested_files": [ + "examples/riscv/example_reduce_sum.py", + "testing/python/riscv/test_riscv_examples.py", + "testing/python/riscv/test_riscv_operator_workflow_examples.py" + ], + "schedule_strategy": [ + "Use T.grid(ROWS, COLS) with one spatial row axis and one reduce axis.", + "Use T.init() to zero the output row before accumulation.", + "Keep the first validation shape small and static so qemu artifact generation has fixed memref sizes." + ], + "compile_commands": [ + "python examples/riscv/example_reduce_sum.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/reduce_sum", + "python examples/riscv/example_reduce_sum.py --run-host --output-dir /tmp/reduce_sum", + "python examples/riscv/example_reduce_sum.py --run-qemu --output-dir /tmp/reduce_sum" + ], + "test_commands": [ + "python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k reduce_sum" + ], + "success_criteria": [ + "Lowered MLIR contains a reduction-friendly structured form such as linalg.reduce or equivalent generic lowering.", + "Host adapter output matches the PyTorch row-sum reference.", + "Artifact export produces one non-empty .s and one non-empty .o file." + ] + }, + { + "name": "matmul", + "category": "matrix", + "task_title": "Implement and validate small matmul for TileLang-RISCV", + "semantics": "For M=2, N=4, K=3, compute C[i, j] = sum(A[i, k] * B[k, j] for k in [0, 3)).", + "inputs": [ + "A: float32 tensor with shape (2, 3)", + "B: float32 tensor with shape (3, 4)" + ], + "outputs": [ + "C: float32 tensor with shape (2, 4)" + ], + "baseline_reference": [ + "PyTorch: lhs @ rhs", + "NumPy: lhs @ rhs" + ], + "baseline_implementation": [ + "PyTorch: expected = lhs @ rhs", + "NumPy: expected = lhs @ rhs" + ], + "current_example": "examples/riscv/example_matmul.py", + "suggested_files": [ + "examples/riscv/example_matmul.py", + "testing/python/riscv/test_riscv_examples.py", + "testing/python/riscv/test_riscv_operator_workflow_examples.py" + ], + "schedule_strategy": [ + "Use T.grid(M, N, K) with two spatial axes and one reduce axis.", + "Use T.init() to zero C[i, j] before the K accumulation.", + "Start with plain TIR loops rather than GPU tile intrinsics; move to T.gemm only after the linalg_riscv path supports the needed pattern." + ], + "compile_commands": [ + "python examples/riscv/example_matmul.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/matmul", + "python examples/riscv/example_matmul.py --run-host --output-dir /tmp/matmul", + "python examples/riscv/example_matmul.py --run-qemu --output-dir /tmp/matmul" + ], + "test_commands": [ + "python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k matmul" + ], + "success_criteria": [ + "Lowered MLIR keeps the matmul semantics visible in structured lowering or produces an equivalent loop form.", + "Host adapter output matches the PyTorch matmul reference.", + "Artifact export produces one non-empty .s and one non-empty .o file." + ] + } + ] +} diff --git a/.agents/riscv_operator_workflow/validation_report.md b/.agents/riscv_operator_workflow/validation_report.md new file mode 100644 index 0000000000..6d43964431 --- /dev/null +++ b/.agents/riscv_operator_workflow/validation_report.md @@ -0,0 +1,228 @@ +# TileLang-RISCV Operator Workflow Validation Report + +This report records the initial Codex operator workflow validation for issue +`RuyiAI-Stack/tilelang-riscv#1`. + +## Scope + +- Host used for this report: x86_64 WSL Ubuntu with LLVM/MLIR 18 and QEMU user mode. +- Native SG2044 status: no new SG2044 run was performed in this validation pass because no SG2044 hardware is available. +- Existing SG2044 data: the documented SG2044 baseline from `docs/get_started/BuildOnSG2044.md` is included below for context. +- RISC-V execution coverage: QEMU user-mode execution of freestanding RISC-V ELFs for all initial operators. +- Workflow target: generate Codex task files, complete three representative operators, compile through the RISC-V lowering path, and run correctness checks. + +The current QEMU validation path checks functional execution of small freestanding RISC-V ELFs, primarily targeting the `rv64gc` correctness path. It does not claim RVV-specific instruction coverage or performance validation. + +## Environment + +- Source checkout: `/home/devcontainers/codex-work/tilelang-riscv` +- Python environment: `/home/devcontainers/codex-work/tilelang-riscv/.venv-riscv` +- LLVM root: `/usr/lib/llvm-18` +- Required exports: + +```bash +source .venv-riscv/bin/activate +export TVM_FFI_DISABLE_TORCH_C_DLPACK=1 +export TILELANG_RISCV_LLVM_ROOT=/usr/lib/llvm-18 +export Z3_ROOT=/usr +``` + +## Workflow Artifacts + +- Operator catalog: `.agents/riscv_operator_workflow/operators.json` +- Task generator: `maint/scripts/generate_riscv_operator_task.py` +- Generated tasks: + - `.agents/tasks/riscv/vector_add.md` + - `.agents/tasks/riscv/reduce_sum.md` + - `.agents/tasks/riscv/matmul.md` +- Repository-local skill: `.agents/skills/riscv-operator-workflow/SKILL.md` +- Workflow guide: `docs/get_started/RiscvOperatorAgentWorkflow.md` +- One-command validation runner: `maint/scripts/run_riscv_operator_workflow_validation.sh` +- Validation runner controls: `TILELANG_RISCV_RUN_QEMU`, `TILELANG_RISCV_OPERATORS`, `TILELANG_RISCV_PYTEST_ARGS`, `TILELANG_RISCV_RUN_PYTEST`, `TILELANG_RISCV_VALIDATION_OUTPUT_ROOT`, `TILELANG_RISCV_RUN_HOST`, and `TILELANG_RISCV_EMIT_ARTIFACTS` + +## Operator Results + +| Operator | Category | Implementation | Host correctness | Artifact export | QEMU correctness | Failure bucket | +| --- | --- | --- | --- | --- | --- | --- | +| `vector_add` | elementwise | `examples/riscv/example_vector_add.py` | passed | `.mlir`, `.ll`, `.s`, `.o` passed | passed | none | +| `reduce_sum` | reduction | `examples/riscv/example_reduce_sum.py` | passed | `.mlir`, `.ll`, `.s`, `.o` passed | passed | none | +| `matmul` | matrix/tensor | `examples/riscv/example_matmul.py` | passed | `.mlir`, `.ll`, `.s`, `.o` passed | passed | none | + +## Existing SG2044 Baseline Data + +`docs/get_started/BuildOnSG2044.md` already records the following validated native SG2044 bring-up results: + +| Test | Documented SG2044 result | +| --- | --- | +| `testing/python/riscv/test_riscv_target_parse.py` | `2 passed` | +| `testing/python/riscv/test_riscv_mlir_codegen.py` | `33 passed` | +| `testing/python/riscv/test_riscv_toolchain.py` | `3 passed` | +| `testing/python/riscv/test_riscv_jit_runtime.py` | `16 passed` | + +These are existing repository baseline results. They were not rerun for this operator workflow validation because no SG2044 hardware is available in the current environment. + +## SG2044 Native Validation Status + +No new SG2044 native run was performed because SG2044 hardware was not available in this environment. + +The existing SG2044 baseline from `docs/get_started/BuildOnSG2044.md` is kept as historical baseline data. It was not rerun as part of this submission. + +The no-hardware validation in this submission covers: + +- host checks +- RISC-V artifact export +- QEMU functional smoke checks + +It does not replace SG2044 native RVV validation. + +## Issue Requirement Audit + +| Requirement | Evidence | +| --- | --- | +| Codex-compatible workflow | `.agents/skills/riscv-operator-workflow/SKILL.md`, `.agents/riscv_operator_workflow/operators.json`, and `docs/get_started/RiscvOperatorAgentWorkflow.md` | +| Task contexts include semantics, references, baseline implementation, files, schedule, compile/test commands, and success criteria | Generated task files in `.agents/tasks/riscv/`; enforced by `testing/python/riscv/test_riscv_operator_workflow.py` | +| Simple task generation flow for a given operator | `maint/scripts/generate_riscv_operator_task.py --operator ` and `--all` | +| 2 to 3 validation operators covering different compute patterns | `vector_add` for elementwise, `reduce_sum` for reduction, `matmul` for matrix/tensor compute | +| Compilation flow and basic correctness tests for each operator | `testing/python/riscv/test_riscv_examples.py`, `testing/python/riscv/test_riscv_operator_workflow_examples.py`, `testing/python/riscv/test_riscv_qemu_smoke.py`, and per-operator manual commands below | +| Results, failure logs, and failure reasons recorded | This report's operator table, existing SG2044 baseline, command results, and failure log | +| Reusable foundation for scaling to more operators | Catalog-driven generator, catalog-driven validation runner, and workflow test coverage | + +## Commands And Results + +Task generation: + +```bash +bash maint/scripts/run_riscv_operator_workflow_validation.sh +# generates tasks, runs testing/python/riscv, and runs the three manual operator checks +# 85 passed, 26 skipped +# vector_add host check passed +# vector_add qemu check passed +# reduce_sum host check passed +# reduce_sum qemu check passed +# matmul host check passed +# matmul qemu check passed +``` + +Individual task generation: + +```bash +python maint/scripts/generate_riscv_operator_task.py --list +# matmul +# reduce_sum +# vector_add + +python maint/scripts/generate_riscv_operator_task.py --all +# .agents/tasks/riscv/matmul.md +# .agents/tasks/riscv/reduce_sum.md +# .agents/tasks/riscv/vector_add.md +``` + +RISC-V backend tests: + +```bash +python -m pytest testing/python/riscv/test_riscv_operator_workflow.py -q +# 7 passed on Linux/WSL +# 5 passed, 2 skipped on Windows where bash validation is not available as a normal shell + +python -m pytest testing/python/riscv/test_riscv_target_parse.py -q +# 2 passed + +python -m pytest testing/python/riscv/test_riscv_toolchain.py -q +# 3 passed + +python -m pytest testing/python/riscv/test_riscv_mlir_codegen.py -q +# 33 passed + +python -m pytest testing/python/riscv/test_riscv_jit_runtime.py -q +# 16 passed + +python -m pytest testing/python/riscv/test_riscv_artifact_export.py -q +# 4 passed + +python -m pytest testing/python/riscv/test_riscv_examples.py -q +# 6 passed, 26 skipped + +python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q +# 6 passed + +python -m pytest testing/python/riscv/test_riscv_qemu_smoke.py -q +# 6 passed + +python -m pytest testing/python/riscv/test_riscv_tladapter_pipeline.py -q +# 1 passed + +python -m pytest testing/python/riscv -q +# 85 passed, 26 skipped +``` + +Coverage note: `testing/python/riscv/test_riscv_examples.py` keeps the original 16-entry RISC-V example list. The current checkout contains the three initial workflow example scripts, so the missing legacy example filenames are reported as skips instead of being removed from the coverage entry point. The focused Issue #1 workflow checks live in `testing/python/riscv/test_riscv_operator_workflow_examples.py`. + +Per-operator manual checks: + +```bash +python examples/riscv/example_vector_add.py --run-host --run-qemu --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/tilelang-riscv-vector_add +# vector_add host check passed +# vector_add qemu check passed + +python examples/riscv/example_reduce_sum.py --run-host --run-qemu --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/tilelang-riscv-reduce_sum +# reduce_sum host check passed +# reduce_sum qemu check passed + +python examples/riscv/example_matmul.py --run-host --run-qemu --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/tilelang-riscv-matmul +# matmul host check passed +# matmul qemu check passed +``` + +## Failure Log + +Current operator validation has no remaining failing operator. + +Two environment and lowering issues were found while preparing the workflow: + +- MLIR 18 API compatibility: Ubuntu MLIR 18 requires the newer `arith::ConstantIntOp`, + `arith::ConstantFloatOp`, and `memref::SubViewOp` result type usage. This was fixed in + `src/target/codegen_linalg_riscv.cc`. Failure bucket: MLIR lowering or structured codegen limitation. +- QEMU link failure from soft-float helpers such as `__addsf3`: the emitted RISC-V object did not + consistently include hard-float features for the freestanding runner. This was fixed by defaulting + `llc` to `-mattr=+m,+a,+f,+d,+c` and the QEMU linker path to `-march=rv64gc`, while preserving + environment overrides. Failure bucket: LLVM/RISC-V toolchain setup issue. + +## Effectiveness Summary + +The workflow is effective for the initial operator loop: + +- The catalog and generator create repeatable Codex task descriptions with operator semantics, shapes, dtypes, baseline references, scheduling guidance, compile commands, tests, success criteria, and failure buckets. +- The three initial operators cover elementwise, reduction, and matrix/tensor patterns. +- The examples can be used both as implementation references and as executable validation targets. +- Host execution and artifact export run on x86_64 with the RISC-V toolchain, which lets the agent loop be pre-validated without SG2044 hardware. +- QEMU checks provide a stronger non-native signal by executing freestanding RISC-V ELFs for all initial operators, but they do not replace SG2044 native RVV validation. +- Existing SG2044 backend baseline data is included from `docs/get_started/BuildOnSG2044.md` and kept separate from the new no-hardware operator validation results. + +Optional follow-up: + +- Replay the same commands on SG2044 for native RVV execution when hardware is available. +- RVV-specific checks can be added later by enabling vector-capable target attributes and checking emitted assembly for vector instructions such as `vsetvli`, `vle`, `vse`, `vfadd`, or `vfmacc`. +- Add more operators only after this initial workflow is reviewed, so new tasks can reuse the same catalog/report structure. + +## PR Summary Draft + +Summary: + +- Add a Codex-compatible TileLang-RISC-V operator workflow catalog, generated tasks, repository-local skill, and workflow guide. +- Add three initial validation operators covering elementwise, reduction, and matrix/tensor patterns. +- Add executable examples, workflow tests, QEMU smoke coverage, and a catalog-driven validation runner. +- Fix LLVM/MLIR 18 compatibility and default RISC-V hard-float codegen flags needed by the QEMU path. +- Record no-hardware validation results and existing SG2044 baseline data separately. + +Tests: + +```bash +bash maint/scripts/run_riscv_operator_workflow_validation.sh +# 85 passed, 26 skipped +# vector_add host check passed +# vector_add qemu check passed +# reduce_sum host check passed +# reduce_sum qemu check passed +# matmul host check passed +# matmul qemu check passed +``` diff --git a/.agents/skills/riscv-operator-workflow/SKILL.md b/.agents/skills/riscv-operator-workflow/SKILL.md new file mode 100644 index 0000000000..4c069366c8 --- /dev/null +++ b/.agents/skills/riscv-operator-workflow/SKILL.md @@ -0,0 +1,62 @@ +# RISC-V Operator Workflow + +Use this skill when creating, completing, or validating TileLang operators for the `riscv` / `linalg_riscv` backend. + +## Workflow + +1. Read the generated operator task from `.agents/tasks/riscv/`. +2. If the task does not exist, generate it with: + +```bash +python maint/scripts/generate_riscv_operator_task.py --operator +``` + +3. Review the context files named in the task before editing code. +4. Implement the operator in `examples/riscv/` or the task's suggested files. +5. Validate on x86/arm first with host execution and artifact export when the Buddy/LLVM toolchain is available. +6. Record failures using one of the task's failure buckets. +7. Update `.agents/riscv_operator_workflow/validation_report.md` with the compile, host, artifact, QEMU, existing SG2044 baseline, and optional SG2044 replay result. +8. Run or hand off the same task on SG2044 only when native hardware is available. + +## Commands + +List available operator tasks: + +```bash +python maint/scripts/generate_riscv_operator_task.py --list +``` + +Generate all initial tasks: + +```bash +python maint/scripts/generate_riscv_operator_task.py --all +``` + +Run the initial validation examples: + +```bash +python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q +python -m pytest testing/python/riscv/test_riscv_qemu_smoke.py -q +``` + +QEMU smoke checks are no-hardware functional validation. They do not replace SG2044 native RVV validation. + +Run an example manually: + +```bash +python examples/riscv/example_vector_add.py --run-host --output-dir /tmp/vector_add +python examples/riscv/example_vector_add.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add +``` + +## Reporting + +Each operator report should include: + +- operator name and category +- files changed +- host execution result +- artifact export result +- QEMU result +- optional SG2044 replay result +- failure bucket if validation failed +- follow-up needed diff --git a/.agents/tasks/riscv/matmul.md b/.agents/tasks/riscv/matmul.md new file mode 100644 index 0000000000..ae75fb70bc --- /dev/null +++ b/.agents/tasks/riscv/matmul.md @@ -0,0 +1,99 @@ +# Implement and validate small matmul for TileLang-RISCV + +## Goal + +Complete or refine the TileLang-RISCV implementation for `matmul` and validate it through task generation, x86/arm host execution, RISC-V artifact export, and QEMU execution when the toolchain is available. + +## Operator Semantics + +For M=2, N=4, K=3, compute C[i, j] = sum(A[i, k] * B[k, j] for k in [0, 3)). + +## Inputs + +- A: float32 tensor with shape (2, 3) +- B: float32 tensor with shape (3, 4) + +## Outputs + +- C: float32 tensor with shape (2, 4) + +## Baseline Reference + +- PyTorch: lhs @ rhs +- NumPy: lhs @ rhs + +## Baseline Implementation + +- PyTorch: expected = lhs @ rhs +- NumPy: expected = lhs @ rhs + +## Context Files + +- examples/riscv/example_matmul.py +- testing/python/riscv/test_riscv_examples.py +- testing/python/riscv/test_riscv_operator_workflow_examples.py +- .agents/riscv_operator_workflow/README.md +- docs/get_started/BuildOnSG2044.md +- docs/get_started/targets.md +- testing/python/riscv/test_riscv_target_parse.py +- testing/python/riscv/test_riscv_mlir_codegen.py +- testing/python/riscv/test_riscv_toolchain.py +- testing/python/riscv/test_riscv_jit_runtime.py +- testing/python/riscv/test_riscv_qemu_smoke.py + +## Scheduling Strategy + +- Use T.grid(M, N, K) with two spatial axes and one reduce axis. +- Use T.init() to zero C[i, j] before the K accumulation. +- Start with plain TIR loops rather than GPU tile intrinsics; move to T.gemm only after the linalg_riscv path supports the needed pattern. + +## Compilation Commands + +- python examples/riscv/example_matmul.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/matmul +- python examples/riscv/example_matmul.py --run-host --output-dir /tmp/matmul +- python examples/riscv/example_matmul.py --run-qemu --output-dir /tmp/matmul + +## Test Commands + +- python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k matmul + +## Success Criteria + +- Lowered MLIR keeps the matmul semantics visible in structured lowering or produces an equivalent loop form. +- Host adapter output matches the PyTorch matmul reference. +- Artifact export produces one non-empty .s and one non-empty .o file. + +## Failure Triage Buckets + +- Generated TileLang/TIR source error +- Unsupported scheduling or TileLang intrinsic for linalg_riscv +- MLIR lowering or structured codegen limitation +- LLVM/RISC-V toolchain setup issue +- Runtime wrapper, host adapter, QEMU, or optional SG2044 environment issue +- Test configuration or tolerance mismatch + +## Existing SG2044 Baseline + +- Existing documented SG2044 native baseline from docs/get_started/BuildOnSG2044.md: +- testing/python/riscv/test_riscv_target_parse.py: 2 passed +- testing/python/riscv/test_riscv_mlir_codegen.py: 33 passed +- testing/python/riscv/test_riscv_toolchain.py: 3 passed +- testing/python/riscv/test_riscv_jit_runtime.py: 16 passed +- This baseline was not rerun during the no-hardware workflow validation pass. + +## Reporting Template + +- Operator: +- Implementation files changed: +- Host compile/run result: +- RISC-V artifact result: +- QEMU result: +- Optional SG2044 replay result: +- Failure bucket, if any: +- Follow-up needed: + +## Target Flow + +TileLang -> MLIR Linalg -> MLIR vector/RVV lowering -> LLVM/RISC-V artifact -> host check -> QEMU smoke check -> optional SG2044 native replay + +The Codex operator workflow can be pre-validated without SG2044 hardware by using x86/arm host execution, RISC-V artifact export, and QEMU. Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available. diff --git a/.agents/tasks/riscv/reduce_sum.md b/.agents/tasks/riscv/reduce_sum.md new file mode 100644 index 0000000000..5f7df08136 --- /dev/null +++ b/.agents/tasks/riscv/reduce_sum.md @@ -0,0 +1,98 @@ +# Implement and validate row-wise reduce_sum for TileLang-RISCV + +## Goal + +Complete or refine the TileLang-RISCV implementation for `reduce_sum` and validate it through task generation, x86/arm host execution, RISC-V artifact export, and QEMU execution when the toolchain is available. + +## Operator Semantics + +For each row i in [0, 4), compute B[i] = sum(A[i, k] for k in [0, 8)). + +## Inputs + +- A: float32 tensor with shape (4, 8) + +## Outputs + +- B: float32 tensor with shape (4,) + +## Baseline Reference + +- PyTorch: data.sum(dim=1) +- NumPy: data.sum(axis=1) + +## Baseline Implementation + +- PyTorch: expected = data.sum(dim=1) +- NumPy: expected = data.sum(axis=1) + +## Context Files + +- examples/riscv/example_reduce_sum.py +- testing/python/riscv/test_riscv_examples.py +- testing/python/riscv/test_riscv_operator_workflow_examples.py +- .agents/riscv_operator_workflow/README.md +- docs/get_started/BuildOnSG2044.md +- docs/get_started/targets.md +- testing/python/riscv/test_riscv_target_parse.py +- testing/python/riscv/test_riscv_mlir_codegen.py +- testing/python/riscv/test_riscv_toolchain.py +- testing/python/riscv/test_riscv_jit_runtime.py +- testing/python/riscv/test_riscv_qemu_smoke.py + +## Scheduling Strategy + +- Use T.grid(ROWS, COLS) with one spatial row axis and one reduce axis. +- Use T.init() to zero the output row before accumulation. +- Keep the first validation shape small and static so qemu artifact generation has fixed memref sizes. + +## Compilation Commands + +- python examples/riscv/example_reduce_sum.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/reduce_sum +- python examples/riscv/example_reduce_sum.py --run-host --output-dir /tmp/reduce_sum +- python examples/riscv/example_reduce_sum.py --run-qemu --output-dir /tmp/reduce_sum + +## Test Commands + +- python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k reduce_sum + +## Success Criteria + +- Lowered MLIR contains a reduction-friendly structured form such as linalg.reduce or equivalent generic lowering. +- Host adapter output matches the PyTorch row-sum reference. +- Artifact export produces one non-empty .s and one non-empty .o file. + +## Failure Triage Buckets + +- Generated TileLang/TIR source error +- Unsupported scheduling or TileLang intrinsic for linalg_riscv +- MLIR lowering or structured codegen limitation +- LLVM/RISC-V toolchain setup issue +- Runtime wrapper, host adapter, QEMU, or optional SG2044 environment issue +- Test configuration or tolerance mismatch + +## Existing SG2044 Baseline + +- Existing documented SG2044 native baseline from docs/get_started/BuildOnSG2044.md: +- testing/python/riscv/test_riscv_target_parse.py: 2 passed +- testing/python/riscv/test_riscv_mlir_codegen.py: 33 passed +- testing/python/riscv/test_riscv_toolchain.py: 3 passed +- testing/python/riscv/test_riscv_jit_runtime.py: 16 passed +- This baseline was not rerun during the no-hardware workflow validation pass. + +## Reporting Template + +- Operator: +- Implementation files changed: +- Host compile/run result: +- RISC-V artifact result: +- QEMU result: +- Optional SG2044 replay result: +- Failure bucket, if any: +- Follow-up needed: + +## Target Flow + +TileLang -> MLIR Linalg -> MLIR vector/RVV lowering -> LLVM/RISC-V artifact -> host check -> QEMU smoke check -> optional SG2044 native replay + +The Codex operator workflow can be pre-validated without SG2044 hardware by using x86/arm host execution, RISC-V artifact export, and QEMU. Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available. diff --git a/.agents/tasks/riscv/vector_add.md b/.agents/tasks/riscv/vector_add.md new file mode 100644 index 0000000000..53f391e5ed --- /dev/null +++ b/.agents/tasks/riscv/vector_add.md @@ -0,0 +1,100 @@ +# Implement and validate vector_add for TileLang-RISCV + +## Goal + +Complete or refine the TileLang-RISCV implementation for `vector_add` and validate it through task generation, x86/arm host execution, RISC-V artifact export, and QEMU execution when the toolchain is available. + +## Operator Semantics + +For each i in [0, 8), compute C[i] = A[i] + B[i]. + +## Inputs + +- A: float32 tensor with shape (8,) +- B: float32 tensor with shape (8,) + +## Outputs + +- C: float32 tensor with shape (8,) + +## Baseline Reference + +- PyTorch: lhs + rhs +- NumPy: lhs + rhs + +## Baseline Implementation + +- PyTorch: expected = lhs + rhs +- NumPy: expected = lhs + rhs + +## Context Files + +- examples/riscv/example_vector_add.py +- testing/python/riscv/test_riscv_examples.py +- testing/python/riscv/test_riscv_operator_workflow_examples.py +- testing/python/riscv/test_riscv_qemu_smoke.py +- .agents/riscv_operator_workflow/README.md +- docs/get_started/BuildOnSG2044.md +- docs/get_started/targets.md +- testing/python/riscv/test_riscv_target_parse.py +- testing/python/riscv/test_riscv_mlir_codegen.py +- testing/python/riscv/test_riscv_toolchain.py +- testing/python/riscv/test_riscv_jit_runtime.py + +## Scheduling Strategy + +- Use a single serial loop over N. +- Use a T.block with one spatial axis. +- Avoid GPU-specific T.Kernel threading, shared memory, warp, or async-copy constructs for this first RISC-V path. + +## Compilation Commands + +- python examples/riscv/example_vector_add.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add +- python examples/riscv/example_vector_add.py --run-host --output-dir /tmp/vector_add +- python examples/riscv/example_vector_add.py --run-qemu --output-dir /tmp/vector_add + +## Test Commands + +- python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q -k vector_add +- python -m pytest testing/python/riscv/test_riscv_qemu_smoke.py -q + +## Success Criteria + +- Host adapter output matches the PyTorch reference. +- Artifact export produces one non-empty .s and one non-empty .o file. +- QEMU smoke emits vector_add.qemu.elf and output matches NumPy reference when a runner is available. + +## Failure Triage Buckets + +- Generated TileLang/TIR source error +- Unsupported scheduling or TileLang intrinsic for linalg_riscv +- MLIR lowering or structured codegen limitation +- LLVM/RISC-V toolchain setup issue +- Runtime wrapper, host adapter, QEMU, or optional SG2044 environment issue +- Test configuration or tolerance mismatch + +## Existing SG2044 Baseline + +- Existing documented SG2044 native baseline from docs/get_started/BuildOnSG2044.md: +- testing/python/riscv/test_riscv_target_parse.py: 2 passed +- testing/python/riscv/test_riscv_mlir_codegen.py: 33 passed +- testing/python/riscv/test_riscv_toolchain.py: 3 passed +- testing/python/riscv/test_riscv_jit_runtime.py: 16 passed +- This baseline was not rerun during the no-hardware workflow validation pass. + +## Reporting Template + +- Operator: +- Implementation files changed: +- Host compile/run result: +- RISC-V artifact result: +- QEMU result: +- Optional SG2044 replay result: +- Failure bucket, if any: +- Follow-up needed: + +## Target Flow + +TileLang -> MLIR Linalg -> MLIR vector/RVV lowering -> LLVM/RISC-V artifact -> host check -> QEMU smoke check -> optional SG2044 native replay + +The Codex operator workflow can be pre-validated without SG2044 hardware by using x86/arm host execution, RISC-V artifact export, and QEMU. Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available. diff --git a/AGENTS.md b/AGENTS.md new file mode 100644 index 0000000000..530c3c4960 --- /dev/null +++ b/AGENTS.md @@ -0,0 +1,19 @@ +# AGENTS.md + +## TileLang-RISC-V Operator Workflow + +When creating, completing, or validating RISC-V operators, first read: + +- `.agents/skills/riscv-operator-workflow/SKILL.md` +- `.agents/riscv_operator_workflow/operators.json` +- `docs/get_started/RiscvOperatorAgentWorkflow.md` +- `docs/get_started/BuildOnSG2044.md` +- `testing/python/riscv/test_riscv_operator_workflow.py` + +Rules: + +- Treat `target="riscv"` as the public spelling and `linalg_riscv` as the internal backend target. +- Use x86/arm for task generation, MLIR/codegen review, host adapter checks, artifact export, and QEMU checks. +- Do not claim new SG2044 native validation unless the workflow was actually rerun on SG2044 hardware. +- Initial operator workflow coverage uses `vector_add`, `reduce_sum`, and `matmul`. +- Record failures using the workflow failure buckets. diff --git a/docs/get_started/BuildOnSG2044.md b/docs/get_started/BuildOnSG2044.md index c02610e3ad..7631a98478 100644 --- a/docs/get_started/BuildOnSG2044.md +++ b/docs/get_started/BuildOnSG2044.md @@ -91,6 +91,40 @@ Observed results from the validated SG2044 bring-up: - `test_riscv_toolchain.py`: `3 passed` - `test_riscv_jit_runtime.py`: `16 passed` +## Operator Agent Workflow Validation + +The Codex operator workflow can be pre-validated without SG2044 hardware by using x86 or arm host execution, RISC-V artifact export, and QEMU. Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available. +See `docs/get_started/RiscvOperatorAgentWorkflow.md` for the generated task format and the initial operator set. + +The validation levels are: + +- x86/arm host check +- RISC-V artifact export plus QEMU smoke check +- SG2044 native hardware replay + +For native SG2044 replay, run the example workflow tests after the native environment is active: + +```bash +cd tilelang-riscv +TILELANG_RISCV_RUN_QEMU=0 bash maint/scripts/run_riscv_operator_workflow_validation.sh +``` + +Or run the focused pytest commands directly: + +```bash +python -m pytest testing/python/riscv/test_riscv_examples.py -q +python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q +python -m pytest testing/python/riscv/test_riscv_qemu_smoke.py -q +``` + +QEMU smoke tests are useful for no-hardware functional validation, but they do not replace SG2044 native RVV validation. + +For a single operator smoke test: + +```bash +python examples/riscv/example_vector_add.py --run-host --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add +``` + ## Troubleshooting - If configuration fails with a Z3 lookup error, check that `Z3_ROOT` points to a prefix containing `include` and `lib` diff --git a/docs/get_started/RiscvOperatorAgentWorkflow.md b/docs/get_started/RiscvOperatorAgentWorkflow.md new file mode 100644 index 0000000000..6e4d809e9b --- /dev/null +++ b/docs/get_started/RiscvOperatorAgentWorkflow.md @@ -0,0 +1,136 @@ +# RISC-V Operator Agent Workflow + +This workflow turns a TileLang-RISCV operator request into a Codex-ready task, then validates the result through the same commands used by the RISC-V backend tests. + +The Codex operator workflow can be pre-validated without SG2044 hardware by using x86 or arm host execution, RISC-V artifact export, and QEMU. +Native SG2044 replay remains the final hardware validation step when SG2044 hardware is available. + +## Files + +- `.agents/riscv_operator_workflow/operators.json`: operator catalog and shared validation context +- `.agents/riscv_operator_workflow/README.md`: short workflow entry point +- `maint/scripts/generate_riscv_operator_task.py`: task generator +- `.agents/tasks/riscv/*.md`: generated Codex task files +- `.agents/skills/riscv-operator-workflow/SKILL.md`: repository-local skill for the workflow +- `examples/riscv/*.py`: executable validation examples +- `testing/python/riscv/test_riscv_operator_workflow_examples.py`: focused tests for the initial workflow examples +- `.agents/riscv_operator_workflow/validation_report.md`: compilation, correctness, and failure-triage report for the initial operators + +## Generate Tasks + +List available operator tasks: + +```bash +python maint/scripts/generate_riscv_operator_task.py --list +``` + +Generate all initial tasks: + +```bash +python maint/scripts/generate_riscv_operator_task.py --all +``` + +Generate one task: + +```bash +python maint/scripts/generate_riscv_operator_task.py --operator vector_add +``` + +## Initial Operators + +| Operator | Pattern | Example | +| --- | --- | --- | +| `vector_add` | elementwise | `examples/riscv/example_vector_add.py` | +| `reduce_sum` | reduction | `examples/riscv/example_reduce_sum.py` | +| `matmul` | matrix/tensor compute | `examples/riscv/example_matmul.py` | + +These three operators are intentionally small and static-shaped. They cover the initial workflow categories while keeping artifact and QEMU validation simple. + +The current QEMU runner path is intended for small static-shaped examples with simple memref-style signatures. More complex dynamic-shape or runtime-heavy operators should be validated separately on SG2044 native hardware. + +## Validation Levels + +Use three separate validation levels when reporting results: + +- x86/arm host check +- RISC-V artifact export plus QEMU smoke check +- SG2044 native hardware replay + +QEMU smoke checks are useful no-hardware functional validation, but they do not replace SG2044 native RVV validation. + +## x86/arm Validation + +On an x86 or arm developer host, use the workflow to validate task generation and the non-native parts of the RISC-V flow: + +```bash +bash maint/scripts/run_riscv_operator_workflow_validation.sh +``` + +The validation script reads `.agents/riscv_operator_workflow/operators.json`, so adding another catalog entry is enough to include that operator in the manual validation loop. +Useful environment variables: + +- `TILELANG_RISCV_RUN_QEMU=0|1|auto`: control qemu execution; `auto` is the default. +- `TILELANG_RISCV_OPERATORS="vector_add matmul"`: validate only selected catalog operators. +- `TILELANG_RISCV_PYTEST_ARGS="testing/python/riscv/test_riscv_examples.py -q"`: override the pytest scope. +- `TILELANG_RISCV_RUN_PYTEST=0`: skip pytest and run only the manual operator loop. +- `TILELANG_RISCV_VALIDATION_OUTPUT_ROOT=/tmp`: choose where per-operator artifacts are written. +- `TILELANG_RISCV_RUN_HOST=0`: skip manual host execution. +- `TILELANG_RISCV_EMIT_ARTIFACTS=0`: skip manual artifact export. + +Or run the checks individually: + +```bash +python maint/scripts/generate_riscv_operator_task.py --all +python -m pytest testing/python/riscv/test_riscv_operator_workflow.py -q +python -m pytest testing/python/riscv/test_riscv_target_parse.py -q +python -m pytest testing/python/riscv/test_riscv_mlir_codegen.py -q +python -m pytest testing/python/riscv/test_riscv_examples.py -q +python -m pytest testing/python/riscv/test_riscv_operator_workflow_examples.py -q +``` + +When the Buddy/LLVM tools are available, also validate artifact export: + +```bash +python examples/riscv/example_vector_add.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add +python examples/riscv/example_reduce_sum.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/reduce_sum +python examples/riscv/example_matmul.py --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/matmul +``` + +If the host adapter toolchain is configured, run host correctness: + +```bash +python examples/riscv/example_vector_add.py --run-host --output-dir /tmp/vector_add +python examples/riscv/example_reduce_sum.py --run-host --output-dir /tmp/reduce_sum +python examples/riscv/example_matmul.py --run-host --output-dir /tmp/matmul +``` + +## Optional SG2044 Replay + +SG2044 hardware is not required for no-hardware functional validation. When an SG2044 machine is available, first follow the native setup in `docs/get_started/BuildOnSG2044.md`, then run: + +```bash +TILELANG_RISCV_RUN_QEMU=0 bash maint/scripts/run_riscv_operator_workflow_validation.sh +``` + +For an individual operator: + +```bash +python examples/riscv/example_vector_add.py --run-host --emit-mlir --emit-llvm --emit-asm --emit-object --output-dir /tmp/vector_add +``` + +## Results Report + +Record each operator result in `.agents/riscv_operator_workflow/validation_report.md`. +The report tracks host correctness, RISC-V artifact export, QEMU execution, optional SG2044 replay status, failure buckets, and follow-up. +It also includes the existing SG2044 baseline results already documented in `docs/get_started/BuildOnSG2044.md`, clearly separated from the new no-hardware operator validation results. + +## Failure Buckets + +Use these buckets in task reports: + +- generated TileLang/TIR source error +- unsupported schedule or TileLang intrinsic for `linalg_riscv` +- MLIR lowering or structured codegen limitation +- LLVM/RISC-V toolchain setup issue +- runtime wrapper, host adapter, QEMU, or optional SG2044 environment issue +- test configuration or tolerance mismatch diff --git a/docs/index.md b/docs/index.md index 1c78ea2f6f..ca7c9b7e36 100644 --- a/docs/index.md +++ b/docs/index.md @@ -15,6 +15,8 @@ low-level optimizations necessary for state-of-the-art performance. get_started/Installation get_started/overview get_started/targets +get_started/BuildOnSG2044 +get_started/RiscvOperatorAgentWorkflow ::: :::{toctree} diff --git a/examples/riscv/__init__.py b/examples/riscv/__init__.py new file mode 100644 index 0000000000..7c6a44d3a5 --- /dev/null +++ b/examples/riscv/__init__.py @@ -0,0 +1 @@ +"""Small RISC-V validation examples.""" diff --git a/examples/riscv/example_matmul.py b/examples/riscv/example_matmul.py new file mode 100644 index 0000000000..709551656b --- /dev/null +++ b/examples/riscv/example_matmul.py @@ -0,0 +1,63 @@ +from __future__ import annotations + +import sys +from pathlib import Path + +import numpy as np +import torch + +REPO_ROOT = Path(__file__).resolve().parents[2] +if str(REPO_ROOT) not in sys.path: + sys.path.insert(0, str(REPO_ROOT)) + +import tilelang.language as T + +try: + from .riscv_example_utils import run_example +except ImportError: + from riscv_example_utils import run_example + + +M = 2 +N = 4 +K = 3 + + +@T.prim_func +def matmul( + A: T.Tensor((M, K), "float32"), + B: T.Tensor((K, N), "float32"), + C: T.Tensor((M, N), "float32"), +): + for i, j, kk in T.grid(M, N, K): + with T.block("matmul"): + vi = T.axis.spatial(M, i) + vj = T.axis.spatial(N, j) + vk = T.axis.reduce(K, kk) + with T.init(): + C[vi, vj] = T.float32(0) + C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vk, vj] + + +def make_torch_case(): + lhs = torch.arange(M * K, dtype=torch.float32).reshape(M, K) + rhs = torch.linspace(-1.0, 2.0, steps=K * N, dtype=torch.float32).reshape(K, N) + return (lhs, rhs), lhs @ rhs + + +def make_numpy_case(): + lhs = np.arange(M * K, dtype=np.float32).reshape(M, K) + rhs = np.linspace(-1.0, 2.0, num=K * N, dtype=np.float32).reshape(K, N) + out = np.zeros((M, N), dtype=np.float32) + return (lhs, rhs, out), lhs @ rhs, 2 + + +if __name__ == "__main__": + run_example( + name="matmul", + description="RISC-V small matmul validation example.", + func=matmul, + out_idx=[2], + make_torch_case=make_torch_case, + make_numpy_case=make_numpy_case, + ) diff --git a/examples/riscv/example_reduce_sum.py b/examples/riscv/example_reduce_sum.py new file mode 100644 index 0000000000..2c2052fc71 --- /dev/null +++ b/examples/riscv/example_reduce_sum.py @@ -0,0 +1,58 @@ +from __future__ import annotations + +import sys +from pathlib import Path + +import numpy as np +import torch + +REPO_ROOT = Path(__file__).resolve().parents[2] +if str(REPO_ROOT) not in sys.path: + sys.path.insert(0, str(REPO_ROOT)) + +import tilelang.language as T + +try: + from .riscv_example_utils import run_example +except ImportError: + from riscv_example_utils import run_example + + +ROWS = 4 +COLS = 8 + + +@T.prim_func +def reduce_sum_rows( + A: T.Tensor((ROWS, COLS), "float32"), + B: T.Tensor((ROWS,), "float32"), +): + for i, k in T.grid(ROWS, COLS): + with T.block("sum"): + vi = T.axis.spatial(ROWS, i) + vk = T.axis.reduce(COLS, k) + with T.init(): + B[vi] = T.float32(0) + B[vi] = B[vi] + A[vi, vk] + + +def make_torch_case(): + data = torch.linspace(-3.0, 5.0, steps=ROWS * COLS, dtype=torch.float32).reshape(ROWS, COLS) + return (data,), data.sum(dim=1) + + +def make_numpy_case(): + data = np.linspace(-3.0, 5.0, num=ROWS * COLS, dtype=np.float32).reshape(ROWS, COLS) + out = np.zeros((ROWS,), dtype=np.float32) + return (data, out), data.sum(axis=1), 1 + + +if __name__ == "__main__": + run_example( + name="reduce_sum", + description="RISC-V row-wise reduce-sum validation example.", + func=reduce_sum_rows, + out_idx=[1], + make_torch_case=make_torch_case, + make_numpy_case=make_numpy_case, + ) diff --git a/examples/riscv/example_vector_add.py b/examples/riscv/example_vector_add.py new file mode 100644 index 0000000000..f26300ad62 --- /dev/null +++ b/examples/riscv/example_vector_add.py @@ -0,0 +1,57 @@ +from __future__ import annotations + +import sys +from pathlib import Path + +import numpy as np +import torch + +REPO_ROOT = Path(__file__).resolve().parents[2] +if str(REPO_ROOT) not in sys.path: + sys.path.insert(0, str(REPO_ROOT)) + +import tilelang.language as T + +try: + from .riscv_example_utils import run_example +except ImportError: + from riscv_example_utils import run_example + + +N = 8 + + +@T.prim_func +def vector_add( + A: T.Tensor((N,), "float32"), + B: T.Tensor((N,), "float32"), + C: T.Tensor((N,), "float32"), +): + for i in T.serial(N): + with T.block("add"): + vi = T.axis.spatial(N, i) + C[vi] = A[vi] + B[vi] + + +def make_torch_case(): + lhs = torch.linspace(-2.0, 2.0, steps=N, dtype=torch.float32) + rhs = torch.linspace(0.25, 1.75, steps=N, dtype=torch.float32) + return (lhs, rhs), lhs + rhs + + +def make_numpy_case(): + lhs = np.linspace(-2.0, 2.0, num=N, dtype=np.float32) + rhs = np.linspace(0.25, 1.75, num=N, dtype=np.float32) + out = np.zeros_like(lhs) + return (lhs, rhs, out), lhs + rhs, 2 + + +if __name__ == "__main__": + run_example( + name="vector_add", + description="RISC-V vector add validation example.", + func=vector_add, + out_idx=[2], + make_torch_case=make_torch_case, + make_numpy_case=make_numpy_case, + ) diff --git a/examples/riscv/riscv_example_utils.py b/examples/riscv/riscv_example_utils.py new file mode 100644 index 0000000000..7a275abfa3 --- /dev/null +++ b/examples/riscv/riscv_example_utils.py @@ -0,0 +1,102 @@ +from __future__ import annotations + +import argparse +from pathlib import Path +from typing import Callable + +import numpy as np +import torch + +import tilelang +from tilelang.jit.adapter.riscv import emit_asm, emit_llvm_ir, emit_mlir, emit_object, run_qemu + + +def build_arg_parser(description: str) -> argparse.ArgumentParser: + parser = argparse.ArgumentParser(description=description) + parser.add_argument("--output-dir", type=Path, default=Path("riscv_artifacts")) + parser.add_argument("--run-host", action="store_true", help="Compile and execute through the native host adapter.") + parser.add_argument("--run-qemu", action="store_true", help="Build and execute a freestanding RISC-V ELF.") + parser.add_argument("--emit-mlir", action="store_true", help="Write the structured MLIR module.") + parser.add_argument("--emit-llvm", action="store_true", help="Write translated LLVM IR.") + parser.add_argument("--emit-asm", action="store_true", help="Write RISC-V assembly.") + parser.add_argument("--emit-object", action="store_true", help="Write a RISC-V object file.") + return parser + + +def lower_riscv(func): + return tilelang.lower(func, target="riscv").rt_mod + + +def emit_requested_artifacts(name: str, func, args: argparse.Namespace) -> None: + if not (args.emit_mlir or args.emit_llvm or args.emit_asm or args.emit_object): + return + + args.output_dir.mkdir(parents=True, exist_ok=True) + rt_mod = lower_riscv(func) + if args.emit_mlir: + emit_mlir(rt_mod, args.output_dir / f"{name}.mlir") + if args.emit_llvm: + emit_llvm_ir(rt_mod, args.output_dir / f"{name}.ll") + if args.emit_asm: + emit_asm(rt_mod, args.output_dir / f"{name}.s") + if args.emit_object: + emit_object(rt_mod, args.output_dir / f"{name}.o") + + +def run_host_check(name: str, func, out_idx: list[int], inputs: tuple[torch.Tensor, ...], expected: torch.Tensor) -> None: + kernel = tilelang.compile(func, out_idx=out_idx, target="riscv") + try: + actual = kernel(*inputs) + torch.testing.assert_close(actual, expected, atol=1e-5, rtol=1e-5) + finally: + kernel.close() + print(f"{name} host check passed") + + +def run_qemu_check( + name: str, + func, + all_args: tuple[np.ndarray, ...], + expected: np.ndarray, + output_arg_index: int, + output_dir: Path, +) -> None: + """Run small static-shaped examples with simple memref-style signatures.""" + output_dir.mkdir(parents=True, exist_ok=True) + run_qemu(lower_riscv(func), *all_args, path=output_dir / f"{name}.qemu.elf") + np.testing.assert_allclose(all_args[output_arg_index], expected, atol=1e-5, rtol=1e-5) + print(f"{name} qemu check passed") + + +def run_example( + *, + name: str, + description: str, + func, + out_idx: list[int], + make_torch_case: Callable[[], tuple[tuple[torch.Tensor, ...], torch.Tensor]], + make_numpy_case: Callable[[], tuple[tuple[np.ndarray, ...], np.ndarray, int]], +) -> None: + parser = build_arg_parser(description) + args = parser.parse_args() + emit_requested_artifacts(name, func, args) + + if args.run_host: + inputs, expected = make_torch_case() + run_host_check(name, func, out_idx, inputs, expected) + + if args.run_qemu: + all_args, expected, output_arg_index = make_numpy_case() + run_qemu_check(name, func, all_args, expected, output_arg_index, args.output_dir) + + if not any( + ( + args.emit_mlir, + args.emit_llvm, + args.emit_asm, + args.emit_object, + args.run_host, + args.run_qemu, + ) + ): + parser.print_help() diff --git a/maint/scripts/generate_riscv_operator_task.py b/maint/scripts/generate_riscv_operator_task.py new file mode 100644 index 0000000000..24dc13beed --- /dev/null +++ b/maint/scripts/generate_riscv_operator_task.py @@ -0,0 +1,144 @@ +#!/usr/bin/env python3 +from __future__ import annotations + +import argparse +import json +from pathlib import Path +from typing import Any + + +REPO_ROOT = Path(__file__).resolve().parents[2] +DEFAULT_CATALOG = REPO_ROOT / ".agents" / "riscv_operator_workflow" / "operators.json" +DEFAULT_OUTPUT_DIR = REPO_ROOT / ".agents" / "tasks" / "riscv" + + +def _load_catalog(path: Path) -> dict[str, Any]: + with path.open(encoding="utf-8") as handle: + return json.load(handle) + + +def _operator_map(catalog: dict[str, Any]) -> dict[str, dict[str, Any]]: + return {item["name"]: item for item in catalog.get("operators", [])} + + +def _render_list(title: str, values: list[str]) -> list[str]: + lines = [f"## {title}", ""] + lines.extend(f"- {value}" for value in values) + lines.append("") + return lines + + +def _unique(values: list[str]) -> list[str]: + seen: set[str] = set() + result: list[str] = [] + for value in values: + if value in seen: + continue + seen.add(value) + result.append(value) + return result + + +def render_task(catalog: dict[str, Any], operator: dict[str, Any]) -> str: + common = catalog["common_context"] + lines = [ + f"# {operator['task_title']}", + "", + "## Goal", + "", + ( + f"Complete or refine the TileLang-RISCV implementation for `{operator['name']}` and validate it through " + "task generation, x86/arm host execution, RISC-V artifact export, and QEMU execution when the toolchain is available." + ), + "", + "## Operator Semantics", + "", + operator["semantics"], + "", + ] + lines.extend(_render_list("Inputs", operator["inputs"])) + lines.extend(_render_list("Outputs", operator["outputs"])) + lines.extend(_render_list("Baseline Reference", operator["baseline_reference"])) + lines.extend(_render_list("Baseline Implementation", operator["baseline_implementation"])) + lines.extend( + _render_list( + "Context Files", + _unique([ + operator["current_example"], + *operator["suggested_files"], + *common["primary_docs"], + *common["reference_tests"], + ]), + ) + ) + lines.extend(_render_list("Scheduling Strategy", operator["schedule_strategy"])) + lines.extend(_render_list("Compilation Commands", operator["compile_commands"])) + lines.extend(_render_list("Test Commands", operator["test_commands"])) + lines.extend(_render_list("Success Criteria", operator["success_criteria"])) + lines.extend(_render_list("Failure Triage Buckets", common["common_failure_buckets"])) + if common.get("sg2044_baseline"): + lines.extend(_render_list("Existing SG2044 Baseline", common["sg2044_baseline"])) + lines.extend( + [ + "## Reporting Template", + "", + "- Operator:", + "- Implementation files changed:", + "- Host compile/run result:", + "- RISC-V artifact result:", + "- QEMU result:", + "- Optional SG2044 replay result:", + "- Failure bucket, if any:", + "- Follow-up needed:", + "", + "## Target Flow", + "", + common["riscv_flow"], + "", + catalog["validation_note"], + "", + ] + ) + return "\n".join(lines) + + +def write_task(catalog: dict[str, Any], operator: dict[str, Any], output_dir: Path) -> Path: + output_dir.mkdir(parents=True, exist_ok=True) + path = output_dir / f"{operator['name']}.md" + path.write_text(render_task(catalog, operator), encoding="utf-8") + return path + + +def main() -> None: + parser = argparse.ArgumentParser(description="Generate Codex task files for TileLang-RISCV operators.") + parser.add_argument("--catalog", type=Path, default=DEFAULT_CATALOG) + parser.add_argument("--operator", help="Operator name to render.") + parser.add_argument("--all", action="store_true", help="Render all operators in the catalog.") + parser.add_argument("--list", action="store_true", help="List available operator names.") + parser.add_argument("--output-dir", type=Path, default=DEFAULT_OUTPUT_DIR) + args = parser.parse_args() + + catalog = _load_catalog(args.catalog) + operators = _operator_map(catalog) + + if args.list: + for name in sorted(operators): + print(name) + return + + if args.all: + selected = [operators[name] for name in sorted(operators)] + elif args.operator: + if args.operator not in operators: + choices = ", ".join(sorted(operators)) + raise SystemExit(f"Unknown operator `{args.operator}`. Available operators: {choices}") + selected = [operators[args.operator]] + else: + raise SystemExit("Pass --operator NAME, --all, or --list.") + + for operator in selected: + print(write_task(catalog, operator, args.output_dir)) + + +if __name__ == "__main__": + main() diff --git a/maint/scripts/run_riscv_operator_workflow_validation.sh b/maint/scripts/run_riscv_operator_workflow_validation.sh new file mode 100644 index 0000000000..7fe3388699 --- /dev/null +++ b/maint/scripts/run_riscv_operator_workflow_validation.sh @@ -0,0 +1,87 @@ +#!/usr/bin/env bash +set -euo pipefail + +repo_root="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" +cd "${repo_root}" + +python_bin="${PYTHON:-python}" +run_qemu="${TILELANG_RISCV_RUN_QEMU:-auto}" +run_host="${TILELANG_RISCV_RUN_HOST:-1}" +run_pytest="${TILELANG_RISCV_RUN_PYTEST:-1}" +emit_artifacts="${TILELANG_RISCV_EMIT_ARTIFACTS:-1}" +validation_output_root="${TILELANG_RISCV_VALIDATION_OUTPUT_ROOT:-/tmp}" + +require_binary_flag() { + local name="$1" + local value="$2" + if [[ "${value}" != "0" && "${value}" != "1" ]]; then + echo "${name} must be 0 or 1." >&2 + exit 2 + fi +} + +if [[ "${run_qemu}" == "auto" ]]; then + if command -v qemu-riscv64 >/dev/null 2>&1 || [[ -n "${TILELANG_RISCV_RUNNER:-}" ]]; then + run_qemu=1 + else + run_qemu=0 + fi +fi + +if [[ "${run_qemu}" != "0" && "${run_qemu}" != "1" ]]; then + echo "TILELANG_RISCV_RUN_QEMU must be 0, 1, or auto." >&2 + exit 2 +fi +require_binary_flag TILELANG_RISCV_RUN_HOST "${run_host}" +require_binary_flag TILELANG_RISCV_RUN_PYTEST "${run_pytest}" +require_binary_flag TILELANG_RISCV_EMIT_ARTIFACTS "${emit_artifacts}" + +"${python_bin}" maint/scripts/generate_riscv_operator_task.py --all + +if [[ "${run_pytest}" == "1" ]]; then + read -r -a pytest_args <<< "${TILELANG_RISCV_PYTEST_ARGS:-testing/python/riscv -q}" + "${python_bin}" -m pytest "${pytest_args[@]}" +fi + +if [[ "${run_host}" != "1" && "${emit_artifacts}" != "1" && "${run_qemu}" != "1" ]]; then + exit 0 +fi + +operator_lines="$( + "${python_bin}" - <<'PY' +import json +import os +from pathlib import Path + +catalog = json.loads(Path(".agents/riscv_operator_workflow/operators.json").read_text(encoding="utf-8")) +selected = set(os.environ.get("TILELANG_RISCV_OPERATORS", "").split()) +known = {operator["name"] for operator in catalog["operators"]} +unknown = sorted(selected - known) +if unknown: + raise SystemExit(f"Unknown TILELANG_RISCV_OPERATORS entries: {', '.join(unknown)}") +for operator in catalog["operators"]: + if selected and operator["name"] not in selected: + continue + print(f"{operator['name']}\t{operator['current_example']}") +PY +)" + +if [[ -z "${operator_lines}" ]]; then + echo "No operators selected for manual validation." >&2 + exit 2 +fi + +while IFS=$'\t' read -r operator example_path; do + args=() + if [[ "${run_host}" == "1" ]]; then + args+=(--run-host) + fi + if [[ "${emit_artifacts}" == "1" ]]; then + args+=(--emit-mlir --emit-llvm --emit-asm --emit-object) + fi + args+=(--output-dir "${validation_output_root%/}/tilelang-riscv-${operator}") + if [[ "${run_qemu}" == "1" ]]; then + args+=(--run-qemu) + fi + "${python_bin}" "${example_path}" "${args[@]}" +done <<< "${operator_lines}" diff --git a/src/target/codegen_linalg_riscv.cc b/src/target/codegen_linalg_riscv.cc index 52fc5eda23..5c1d2f5eed 100644 --- a/src/target/codegen_linalg_riscv.cc +++ b/src/target/codegen_linalg_riscv.cc @@ -650,7 +650,7 @@ class TIRToMLIRLowerer final : private tir::StmtFunctor, if (type.isIndex()) { return builder_.create(loc_, value); } - return builder_.create(loc_, type, value); + return builder_.create(loc_, value, type); } mlir::Value ZeroIndex() { return ConstantIntLike(0, builder_.getIndexType()); } @@ -1135,7 +1135,7 @@ class TIRToMLIRLowerer final : private tir::StmtFunctor, if (mlir::isa(value_type)) { mlir::Type zero_type = value_type; mlir::Value zero = builder_.create( - loc_, mlir::cast(zero_type), llvm::APFloat(0.0)); + loc_, llvm::APFloat(0.0), mlir::cast(zero_type)); return builder_.create(loc_, mlir::arith::CmpFPredicate::UNE, value, zero); } @@ -1222,10 +1222,11 @@ class TIRToMLIRLowerer final : private tir::StmtFunctor, mlir::MemRefType result_type; if (target_buffer->shape.size() == source_region->region.size()) { - result_type = mlir::memref::SubViewOp::inferResultType(source_type, offsets, sizes, strides); + result_type = mlir::cast( + mlir::memref::SubViewOp::inferResultType(source_type, offsets, sizes, strides)); } else { - result_type = mlir::memref::SubViewOp::inferRankReducedResultType( - LowerStaticShape(target_buffer->shape), source_type, offsets, sizes, strides); + result_type = mlir::cast(mlir::memref::SubViewOp::inferRankReducedResultType( + LowerStaticShape(target_buffer->shape), source_type, offsets, sizes, strides)); } return builder_ @@ -1240,8 +1241,8 @@ class TIRToMLIRLowerer final : private tir::StmtFunctor, llvm::SmallVector offsets = LowerRegionOffsets(region->region); llvm::SmallVector sizes = LowerRegionSizes(region->region); llvm::SmallVector strides = UnitStrides(region->region.size()); - mlir::MemRefType result_type = - mlir::memref::SubViewOp::inferResultType(source_type, offsets, sizes, strides); + mlir::MemRefType result_type = mlir::cast( + mlir::memref::SubViewOp::inferResultType(source_type, offsets, sizes, strides)); return builder_ .create(loc_, result_type, source, offsets, sizes, strides) .getResult(); @@ -1264,8 +1265,9 @@ class TIRToMLIRLowerer final : private tir::StmtFunctor, llvm::SmallVector offsets = LowerRegionOffsets(region->region); llvm::SmallVector sizes = LowerRegionSizes(region->region); llvm::SmallVector strides = UnitStrides(region->region.size()); - mlir::MemRefType result_type = mlir::memref::SubViewOp::inferRankReducedResultType( - LowerStaticShape(logical_extents), source_type, offsets, sizes, strides); + mlir::MemRefType result_type = mlir::cast( + mlir::memref::SubViewOp::inferRankReducedResultType(LowerStaticShape(logical_extents), + source_type, offsets, sizes, strides)); return builder_ .create(loc_, result_type, source, offsets, sizes, strides) .getResult(); diff --git a/testing/python/riscv/test_riscv_examples.py b/testing/python/riscv/test_riscv_examples.py index 0b8e5adcc5..a0461a5726 100644 --- a/testing/python/riscv/test_riscv_examples.py +++ b/testing/python/riscv/test_riscv_examples.py @@ -46,7 +46,17 @@ def _env() -> dict[str, str]: return env -@pytest.mark.parametrize("example_name", EXAMPLES) +def _example_params(): + params = [] + for example_name in EXAMPLES: + marks = [] + if not (EXAMPLE_ROOT / example_name).is_file(): + marks.append(pytest.mark.skip(reason=f"{example_name} is not present in this checkout")) + params.append(pytest.param(example_name, marks=marks, id=example_name)) + return params + + +@pytest.mark.parametrize("example_name", _example_params()) def test_riscv_examples_run_on_host(example_name, tmp_path): result = subprocess.run( [sys.executable, str(EXAMPLE_ROOT / example_name), "--run-host", "--output-dir", str(tmp_path / example_name)], @@ -61,7 +71,7 @@ def test_riscv_examples_run_on_host(example_name, tmp_path): assert "host check passed" in result.stdout -@pytest.mark.parametrize("example_name", EXAMPLES) +@pytest.mark.parametrize("example_name", _example_params()) def test_riscv_examples_emit_riscv_artifacts(example_name, tmp_path): output_dir = tmp_path / example_name result = subprocess.run( diff --git a/testing/python/riscv/test_riscv_operator_workflow.py b/testing/python/riscv/test_riscv_operator_workflow.py new file mode 100644 index 0000000000..3a5a8dac08 --- /dev/null +++ b/testing/python/riscv/test_riscv_operator_workflow.py @@ -0,0 +1,189 @@ +from __future__ import annotations + +import importlib.util +import json +import os +import shutil +import subprocess +import sys +from pathlib import Path + +import pytest + + +REPO_ROOT = Path(__file__).resolve().parents[3] +CATALOG_PATH = REPO_ROOT / ".agents" / "riscv_operator_workflow" / "operators.json" +TASK_ROOT = REPO_ROOT / ".agents" / "tasks" / "riscv" +GENERATOR_PATH = REPO_ROOT / "maint" / "scripts" / "generate_riscv_operator_task.py" +VALIDATION_SCRIPT = REPO_ROOT / "maint" / "scripts" / "run_riscv_operator_workflow_validation.sh" + +REQUIRED_OPERATOR_FIELDS = { + "name", + "category", + "task_title", + "semantics", + "inputs", + "outputs", + "baseline_reference", + "baseline_implementation", + "current_example", + "suggested_files", + "schedule_strategy", + "compile_commands", + "test_commands", + "success_criteria", +} +REQUIRED_TASK_SECTIONS = ( + "## Goal", + "## Operator Semantics", + "## Inputs", + "## Outputs", + "## Baseline Reference", + "## Baseline Implementation", + "## Context Files", + "## Scheduling Strategy", + "## Compilation Commands", + "## Test Commands", + "## Success Criteria", + "## Failure Triage Buckets", + "## Existing SG2044 Baseline", + "## Reporting Template", + "## Target Flow", +) + + +def _load_catalog() -> dict: + with CATALOG_PATH.open(encoding="utf-8") as handle: + return json.load(handle) + + +def _load_generator(): + spec = importlib.util.spec_from_file_location("generate_riscv_operator_task", GENERATOR_PATH) + assert spec is not None + assert spec.loader is not None + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + + +def test_riscv_operator_catalog_covers_initial_validation_patterns(): + catalog = _load_catalog() + operators = catalog["operators"] + + assert catalog["target"] == "riscv" + assert catalog["target_alias"] == "linalg_riscv" + assert {operator["name"] for operator in operators} == {"vector_add", "reduce_sum", "matmul"} + assert {operator["category"] for operator in operators} == {"elementwise", "reduction", "matrix"} + + for operator in operators: + assert REQUIRED_OPERATOR_FIELDS <= set(operator) + for field in REQUIRED_OPERATOR_FIELDS - {"name", "category", "task_title", "semantics", "current_example"}: + assert operator[field], f"{operator['name']} has empty {field}" + + +def test_riscv_operator_context_files_exist(): + catalog = _load_catalog() + common = catalog["common_context"] + common_paths = [*common["primary_docs"], *common["reference_tests"]] + + for operator in catalog["operators"]: + paths = [operator["current_example"], *operator["suggested_files"], *common_paths] + for path in paths: + assert (REPO_ROOT / path).is_file(), f"{operator['name']} references missing context file {path}" + + +def test_riscv_operator_tasks_are_up_to_date_with_generator(): + catalog = _load_catalog() + generator = _load_generator() + + for operator in catalog["operators"]: + expected = generator.render_task(catalog, operator) + actual_path = TASK_ROOT / f"{operator['name']}.md" + assert actual_path.is_file() + assert actual_path.read_text(encoding="utf-8") == expected + + +def test_riscv_operator_generated_tasks_include_required_sections(tmp_path): + result = subprocess.run( + [sys.executable, str(GENERATOR_PATH), "--all", "--output-dir", str(tmp_path)], + cwd=REPO_ROOT, + text=True, + capture_output=True, + check=False, + ) + assert result.returncode == 0, result.stderr + + generated = sorted(path.name for path in tmp_path.glob("*.md")) + assert generated == ["matmul.md", "reduce_sum.md", "vector_add.md"] + for path in tmp_path.glob("*.md"): + content = path.read_text(encoding="utf-8") + for section in REQUIRED_TASK_SECTIONS: + assert section in content, f"{path.name} is missing {section}" + + +def test_riscv_operator_generator_rejects_unknown_operator(): + result = subprocess.run( + [ + sys.executable, + str(GENERATOR_PATH), + "--operator", + "not_a_real_operator", + ], + cwd=REPO_ROOT, + text=True, + capture_output=True, + check=False, + ) + + assert result.returncode != 0 + assert "Unknown operator" in result.stderr + assert "vector_add" in result.stderr + assert "reduce_sum" in result.stderr + assert "matmul" in result.stderr + + +def test_riscv_operator_validation_script_is_bash_parseable(): + if os.name == "nt": + pytest.skip("bash -n validation is covered on Linux/WSL") + + bash = shutil.which("bash") + if bash is None: + pytest.skip("bash is not available") + + result = subprocess.run( + [bash, "-n", str(VALIDATION_SCRIPT)], + cwd=REPO_ROOT, + text=True, + capture_output=True, + check=False, + ) + assert result.returncode == 0, result.stderr + + +def test_riscv_operator_validation_script_rejects_unknown_operator(): + if os.name == "nt": + pytest.skip("bash validation is covered on Linux/WSL") + + bash = shutil.which("bash") + if bash is None: + pytest.skip("bash is not available") + + env = os.environ.copy() + env.update({ + "PYTHON": sys.executable, + "TILELANG_RISCV_RUN_PYTEST": "0", + "TILELANG_RISCV_RUN_QEMU": "0", + "TILELANG_RISCV_RUN_HOST": "1", + "TILELANG_RISCV_EMIT_ARTIFACTS": "0", + "TILELANG_RISCV_OPERATORS": "not_a_real_operator", + }) + result = subprocess.run( + [bash, str(VALIDATION_SCRIPT)], + cwd=REPO_ROOT, + text=True, + capture_output=True, + check=False, + env=env, + ) + assert result.returncode != 0 + assert "Unknown TILELANG_RISCV_OPERATORS entries: not_a_real_operator" in result.stderr diff --git a/testing/python/riscv/test_riscv_operator_workflow_examples.py b/testing/python/riscv/test_riscv_operator_workflow_examples.py new file mode 100644 index 0000000000..ee9ba0bc7f --- /dev/null +++ b/testing/python/riscv/test_riscv_operator_workflow_examples.py @@ -0,0 +1,95 @@ +from __future__ import annotations + +import os +import subprocess +import sys +from pathlib import Path + +import pytest + + +pytest.importorskip("tilelang.tladapter._native") + + +REPO_ROOT = Path(__file__).resolve().parents[3] +EXAMPLE_ROOT = REPO_ROOT / "examples" / "riscv" +INITIAL_WORKFLOW_EXAMPLES = ( + "example_vector_add.py", + "example_reduce_sum.py", + "example_matmul.py", +) + + +def _env() -> dict[str, str]: + env = os.environ.copy() + env["PYTHONPATH"] = str(REPO_ROOT) + ld_parts = [ + str(REPO_ROOT / "build" / "lib"), + str(REPO_ROOT / "3rdparty" / "llvm-project" / "install" / "lib"), + ] + if env.get("LD_LIBRARY_PATH"): + ld_parts.append(env["LD_LIBRARY_PATH"]) + env["LD_LIBRARY_PATH"] = ":".join(ld_parts) + return env + + +@pytest.mark.parametrize("example_name", INITIAL_WORKFLOW_EXAMPLES) +def test_riscv_operator_workflow_examples_run_on_host(example_name, tmp_path): + example_path = EXAMPLE_ROOT / example_name + assert example_path.is_file(), example_name + + result = subprocess.run( + [sys.executable, str(example_path), "--run-host", "--output-dir", str(tmp_path / example_name)], + cwd=REPO_ROOT, + env=_env(), + text=True, + capture_output=True, + check=False, + ) + if result.returncode != 0: + raise AssertionError(f"{example_name} failed:\nstdout:\n{result.stdout}\nstderr:\n{result.stderr}") + assert "host check passed" in result.stdout + + +@pytest.mark.parametrize("example_name", INITIAL_WORKFLOW_EXAMPLES) +def test_riscv_operator_workflow_examples_emit_full_artifacts(example_name, tmp_path): + example_path = EXAMPLE_ROOT / example_name + assert example_path.is_file(), example_name + + output_dir = tmp_path / example_name + result = subprocess.run( + [ + sys.executable, + str(example_path), + "--emit-mlir", + "--emit-llvm", + "--emit-asm", + "--emit-object", + "--output-dir", + str(output_dir), + ], + cwd=REPO_ROOT, + env=_env(), + text=True, + capture_output=True, + check=False, + ) + if result.returncode != 0: + raise AssertionError(f"{example_name} failed:\nstdout:\n{result.stdout}\nstderr:\n{result.stderr}") + + mlir_candidates = list(output_dir.glob("*.mlir")) + llvm_candidates = list(output_dir.glob("*.ll")) + asm_candidates = list(output_dir.glob("*.s")) + obj_candidates = list(output_dir.glob("*.o")) + if len(mlir_candidates) != 1: + raise AssertionError(f"{example_name} should emit exactly one MLIR file, found {mlir_candidates}") + if len(llvm_candidates) != 1: + raise AssertionError(f"{example_name} should emit exactly one LLVM IR file, found {llvm_candidates}") + if len(asm_candidates) != 1: + raise AssertionError(f"{example_name} should emit exactly one asm file, found {asm_candidates}") + if len(obj_candidates) != 1: + raise AssertionError(f"{example_name} should emit exactly one object file, found {obj_candidates}") + assert "func.func" in mlir_candidates[0].read_text() + assert "define void" in llvm_candidates[0].read_text() + assert asm_candidates[0].stat().st_size > 0 + assert obj_candidates[0].stat().st_size > 0 diff --git a/testing/python/riscv/test_riscv_qemu_smoke.py b/testing/python/riscv/test_riscv_qemu_smoke.py index ed294f5c18..3fc9510069 100644 --- a/testing/python/riscv/test_riscv_qemu_smoke.py +++ b/testing/python/riscv/test_riscv_qemu_smoke.py @@ -11,7 +11,11 @@ REPO_ROOT = Path(__file__).resolve().parents[3] -EXAMPLE = REPO_ROOT / "examples" / "riscv" / "example_vector_add.py" +EXAMPLES = ( + ("vector_add", REPO_ROOT / "examples" / "riscv" / "example_vector_add.py"), + ("reduce_sum", REPO_ROOT / "examples" / "riscv" / "example_reduce_sum.py"), + ("matmul", REPO_ROOT / "examples" / "riscv" / "example_matmul.py"), +) def _env() -> dict[str, str]: @@ -27,18 +31,20 @@ def _env() -> dict[str, str]: return env -def test_riscv_qemu_smoke_example_is_present(): - assert EXAMPLE.is_file() +@pytest.mark.parametrize("example_name,example_path", EXAMPLES) +def test_riscv_qemu_smoke_example_is_present(example_name, example_path): + assert example_path.is_file(), example_name -def test_riscv_example_vector_add_runs_on_qemu(tmp_path): +@pytest.mark.parametrize("example_name,example_path", EXAMPLES) +def test_riscv_examples_run_on_qemu(example_name, example_path, tmp_path): pytest.importorskip("tilelang.tladapter._native") if resolve_riscv_runner(required=False) is None: pytest.skip("qemu/spike runner not available on this machine") - output_dir = tmp_path / "vector_add_qemu" + output_dir = tmp_path / f"{example_name}_qemu" result = subprocess.run( - [sys.executable, str(EXAMPLE), "--run-qemu", "--output-dir", str(output_dir)], + [sys.executable, str(example_path), "--run-qemu", "--output-dir", str(output_dir)], cwd=REPO_ROOT, env=_env(), text=True, @@ -46,7 +52,7 @@ def test_riscv_example_vector_add_runs_on_qemu(tmp_path): check=False, ) if result.returncode != 0: - raise AssertionError(f"qemu smoke failed:\nstdout:\n{result.stdout}\nstderr:\n{result.stderr}") + raise AssertionError(f"{example_name} qemu smoke failed:\nstdout:\n{result.stdout}\nstderr:\n{result.stderr}") assert "qemu check passed" in result.stdout - assert (output_dir / "vector_add.qemu.elf").is_file() + assert (output_dir / f"{example_name}.qemu.elf").is_file() diff --git a/tilelang/jit/adapter/riscv/libgen.py b/tilelang/jit/adapter/riscv/libgen.py index 1ab36914d9..48f3b99593 100644 --- a/tilelang/jit/adapter/riscv/libgen.py +++ b/tilelang/jit/adapter/riscv/libgen.py @@ -3,6 +3,7 @@ from __future__ import annotations import os +import shlex import subprocess import tempfile from pathlib import Path @@ -13,6 +14,7 @@ DEFAULT_RISCV_TRIPLE = os.environ.get("TILELANG_RISCV_TRIPLE", "riscv64-unknown-linux-gnu") +DEFAULT_RISCV_LLC_MATTR = os.environ.get("TILELANG_RISCV_LLC_MATTR", "+m,+a,+f,+d,+c") def _extract_mlir_source(value: Any) -> str: @@ -111,16 +113,17 @@ def _emit_llc_artifact( ll_path = temp_dir_path / "kernel.ll" out_path = temp_dir_path / f"kernel.{ 's' if filetype == 'asm' else 'o' }" ll_path.write_text(llvm_ir) - _run_checked( - [ - str(resolve_tool("llc")), - f"-mtriple={triple}", - f"-filetype={filetype}", - str(ll_path), - "-o", - str(out_path), - ] - ) + cmd = [ + str(resolve_tool("llc")), + f"-mtriple={triple}", + f"-filetype={filetype}", + ] + if DEFAULT_RISCV_LLC_MATTR: + cmd.append(f"-mattr={DEFAULT_RISCV_LLC_MATTR}") + if os.environ.get("TILELANG_RISCV_LLC_FLAGS"): + cmd.extend(shlex.split(os.environ["TILELANG_RISCV_LLC_FLAGS"])) + cmd.extend([str(ll_path), "-o", str(out_path)]) + _run_checked(cmd) if filetype == "asm": text = out_path.read_text() if path is not None: diff --git a/tilelang/jit/adapter/riscv/wrapper.py b/tilelang/jit/adapter/riscv/wrapper.py index 36422eac7c..c521c8038e 100644 --- a/tilelang/jit/adapter/riscv/wrapper.py +++ b/tilelang/jit/adapter/riscv/wrapper.py @@ -536,6 +536,8 @@ def build_qemu_executable( str(out_path), ] cmd.extend(_infer_riscv_abi_flag(obj_path)) + if not os.environ.get("TILELANG_RISCV_MARCH"): + cmd.append("-march=rv64gc") cmd.extend(_riscv_clang_flags()) if clang_flags: cmd.extend(clang_flags)