[BugFix] Scalarize vectorized bool with more than 4 lanes

[jjppp]

• Scalarize vectorized bool with more than 4 lanes to avoid ICE during downstream codegen
• Use dedicated init values for boolean reductions

Fixes #2568, fixes #2206

Summary

• Scalarize vectorized boolean values with more than four lanes to prevent CUDA codegen failures.
• Use boolean identity values for reductions:
  • 0 for max and logical OR.
  • 1 for min, bitwise AND, and logical AND.
• Preserve existing initialization behavior for integer and floating-point reductions.
• Add CUDA regression tests for boolean reductions and coverage for int32 and float32 reductions.

C++ style / lint notes

• The PR changes C++ code in reduce.h and vectorize_loop.cc.
• No changes to the rules documented in docs/developer_guide/cpp_style.md are indicated.
• The “C++ API Style Audit (warning only)” CI step is relevant if configured for these files.
• No correctness or build issue is reported from style warnings. Advisory TLCPP003/TLCPP004 findings should not block merge unless they indicate a new API, FFI, or maintainability risk.

[github-actions]

👋 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]

📝 Walkthrough

Walkthrough

The change fixes boolean reduction identities, scalarizes wide boolean vectors during vectorization, and adds CUDA regression tests for boolean reductions and numeric reduce_max.

Changes

Boolean CUDA code generation

Layer / File(s)	Summary
Boolean reduction identities and regression coverage src/backend/common/op/reduce.h, testing/python/issue/test_tilelang_issue_2568.py	Boolean reduce_max, reduce_min, and reduce_bitand use domain-specific identity values. CUDA tests compare boolean reductions with logical reductions and PyTorch, and verify integer and floating-point reduce_max compilation.
Wide boolean vector scalarization src/transform/vectorize_loop.cc	VisitExpr requests scalarization for boolean vectors with more than four lanes and returns the original expression after scalarization is required.

Estimated code review effort: 3 (Moderate) | ~20 minutes

Possibly related PRs

• tile-ai/tilelang#2807: Both changes address boolean logical-reduction correctness, although that PR changes layout-aware reduction handling.
• tile-ai/tilelang#2845: Both changes update reduction initialization logic and add dtype-focused regression tests.

Suggested reviewers: leiwang1999, morluto

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly describes the primary change: scalarizing vectorized boolean values with more than four lanes.
Linked Issues check	✅ Passed	The changes address both linked issues by scalarizing wide boolean vectors, fixing boolean reduction identities, and adding regression coverage.
Out of Scope Changes check	✅ Passed	All code and test changes directly support the linked issues and stated objectives.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 1

🧹 Nitpick comments (1)

src/transform/vectorize_loop.cc (1)

307-314: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Use a descriptive name for the transformed expression.

Replace ret with a name such as vectorized_expr. The name must identify the expression state at the scalarization check.

As per path instructions, parameters and local variables should use descriptive lower_snake names.

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@src/transform/vectorize_loop.cc` around lines 307 - 314, Rename the local
variable ret in the ExprFunctor::VisitExpr flow to a descriptive
lower_snake_case name such as vectorized_expr, and update its uses in the
vectorized boolean scalarization check and return statement.

Source: Path instructions

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@testing/python/issue/test_tilelang_issue_2568.py`:
- Line 24: Rename the unused kernel bindings in the T.Kernel context manager
from bx and by to _bx and _by, preserving the existing kernel body and test
behavior.

---

Nitpick comments:
In `@src/transform/vectorize_loop.cc`:
- Around line 307-314: Rename the local variable ret in the
ExprFunctor::VisitExpr flow to a descriptive lower_snake_case name such as
vectorized_expr, and update its uses in the vectorized boolean scalarization
check and return statement.

🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

• Push a commit to this branch (recommended)
• Create a new PR with the fixes

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 58a28e29-ec6d-4bb7-8985-1cbba9831e26

📥 Commits

Reviewing files that changed from the base of the PR and between fef9f76 and f4b01a5.

📒 Files selected for processing (3)

• src/backend/common/op/reduce.h
• src/transform/vectorize_loop.cc
• testing/python/issue/test_tilelang_issue_2568.py

[coderabbitai]

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

Rename the unused kernel bindings.

Ruff reports bx and by as unused. Rename them to _bx and _by to keep this regression test lint-clean.

Proposed fix

-        with T.Kernel(1, 1, threads=THREADS) as (bx, by):
+        with T.Kernel(1, 1, threads=THREADS) as (_bx, _by):

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change

	with T.Kernel(1, 1, threads=THREADS) as (bx, by):
	with T.Kernel(1, 1, threads=THREADS) as (_bx, _by):

🧰 Tools

🪛 Ruff (0.16.1)

[warning] 24-24: Unpacked variable bx is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

---

[warning] 24-24: Unpacked variable by is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@testing/python/issue/test_tilelang_issue_2568.py` at line 24, Rename the
unused kernel bindings in the T.Kernel context manager from bx and by to _bx and
_by, preserving the existing kernel body and test behavior.

Source: Linters/SAST tools

[jjppp]

I think there might be multiple (at least two different) ways to support vectorized bools (e.g., boolx8).

1. Scalarize for specific lanes
2. Extend codegen to delegate to existing types (e.g., int8x8) via bit-packing

For me, the first approach seems most straightforward and simplest to implement.
I'm not a kernel expert, so I'm not certain about its performance impact (though it is not originally supported, ;-P).
Would be glad to hear some thoughts on this.

---

Some notes on the CI failure: since #2874 we have ROCm CI re-enabled, and #2889 fixes an ROCm bug that is missed by #2843

[andyluo7]

@jjppp , the failing ROCm check here isn't caused by this PR. test_tilelang_issue_2682.py::test_vectorized_select is currently failing on main for everyone on gfx942 (1 failed, 1830 passed in your run) — it's a HIP codegen gap from
#2843. #2889 fixes it and is green on the gfx942 runner. Once that merges, please rebase and the AMD leg should go clean.