Skip to content

[BugFix] Resolve partial scalar reduce barrier participation - #2777

Merged
SiriusNEO merged 5 commits into
tile-ai:mainfrom
KellyFrog:fix/deferred-scalar-allreduce-barrier
Jul 30, 2026
Merged

SiriusNEO merged 5 commits into
tile-ai:mainfrom
KellyFrog:fix/deferred-scalar-allreduce-barrier

Conversation

@KellyFrog

@KellyFrog KellyFrog commented Jul 27, 2026 •

Copy link
Copy Markdown
Collaborator

T.reduce will produce wrong barrier counts in SM >= 90.

This can be reproduced with

"""Reproduce a conditional 128-thread barrier generated by TileLang."""

from __future__ import annotations

from typing import Any

import tilelang
import tilelang.language as T

BLOCK_THREADS = 128
FRAGMENT_THREADS = 64
ROWS = 4
WIDTH = 512
VECTOR_SIZE = 8


@tilelang.jit(
    out_idx=1,
    target="cuda",
    pass_configs={
        tilelang.PassConfigKey.TL_DISABLE_DATA_RACE_CHECK: True,
        tilelang.PassConfigKey.TL_DISABLE_THREAD_STORAGE_SYNC: True,
    },
)
def make_repro_kernel() -> Any:
    def x_layout(i: int, j: int) -> tuple[int, int]:
        # j // 8 is in [0, 64), so only threads 0..63 own x_frag.
        thread_id = j // VECTOR_SIZE
        local_id = i * VECTOR_SIZE + j % VECTOR_SIZE
        return thread_id, local_id

    @T.prim_func
    def fragment_reduce(
        x: T.Tensor((ROWS, WIDTH), "float32"),
        out: T.Tensor((ROWS, WIDTH), "float32"),
    ) -> None:
        with T.Kernel(1, threads=BLOCK_THREADS):
            x_frag = T.alloc_fragment((ROWS, WIDTH), "float32")
            sum_frag = T.alloc_fragment((ROWS,), "float32")

            T.annotate_layout(
                {
                    x_frag: T.Fragment(x_frag.shape, forward_fn=x_layout),
                    sum_frag: T.Fragment(
                        sum_frag.shape,
                        forward_fn=lambda i, rep: (rep, i),
                        replicate=FRAGMENT_THREADS,
                    ),
                }
            )

            # T.Parallel -> T.reduce_sum -> T.Parallel is the trigger.
            for i, j in T.Parallel(ROWS, WIDTH):
                x_frag[i, j] = x[i, j]

            T.reduce_sum(x_frag, sum_frag, dim=1)

            for i, j in T.Parallel(ROWS, WIDTH):
                out[i, j] = x_frag[i, j] / sum_frag[i]

    return fragment_reduce


def main() -> None:
    kernel = make_repro_kernel()
    print(kernel.get_kernel_source())


if __name__ == "__main__":
    main()

Now this issue is fixed by introducing intermediate representation in pass LowerTileOp.

For convince, reduces that use non-uniform registers layouts in a warp (e.g. using registers only in threads that tx % 2 == 0) or use registers in non-consecutive warps (e.g. tx [0, 32) and tx [64, 92)) will now be rejected by the compiler.

Extra tests over T.reduce are added.

2 bugs are observed but not fixed in current PR:

  • The same issue also occurs in SM < 90, where the default sync method is __syncthreads and will require all threads in a block to participate in T.reduce.
  • Artificial warp specialization (e.g. if tx < 32: T.reduce(A) elif tx < 64 T.reduce(B)) calling T.reduce at the same time will interfere each other's barriers as the current code hardcodes barrier 1 and 2 in AllReduce.

They are expected to be fixed in other 2 PRs.

Summary

  • Fixes incorrect T.reduce barrier counts on SM90+ for reductions involving only part of a scalar fragment.
  • Resolves participating threads during lowering and supports warp-aligned contiguous partial-CTA reductions.
  • Rejects non-uniform warp layouts, non-consecutive warp allocations, invalid thread ranges, and multiple partial scalar reductions that would collide on shared barriers/workspace.
  • Updates CUDA barrier generation and keeps ROCm interfaces compatible.
  • Adds extensive correctness and rejection tests, including coverage for partial warps, offsets, disjoint groups, and pre-SM90 behavior.

C++ style / lint notes

  • C++ sources and public lowering interfaces are modified; no changes to docs/developer_guide/cpp_style.md are indicated.
  • The C++ API Style Audit (warning only) remains advisory. Any TLCPP003/TLCPP004 findings should be treated as warning-only unless they reveal a concrete API or maintainability risk.
  • Correctness limitations remain for SM<90 block-wide synchronization and artificial warp specialization with concurrent reductions.

@github-actions

Copy link
Copy Markdown

👋 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

coderabbitai Bot commented Jul 27, 2026 •

Copy link
Copy Markdown
Contributor

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

Scalar AllReduce lowering now resolves partial warp-aligned CTA participation on SM90+, tracks per-kernel partial reductions, and emits matching CUDA named-barrier participant counts. ROCm signatures remain compatible, while CUDA tests cover supported and rejected layouts.

Changes

Partial scalar AllReduce barriers

Layer / File(s) Summary
Resolve and propagate partial barriers
src/backend/common/op/reduce.h, src/op/operator.h, src/transform/lower_tile_op.cc
Scalar AllReduce lowering resolves contiguous warp-aligned participation, adjusts thread offsets, tracks partial reductions, and forwards participant counts.
Emit backend barrier parameters
src/cuda/op/reduce.cc, src/rocm/op/reduce.cc
CUDA emits participant-specific named barriers; ROCm accepts the extended scalar AllReduce signature.
Validate partial reduction behavior
testing/python/language/test_tilelang_language_reduce.py
CUDA tests cover partial, grouped, offset, and partial-warp reductions, rejected layouts, and pre-SM90 behavior.

Estimated code review effort: 4 (Complex) | ~45 minutes

Possibly related PRs

Suggested reviewers: leiwang1999, siriusneo

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 18.33% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title clearly matches the main change: fixing partial scalar reduce barrier participation on lowering.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
✨ 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

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

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

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 (1)
src/backend/common/op/reduce.h (1)

1121-1127: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Attribute key is a cross-file contract duplicated as a raw literal. The producer writes the raw string while the consumer declares its own constexpr copy; a typo on either side silently disables the barrier fix instead of failing loudly.

  • src/backend/common/op/reduce.h#L1121-L1127: use a shared constant instead of the inline "tl.deferred_scalar_allreduce_barrier" literal.
  • src/transform/lower_tile_op.cc#L45-L46: move kDeferredScalarAllReduceBarrier into a shared header (e.g. alongside the other tl::attr keys) and reference it from both sites.
🤖 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/backend/common/op/reduce.h` around lines 1121 - 1127, The attribute key
is duplicated as a raw literal and a local constexpr, allowing producer and
consumer names to diverge. In src/transform/lower_tile_op.cc lines 45-46, move
kDeferredScalarAllReduceBarrier into a shared header alongside the existing
tl::attr keys; in src/backend/common/op/reduce.h lines 1121-1127, replace the
inline string with that shared constant, and update the consumer to reference
the same definition.
🤖 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 `@src/transform/lower_tile_op.cc`:
- Around line 55-73: The deferred resolver currently binds thread_var_ to [0,
block_threads) instead of the actual threadBounds range, breaking nonzero
threadBounds.min cases. Update the constructor and its call from the resolver
setup to accept the complete thread_bounds range, bind analyzer_ using that
range, and validate the computed lower bound against thread_bounds.min.

---

Nitpick comments:
In `@src/backend/common/op/reduce.h`:
- Around line 1121-1127: The attribute key is duplicated as a raw literal and a
local constexpr, allowing producer and consumer names to diverge. In
src/transform/lower_tile_op.cc lines 45-46, move kDeferredScalarAllReduceBarrier
into a shared header alongside the existing tl::attr keys; in
src/backend/common/op/reduce.h lines 1121-1127, replace the inline string with
that shared constant, and update the consumer to reference the same definition.
🪄 Autofix (Beta)

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: 2e0adc97-a075-4294-ac40-72a686375a0f

📥 Commits

Reviewing files that changed from the base of the PR and between fc517bd and eeba659.

📒 Files selected for processing (3)
  • src/backend/common/op/reduce.h
  • src/transform/lower_tile_op.cc
  • testing/python/language/test_tilelang_language_reduce.py

Comment thread src/transform/lower_tile_op.cc Outdated
@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from eeba659 to 2d4bcd7 Compare July 27, 2026 08:00

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Actionable comments posted: 1

🤖 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/language/test_tilelang_language_reduce.py`:
- Around line 279-287: Update
test_reduce_partial_thread_barrier_rejects_disjoint_groups to escape the dot in
the pytest.raises match pattern, using a raw or equivalent escaped regex so it
matches the literal “threadIdx.x” text emitted by LowerTileOp.
🪄 Autofix (Beta)

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: 86384193-9dc5-4940-bffc-735ee7c4ca2f

📥 Commits

Reviewing files that changed from the base of the PR and between eeba659 and 2d4bcd7.

📒 Files selected for processing (3)
  • src/backend/common/op/reduce.h
  • src/transform/lower_tile_op.cc
  • testing/python/language/test_tilelang_language_reduce.py
🚧 Files skipped from review as they are similar to previous changes (2)
  • src/backend/common/op/reduce.h
  • src/transform/lower_tile_op.cc

Comment thread testing/python/language/test_tilelang_language_reduce.py Outdated
@LeiWang1999
LeiWang1999 self-requested a review July 27, 2026 11:42
Comment thread testing/python/language/test_tilelang_language_reduce.py Outdated
@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from 2d4bcd7 to 72e192f Compare July 28, 2026 07:21

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Actionable comments posted: 3

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)
testing/python/language/test_tilelang_language_reduce.py (1)

1109-1109: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

Use raw strings for the match= patterns to clear RUF043.

The .* is intentional here, but Ruff still flags the non-raw literal; making it raw keeps CI clean with no behavior change. Same rule was raised on an earlier commit for another match= pattern in this file.

🔧 Proposed fix
-    with pytest.raises(Exception, match="logical_width.*positive power of two"):
+    with pytest.raises(Exception, match=r"logical_width.*positive power of two"):

Also applies to: 1137-1137

🤖 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/language/test_tilelang_language_reduce.py` at line 1109,
Update the pytest.raises match patterns at the affected locations in the
reduction tests to use raw string literals, preserving the existing regular
expressions and exception behavior while clearing RUF043.

Source: Linters/SAST tools

🧹 Nitpick comments (4)
src/config.h (1)

52-54: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Use kMaxNamedBarrier instead of the literal 15.

This PR introduces tl::kMaxNamedBarrier in src/transform/common/thread_sync_types.h; referencing it here keeps the hardware bound in one place.

♻️ Proposed change
-  if (start < 2 || start > 15) {
-    LOG(FATAL) << "tl.named_barrier_start must be in [2, 15], got " << start;
+  if (start < 2 || start > kMaxNamedBarrier) {
+    LOG(FATAL) << "tl.named_barrier_start must be in [2, " << kMaxNamedBarrier
+               << "], got " << start;
   }

Requires including transform/common/thread_sync_types.h here.

🤖 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/config.h` around lines 52 - 54, Update the start-range validation in the
configuration check to compare against tl::kMaxNamedBarrier instead of the
literal 15, and include transform/common/thread_sync_types.h so the symbol is
available. Keep the existing lower bound and fatal error behavior unchanged.
src/backend/common/op/reduce.h (1)

207-238: 🚀 Performance & Scalability | 🔵 Trivial | ⚖️ Poor tradeoff

Exhaustive enumeration can cost up to ~1M Substitute + Simplify calls at compile time.

kMaxCombos = 1 << 20 bounds correctness but not compile time; each combination runs a full TIR substitution and simplification. For wide fragment layouts this can dominate lowering. Consider deriving the image analytically (or memoizing/short-circuiting once contiguity and span are established), or lowering the cap to something defensible with a clear error.

🤖 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/backend/common/op/reduce.h` around lines 207 - 238, Replace the
exhaustive enumeration in the reduce-layout image analysis with a bounded-cost
approach: derive the image analytically or memoize/short-circuit once contiguity
and span are established, rather than invoking Substitute and Simplify for every
combination. Update the kMaxCombos handling to enforce a compile-time-defensible
limit and retain a clear fatal error when the layout cannot be resolved safely.
testing/python/transform/test_tilelang_transform_thread_sync.py (1)

821-822: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

requires_cuda looks unnecessary here.

This test only runs the ThreadSync pass over an IRModule with no target attribute and no device execution, unlike the sibling tests at Lines 767 and 795 which are ungated. Dropping the gate keeps this coverage alive on CPU-only CI. The duplicated func body with test_partial_sync_warp_multiple_still_lowered could also be shared via a small helper.

🤖 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/transform/test_tilelang_transform_thread_sync.py` around lines
821 - 822, Remove the requires_cuda decorator from
test_partial_sync_honors_named_barrier_start_attr so the IR-only ThreadSync
coverage runs on CPU-only CI. Keep the test’s existing behavior unchanged;
optionally factor its duplicated func setup with
test_partial_sync_warp_multiple_still_lowered only if a small local helper
preserves readability.
testing/python/language/test_tilelang_language_reduce.py (1)

592-592: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Barrier-ID regex assumes a literal single space; line 942 in this same file uses ,\s*.

If codegen emits NamedBarrier<64,1> (no space), re.findall returns [] and these assertions fail for the wrong reason. Align with the tolerant pattern already used in test_finalize_reducer_codegen, and hoist it into one module-level constant.

♻️ Proposed change
-    ids = [int(value) for value in re.findall(r"NamedBarrier<\d+, (\d+)>", src)]
+    ids = [int(value) for value in re.findall(_BARRIER_ID_RE, src)]

Add near the top of the module:

_BARRIER_ID_RE = r"NamedBarrier<\s*\d+\s*,\s*(\d+)\s*>"

Also applies to: 605-605, 610-610

🤖 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/language/test_tilelang_language_reduce.py` at line 592, Define
a module-level _BARRIER_ID_RE using whitespace-tolerant matching for
NamedBarrier IDs, then replace the inline regex usages in the affected reduce
tests and test_finalize_reducer_codegen with this shared constant. Preserve the
existing capture of the barrier ID while accepting codegen output with or
without spaces.
🤖 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 `@src/backend/common/op/reduce.h`:
- Around line 223-242: Update the terminal branch of the enumerate lambda so any
substituted thread expression that does not simplify to a constant causes a
compile-time fatal error instead of being omitted from image. Preserve insertion
of valid constant values, and include sufficient context in the diagnostic to
identify the unresolved thread-image expression.
- Around line 1178-1186: Update the batched reduction path around
Impl::MakeBatchAllReduce to set reduce::AllReduceBarrier participants for
partial CTA thread images, matching the scalar reduction handling and using the
actual reducing-thread count. Ensure named-barrier generation consumes this
resolved participant count rather than defaulting to the full thread image,
while preserving existing behavior for full CTAs and unsupported targets.

In `@src/transform/thread_storage_sync.cc`:
- Around line 330-338: Update the barrier-ID validation in the thread storage
synchronization flow to reject barrier_id == 0 in addition to IDs exceeding
kMaxNamedBarrier. Preserve the existing fatal error path and message context,
ensuring the reserved __syncthreads barrier cannot be selected.

---

Outside diff comments:
In `@testing/python/language/test_tilelang_language_reduce.py`:
- Line 1109: Update the pytest.raises match patterns at the affected locations
in the reduction tests to use raw string literals, preserving the existing
regular expressions and exception behavior while clearing RUF043.

---

Nitpick comments:
In `@src/backend/common/op/reduce.h`:
- Around line 207-238: Replace the exhaustive enumeration in the reduce-layout
image analysis with a bounded-cost approach: derive the image analytically or
memoize/short-circuit once contiguity and span are established, rather than
invoking Substitute and Simplify for every combination. Update the kMaxCombos
handling to enforce a compile-time-defensible limit and retain a clear fatal
error when the layout cannot be resolved safely.

In `@src/config.h`:
- Around line 52-54: Update the start-range validation in the configuration
check to compare against tl::kMaxNamedBarrier instead of the literal 15, and
include transform/common/thread_sync_types.h so the symbol is available. Keep
the existing lower bound and fatal error behavior unchanged.

In `@testing/python/language/test_tilelang_language_reduce.py`:
- Line 592: Define a module-level _BARRIER_ID_RE using whitespace-tolerant
matching for NamedBarrier IDs, then replace the inline regex usages in the
affected reduce tests and test_finalize_reducer_codegen with this shared
constant. Preserve the existing capture of the barrier ID while accepting
codegen output with or without spaces.

In `@testing/python/transform/test_tilelang_transform_thread_sync.py`:
- Around line 821-822: Remove the requires_cuda decorator from
test_partial_sync_honors_named_barrier_start_attr so the IR-only ThreadSync
coverage runs on CPU-only CI. Keep the test’s existing behavior unchanged;
optionally factor its duplicated func setup with
test_partial_sync_warp_multiple_still_lowered only if a small local helper
preserves readability.
🪄 Autofix (Beta)

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: 65af5ad7-1a54-446e-a422-3d8a9e1e8ed3

📥 Commits

Reviewing files that changed from the base of the PR and between 2d4bcd7 and 72e192f.

📒 Files selected for processing (19)
  • src/backend/common/op/finalize_reducer.h
  • src/backend/common/op/reduce.h
  • src/config.h
  • src/cuda/op/finalize_reducer.cc
  • src/cuda/op/reduce.cc
  • src/op/builtin.cc
  • src/op/builtin.h
  • src/op/operator.h
  • src/rocm/op/finalize_reducer.cc
  • src/rocm/op/reduce.cc
  • src/tl_templates/cuda/reduce.h
  • src/transform/common/thread_sync_types.h
  • src/transform/lower_tile_op.cc
  • src/transform/thread_storage_sync.cc
  • testing/python/language/test_tilelang_language_reduce.py
  • testing/python/target/test_tilelang_codegen_cutedsl_reduce.py
  • testing/python/transform/test_tilelang_transform_thread_sync.py
  • tilelang/contrib/cutedsl/reduce.py
  • tilelang/transform/pass_config.py

Comment thread src/backend/common/op/reduce.h Outdated
Comment thread src/backend/common/op/reduce.h Outdated
Comment thread src/transform/thread_storage_sync.cc Outdated

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Actionable comments posted: 1

♻️ Duplicate comments (1)
src/backend/common/op/reduce.h (1)

1008-1016: 🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift

Batch AllReduce path still doesn't resolve partial-participation barriers.

Unlike the scalar path (lines 1199-1208) which calls ResolveAllReduceBarrier to compute partial/base/participants for the actual participating range, this batch path only claims a barrier_id and leaves barrier.partial=false/participants=0. If a batched reduction can also execute with only part of the CTA participating (partial scalar fragment), the named barrier emitted downstream would use the whole-CTA arrival count while only a subset of threads arrive — the same hang/premature-release hazard this PR fixes for the scalar path.

If batched reductions are guaranteed to always cover the whole CTA by construction (e.g., rejected earlier during layout validation), this is safe as-is and worth a short comment noting why; otherwise this path needs the same ResolveAllReduceBarrier treatment.

#!/bin/bash
# Check whether the batch AllReduce path can ever see partial (non-whole-CTA) participation,
# and whether MakeBatchAllReduce differentiates on barrier.partial like the scalar codegen does.
rg -n "MakeBatchAllReduce" -A 20 src/cuda/op/reduce.cc src/rocm/op/reduce.cc src/tl_templates/cuda/reduce.h 2>/dev/null
rg -n "barrier.partial|barrier.participants" src/cuda/op src/tl_templates/cuda 2>/dev/null
🤖 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/backend/common/op/reduce.h` around lines 1008 - 1016, Validate whether
the batch AllReduce path around MakeBatchAllReduce can receive partial CTA
participation; if it can, apply ResolveAllReduceBarrier to the barrier before
claiming the named barrier so partial, base, and participants reflect the actual
range. If batch reductions are guaranteed to cover the whole CTA by prior
validation, preserve the current logic and add a concise comment documenting
that invariant.
🤖 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 `@src/backend/common/op/reduce.cc`:
- Around line 119-127: Update the terminal branch of enumerate in the reduction
logic to reject substituted thread-image expressions that do not simplify to
integer constants, rather than silently skipping them. Emit the established
compile-time diagnostic or propagate failure so contiguity and warp-alignment
checks cannot proceed with an incomplete image; preserve insertion of valid
constants into image.

---

Duplicate comments:
In `@src/backend/common/op/reduce.h`:
- Around line 1008-1016: Validate whether the batch AllReduce path around
MakeBatchAllReduce can receive partial CTA participation; if it can, apply
ResolveAllReduceBarrier to the barrier before claiming the named barrier so
partial, base, and participants reflect the actual range. If batch reductions
are guaranteed to cover the whole CTA by prior validation, preserve the current
logic and add a concise comment documenting that invariant.
🪄 Autofix (Beta)

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: c947aa1d-e56d-4983-88a8-b50a2e10d18b

📥 Commits

Reviewing files that changed from the base of the PR and between 72e192f and 598c660.

📒 Files selected for processing (3)
  • CMakeLists.txt
  • src/backend/common/op/reduce.cc
  • src/backend/common/op/reduce.h

Comment thread src/backend/common/op/reduce.cc Outdated
@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from 598c660 to 1ee8bec Compare July 28, 2026 09:22

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

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)
testing/python/language/test_tilelang_language_reduce.py (1)

380-389: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Use subprocess.run(timeout=...) instead of shelling out to the timeout binary.

timeout is GNU coreutils and is absent on macOS/Windows, where the call raises FileNotFoundError. Under xfail(strict=True) that failure is indistinguishable from the hang this test is meant to observe, so the test silently reports "expected failure" for the wrong reason. The Python-level timeout is portable and makes the expectation explicit. (Ruff's S603/S607 hits here are benign — all inputs come from sys.executable/__file__.)

♻️ Proposed change
-    result = subprocess.run(
-        ["timeout", "15s", sys.executable, str(Path(__file__).resolve())],
-        env={**os.environ, "TILELANG_PRE_SM90_PARTIAL_REDUCE_CHILD": "1"},
-        capture_output=True,
-        text=True,
-    )
-    assert result.returncode == 0, (
-        f"partial scalar reduction did not complete on pre-SM90 CUDA (exit code {result.returncode}):\n{result.stderr}"
-    )
+    try:
+        result = subprocess.run(
+            [sys.executable, str(Path(__file__).resolve())],
+            env={**os.environ, "TILELANG_PRE_SM90_PARTIAL_REDUCE_CHILD": "1"},
+            capture_output=True,
+            text=True,
+            timeout=15,
+        )
+    except subprocess.TimeoutExpired:
+        pytest.fail("partial scalar reduction hung on pre-SM90 CUDA")
+    assert result.returncode == 0, (
+        f"partial scalar reduction did not complete on pre-SM90 CUDA (exit code {result.returncode}):\n{result.stderr}"
+    )
🤖 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/language/test_tilelang_language_reduce.py` around lines 380 -
389, Replace the external "timeout" wrapper in
test_reduce_partial_thread_barrier_pre_sm90_runtime with subprocess.run’s
Python-level timeout argument, keeping sys.executable, the resolved test path,
environment, and output capture unchanged. Preserve the 15-second timeout
behavior while avoiding reliance on the platform-specific timeout binary.

Source: Linters/SAST tools

src/backend/common/op/reduce.h (2)

232-243: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick win

Enumeration is authoritative; the const_int_bound cross-checks can misfire as "internal error".

analyzer.const_int_bound is an over-approximation, so a layout whose exact image is contiguous but narrower than the analyzer's bound (e.g. any thread expression the analyzer cannot tighten) trips Lines 238/243 with an unhelpful "internal error" message instead of the descriptive contiguity/alignment diagnostics. Since image is computed exactly, consider dropping these two equality checks (keeping the earlier ±inf reachability check) or demoting them to DLOG.

🤖 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/backend/common/op/reduce.h` around lines 232 - 243, Remove the ICHECK_EQ
validations comparing image_min/image_max with bound->min_value/bound->max_value
in the partial scalar AllReduce barrier logic. Treat the exactly enumerated
image range as authoritative, retain the existing image_count contiguity check
and earlier ±inf reachability validation, and preserve their descriptive
diagnostics.

193-227: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low value

Compile-time cost of the enumeration cap.

kMaxCombos = 1 << 20 permits up to ~1M Substitute + analyzer.Simplify calls per T.reduce, each allocating PrimExpr nodes — that can dominate lowering time for wide fragments. A much smaller cap (e.g. 2^16) would still cover realistic reduce layouts while failing fast on pathological ones; alternatively hoist a single Substitute of the constant map per combo without re-running full simplification.

🤖 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/backend/common/op/reduce.h` around lines 193 - 227, Reduce the
enumeration limit in the thread-image resolution logic around kMaxCombos from
2^20 to a substantially smaller bound such as 2^16, while preserving the
existing fatal error for layouts exceeding the cap. Keep the enumerate callback
and its Substitute/analyzer.Simplify behavior unchanged unless a safe
single-substitution optimization is already supported.
🤖 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/language/test_tilelang_language_reduce.py`:
- Around line 392-396: Make the child-process branch containing
_make_partial_reduce_kernel mutually exclusive with the normal
tilelang.testing.main() runner by adding an else around the original __main__
runner or exiting after the child assertions. Ensure child execution runs only
the partial-reduce validation and normal execution still invokes
tilelang.testing.main().

---

Nitpick comments:
In `@src/backend/common/op/reduce.h`:
- Around line 232-243: Remove the ICHECK_EQ validations comparing
image_min/image_max with bound->min_value/bound->max_value in the partial scalar
AllReduce barrier logic. Treat the exactly enumerated image range as
authoritative, retain the existing image_count contiguity check and earlier ±inf
reachability validation, and preserve their descriptive diagnostics.
- Around line 193-227: Reduce the enumeration limit in the thread-image
resolution logic around kMaxCombos from 2^20 to a substantially smaller bound
such as 2^16, while preserving the existing fatal error for layouts exceeding
the cap. Keep the enumerate callback and its Substitute/analyzer.Simplify
behavior unchanged unless a safe single-substitution optimization is already
supported.

In `@testing/python/language/test_tilelang_language_reduce.py`:
- Around line 380-389: Replace the external "timeout" wrapper in
test_reduce_partial_thread_barrier_pre_sm90_runtime with subprocess.run’s
Python-level timeout argument, keeping sys.executable, the resolved test path,
environment, and output capture unchanged. Preserve the 15-second timeout
behavior while avoiding reliance on the platform-specific timeout binary.
🪄 Autofix (Beta)

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: 2c511773-b003-4270-b615-68b2c3d17ce1

📥 Commits

Reviewing files that changed from the base of the PR and between 598c660 and 1ee8bec.

📒 Files selected for processing (6)
  • src/backend/common/op/reduce.h
  • src/cuda/op/reduce.cc
  • src/op/operator.h
  • src/rocm/op/reduce.cc
  • src/transform/lower_tile_op.cc
  • testing/python/language/test_tilelang_language_reduce.py

Comment thread testing/python/language/test_tilelang_language_reduce.py Outdated
@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from 8829249 to 857b881 Compare July 28, 2026 10:00
@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from 857b881 to 0f7b2b2 Compare July 28, 2026 10:11
… allreduce_barriers doc

- Inline ResolveAllReduceThreadRange body from reduce.cc into reduce.h,
  add #include <tvm/ir/with_context.h> for With<ConstraintContext> usage.
- Delete src/backend/common/op/reduce.cc (function now in header).
- Remove reduce.cc from CMakeLists.txt TILE_LANG_SRCS.
- Delete docs/compiler_internals/allreduce_barriers.md.
- Remove toctree entry from docs/index.md.
@KellyFrog
KellyFrog requested a review from SiriusNEO July 28, 2026 23:07
LeiWang1999
LeiWang1999 previously approved these changes Jul 29, 2026

@LeiWang1999 LeiWang1999 left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Overall LGTM, but likely we should fix batch reduce as well.

@KellyFrog
KellyFrog force-pushed the fix/deferred-scalar-allreduce-barrier branch from 44762c7 to 16d091e Compare July 29, 2026 05:35
LeiWang1999
LeiWang1999 previously approved these changes Jul 29, 2026
Comment thread src/backend/common/op/reduce.h

@SiriusNEO SiriusNEO left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM

@SiriusNEO
SiriusNEO merged commit 0bc1913 into tile-ai:main Jul 30, 2026
6 checks passed
@KellyFrog
KellyFrog deleted the fix/deferred-scalar-allreduce-barrier branch July 30, 2026 07:46
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants