Skip to content

Conversation

@LeiWang1999
Copy link
Member

@LeiWang1999 LeiWang1999 commented Oct 13, 2025

This pull request introduces a new override for handling assignment statements in the tilelang parser, specifically to support advanced assignment features such as chained writes and local.var buffers. The main change is the addition of a custom tilelang_visit_assign method for the Assign node type, which enhances assignment handling to align with the requirements of the tilelang language.

Assignment node handling improvements:

  • Added a custom tilelang_visit_assign method to override assignment behavior, supporting chained writes and assignments to local.var buffers. This includes logic for slice handling, buffer stores, and context-aware assignment to variables and buffer indices.

Summary by CodeRabbit

  • New Features

    • Improved TileLang assignment semantics: supports chained assignments and in-place updates for local variable buffers, enabling more intuitive and concise kernel code.
  • Tests

    • Added a CUDA-based test that compiles and runs a sample kernel to validate chained assignment behavior and ensure outputs match expected results.

@github-actions
Copy link

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

@coderabbitai
Copy link
Contributor

coderabbitai bot commented Oct 13, 2025

Walkthrough

Adds a new CUDA-backed TileLang test exercising chained equal assignments, and extends the TileLang parser override to handle Assign nodes with chained writes and local.var buffers, including subscript handling and read-then-write semantics. No existing public APIs are modified beyond introducing new test helper functions and a new Assign visitor.

Changes

Cohort / File(s) Summary of changes
Test: chain equal on CUDA
testing/python/language/test_tilelang_laguange_chain_equal.py
Adds chain_equal, run_chain_equal, and test_chain_equal to compile and run a kernel that assigns 1 to A[idx], B[idx], C[idx]; includes CUDA guard and script entrypoint.
Parser override: Assign handling
tilelang/language/overrides/parser.py
Introduces tilelang_visit_assign to support chained Assign targets, Subscript default slice step insertion, buffer_store for Subscript, read-then-write for Name in var table, and direct local.var buffer_store optimization for single-index [0]; falls back to general assign evaluation otherwise.

Sequence Diagram(s)

sequenceDiagram
  autonumber
  actor User
  participant Parser as Parser Override
  participant VarTbl as Variable Table
  participant Eval as RHS Evaluator
  participant Buf as Buffer Ops

  User->>Parser: Visit Assign (targets, value)
  Parser->>Parser: Validate targets / normalize slices
  Parser->>Eval: Evaluate RHS once
  loop For each LHS target
    alt LHS is Subscript
      Parser->>Buf: buffer_store(base, indices, RHS)
    else LHS is Name and in VarTbl
      Parser->>Buf: buffer_load(var)  Note over Parser,Buf: Read-then-write semantics
      alt Loaded is local.var BufferLoad with idx == [0]
        Parser->>Buf: buffer_store(local.var, [0], RHS)
      else
        Parser->>Parser: General in-place update path
      end
    else
      Parser->>Parser: Fallback eval_assign(LHS, RHS)
    end
  end
Loading

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~45 minutes

Possibly related PRs

Poem

A nibble, a hop, through lines I roam,
I chain three equals, make ones their home.
The parser learns a brand-new trick,
To store and load both fast and slick.
CUDA hums—kernels fly—
Carrots compiled, we multiply! 🥕🚀

Pre-merge checks and finishing touches

