Skip to content

[BugFix] Fix absmax and abssum for uint dtypes - #2845

Open
jjppp wants to merge 3 commits into
tile-ai:mainfrom
jjppp:fix-reduce-absmax-abssum
Open

[BugFix] Fix absmax and abssum for uint dtypes#2845
jjppp wants to merge 3 commits into
tile-ai:mainfrom
jjppp:fix-reduce-absmax-abssum

Conversation

@jjppp

@jjppp jjppp commented Aug 2, 2026

Copy link
Copy Markdown
Contributor

The abs function during reduce_absmax/reduce_abssum is implemented as max(rhs, -rhs), which silently fails when rhs is unsigned.
Simply use rhs if we know it is unsigned, and fall back to max(rhs, -rhs) otherwise.

Fixes #2624

Summary

  • Fix unsigned uint32 and uint64 handling in AbsMax and AbsSum reductions.
  • Use the operand directly for unsigned dtypes.
  • Preserve Max(rhs, -rhs) for signed and floating-point dtypes.
  • Add CUDA tests for absmax and abssum against max and sum.

C++ style / lint notes

  • The change touches C++ reduction lowering.
  • It does not touch rules documented in docs/developer_guide/cpp_style.md.
  • No CI or lint configuration changes are included.
  • The C++ API Style Audit remains warning-only. No new correctness or build issue is identified.

@github-actions

github-actions Bot commented Aug 2, 2026

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 Aug 2, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 8625db07-3df2-461b-9f6e-6b1c648b5e6d

📥 Commits

Reviewing files that changed from the base of the PR and between b843ac1 and 4f1dd57.

📒 Files selected for processing (1)
  • testing/python/issue/test_tilelang_issue_2624.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • testing/python/issue/test_tilelang_issue_2624.py

📝 Walkthrough

Walkthrough

Scalar AbsSum and AbsMax reductions now avoid unsigned negation. Parameterized CUDA tests verify uint32 and uint64 absolute reductions against their non-absolute equivalents.

Changes

Unsigned absolute reductions

Layer / File(s) Summary
Unsigned absolute reduction handling and regression coverage
src/backend/common/op/reduce.h, testing/python/issue/test_tilelang_issue_2624.py
Unsigned operands are used directly for scalar absolute reductions. Signed and floating-point operands retain Max(rhs, -rhs). CUDA tests cover uint32 and uint64 absmax and abssum results.

Estimated code review effort: 2 (Simple) | ~10 minutes

Possibly related PRs

  • tile-ai/tilelang#2782: Both PRs modify reduction handling for unsigned 64-bit operands and add related CUDA reduction tests, but they target different reduction implementations.

Suggested reviewers: leiwang1999

🚥 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 identifies the fix for unsigned integer absmax and abssum reductions.
Linked Issues check ✅ Passed The lowering fix and uint32/uint64 tests satisfy the linked issue requirements for correct unsigned absmax and abssum behavior.
Out of Scope Changes check ✅ Passed All changes directly support the unsigned reduction bug fix and its regression coverage.
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

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

🤖 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_2624.py`:
- Line 17: Update the reducer selection in make_kernel to map absmax to
T.reduce_absmax, abssum to T.reduce_abssum, and sum to T.reduce_sum, while
retaining T.reduce_max as the fallback for other operations.
🪄 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: bca65cf9-fcbf-447a-9b1d-b4343dacc164

📥 Commits

Reviewing files that changed from the base of the PR and between 6b81bb8 and b843ac1.

📒 Files selected for processing (2)
  • src/backend/common/op/reduce.h
  • testing/python/issue/test_tilelang_issue_2624.py

As = T.alloc_fragment((N,), dt)
Bs = T.alloc_fragment((1,), dt)
T.copy(A, As)
(T.reduce_absmax if op == "absmax" else T.reduce_max)(As, Bs, dim=0)

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.

🎯 Functional Correctness | 🟠 Major | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail
rg -n -C 2 'reduce_(absmax|abssum|max|sum)' testing/python/issue/test_tilelang_issue_2624.py src

Repository: tile-ai/tilelang

Length of output: 6351


🏁 Script executed:

#!/bin/bash
set -euo pipefail

echo "== test file =="
cat -n testing/python/issue/test_tilelang_issue_2624.py

echo
echo "== reducer symbol definitions/usages =="
rg -n -C 3 'reduce_absmax|reduce_abssum|reduce_max|reduce_sum|warp_reduce_(abs)?(abssum|max|sum|min)' src testing/python/issue/test_tilelang_issue_2624.py python tilelang | head -n 220

Repository: tile-ai/tilelang

Length of output: 15370


🏁 Script executed:

#!/bin/bash
set -euo pipefail

echo "== tilelang/language/reduce_op.py outline/slice =="
sed -n '1,220p' tilelang/language/reduce_op.py

echo
echo "== search abssum symbols =="
rg -n -C 3 'reduce_abssum|abssum|warp_reduce_abssum|warp_reduce.*abs' .

Repository: tile-ai/tilelang

Length of output: 20569


Route each operation to the matching reducer.

abssum cases call make_kernel(op.replace("abs", ""), ...) with "sum", but make_kernel only distinguishes "absmax" and falls back to T.reduce_max for everything else. Map "absmax" to T.reduce_absmax, "abssum" to T.reduce_abssum, "sum" to T.reduce_sum, and keep the fallback on T.reduce_max.

🤖 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_2624.py` at line 17, Update the
reducer selection in make_kernel to map absmax to T.reduce_absmax, abssum to
T.reduce_abssum, and sum to T.reduce_sum, while retaining T.reduce_max as the
fallback for other operations.

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.

[BUG][Fuzzer][wrong-code] T.reduce_absmax/reduce_abssum on uint32/uint64 silently returns a wrapped 2^bits value

1 participant