❌ Failed checks (1 warning)
Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 33.33% which is insufficient. The required threshold is 80.00%. You can run @coderabbitai generate docstrings to improve docstring coverage.
✅ Passed checks (2 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title Check ✅ Passed The title concisely and accurately summarizes the primary change by indicating that the tilelang parser now supports chained assignments like a = b = c = 1, directly reflecting the new tilelang_visit_assign override and its intent; it is clear, specific, and free of unnecessary detail or generic phrasing.
✨ Finishing touches
  • 📝 Generate docstrings
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Post copyable unit tests in a comment

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands and usage tips.

Copy link
Contributor

@coderabbitai coderabbitai bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 1

🧹 Nitpick comments (3)
tilelang/language/overrides/parser.py (1)

49-55: Align indices handling with AugAssign (wrap non-tuple slice in list).

AugAssign wraps the single index in a list; Assign currently passes a scalar. While buffer_store coerces scalars, keeping both consistent improves readability and avoids surprises.

-            else:
-                indices = self.eval_expr(lhs.slice)
+            else:
+                indices = [self.eval_expr(lhs.slice)]
testing/python/language/test_tilelang_laguange_chain_equal.py (2)

30-32: Honor dtype parameter when allocating tensors.

dtype arg is ignored; tensors are hardcoded to torch.float32. Use the provided dtype for correctness and future coverage.

-    A = torch.zeros((N,), dtype=torch.float32, device="cuda")
-    B = torch.zeros((N,), dtype=torch.float32, device="cuda")
-    C = torch.zeros((N,), dtype=torch.float32, device="cuda")
+    torch_dtype = getattr(torch, dtype)  # e.g., "float32" -> torch.float32
+    A = torch.zeros((N,), dtype=torch_dtype, device="cuda")
+    B = torch.zeros((N,), dtype=torch_dtype, device="cuda")
+    C = torch.zeros((N,), dtype=torch_dtype, device="cuda")

1-46: Nit: fix filename typo for test discoverability and consistency.

Consider renaming test_tilelang_laguange_chain_equal.py -> test_tilelang_language_chain_equal.py.

📜 Review details

Configuration used: CodeRabbit UI

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 4a229dd and 704af76.

📒 Files selected for processing (2)
  • testing/python/language/test_tilelang_laguange_chain_equal.py (1 hunks)
  • tilelang/language/overrides/parser.py (1 hunks)
🧰 Additional context used
🧬 Code graph analysis (2)
testing/python/language/test_tilelang_laguange_chain_equal.py (2)
tilelang/jit/__init__.py (1)
  • jit (244-317)
tilelang/transform/pass_config.py (1)
  • PassConfigKey (6-104)
tilelang/language/overrides/parser.py (2)
tilelang/language/ast/ir.py (1)
  • buffer_store (1263-1300)
tilelang/language/parser/parser.py (1)
  • bind_assign_value (114-160)
🪛 Ruff (0.13.3)
tilelang/language/overrides/parser.py

20-20: Possible hardcoded password assigned to argument: "token"

(S106)

Comment on lines +26 to +46
if isinstance(node.value, doc.Subscript):
check_slices = []
if isinstance(node.value.slice, doc.Slice):
check_slices = [node.value.slice]
elif isinstance(node.value.slice, doc.Tuple):
for part in node.value.slice.elts:
if isinstance(part, doc.Slice):
check_slices.append(part)
for slice_node in check_slices:
if not slice_node.step and slice_node.upper and slice_node.lower:
slice_node.step = doc.Constant(
1,
None,
1,
1,
slice_node.upper.lineno,
slice_node.upper.end_col_offset + 1,
slice_node.upper.lineno,
slice_node.upper.end_col_offset + 2,
)

Copy link
Contributor

Choose a reason for hiding this comment

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

⚠️ Potential issue | 🔴 Critical

doc.Constant constructed with invalid arg list (likely runtime error).

The doc.Constant(...) call passes 9 positional args. TVM’s doc nodes expect Constant(value, kind, lineno, col, end_lineno, end_col) (value + kind + 4 positions). The extra three 1 literals make the signature invalid and will crash when this path is hit.

Fix by removing the extraneous args and keeping only the positional span:

-                slice_node.step = doc.Constant(
-                    1,
-                    None,
-                    1,
-                    1,
-                    1,
-                    slice_node.upper.lineno,
-                    slice_node.upper.end_col_offset + 1,
-                    slice_node.upper.lineno,
-                    slice_node.upper.end_col_offset + 2,
-                )
+                slice_node.step = doc.Constant(
+                    1,
+                    None,
+                    slice_node.upper.lineno,
+                    slice_node.upper.end_col_offset + 1,
+                    slice_node.upper.lineno,
+                    slice_node.upper.end_col_offset + 2,
+                )
🤖 Prompt for AI Agents
In tilelang/language/overrides/parser.py around lines 26 to 46, the
doc.Constant(...) call is constructed with nine positional arguments (including
three extra numeric 1s) which does not match the expected signature and will
raise at runtime; replace the call so it only passes value, kind, lineno, col,
end_lineno, end_col (e.g. value=1, kind=None, lineno=slice_node.upper.lineno,
col=slice_node.upper.end_col_offset+1, end_lineno=slice_node.upper.lineno,
end_col=slice_node.upper.end_col_offset+2) — remove the extraneous 1 literals
and supply the six correct arguments in that order.

@LeiWang1999 LeiWang1999 merged commit e59e7f9 into tile-ai:main Oct 14, 2025
7 of 10 checks passed
RubiaCx pushed a commit to RubiaCx/tilelang that referenced this pull request Nov 24, 2025
…le-ai#992)

* chained assignments

* test update

* [Lint]: [pre-commit.ci] auto fixes [...]

---------

Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants