Skip to content

fix(legacy): avoid undefined 32-bit shift - #701

Open
youchengsong wants to merge 1 commit into
deepseek-ai:mainfrom
youchengsong:fix-legacy-shift32
Open

fix(legacy): avoid undefined 32-bit shift#701
youchengsong wants to merge 1 commit into
deepseek-ai:mainfrom
youchengsong:fix-legacy-shift32

Conversation

@youchengsong

@youchengsong youchengsong commented Jul 27, 2026

Copy link
Copy Markdown

Summary

  • avoid shifting a 32-bit dispatch transaction window by 32
  • explicitly clear the window after all 32 pending slots become contiguous

Why

When window == 0xffffffffu, num_empty_slots is 32. Shifting a
32-bit value by its width is undefined behavior in C++, so the compiler is
not required to produce the intended zero value. If the bitmap is not
cleared, the dispatch transaction window can retain stale completion state.

The full-window case has no remaining pending bits, so assigning zero is the
direct representation of the intended state. Other shift counts retain the
existing behavior.

Fixes #366.

Validation

  • standalone GCC 15.2 UBSan regression:
    • the old expression fails with shift exponent 32 is too large
    • the fixed expression returns 0u for the full-window case
    • shift counts 1 through 31 retain the previous result
  • bash ./format.sh
  • git diff --check upstream/main...HEAD

format.sh also reformats several pre-existing long lines in the same
legacy source file; the only behavioral change is the transaction-window
assignment.

@@ -746,7 +750,7 @@ __global__ void __launch_bounds__(((kNumDispatchRDMASenderWarps + 1 + LEGACY_NUM
if (offset == 0) {

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.

🔵 suggestion: The MR mixes the functional fix with several formatting-only line-wrapping changes across the file (e.g., notify_dispatch, combine). These appear to come from ./format.sh and are harmless, but bundling unrelated reformatting with a bugfix increases review surface and makes bisecting harder. Consider isolating the functional change from the reformatting in future MRs.

🤖 v3

@ds-review-bot

Copy link
Copy Markdown
Collaborator

🤖 ds-review-bot Code Review

v6

The substantive change correctly avoids shifting a 32-bit value by 32 while preserving the intended zero result when the transaction window is full. Other changes are formatting-only.

v4

⚠️ 未完成评审(no_result_file:模型未产出结果文件)

v3

This MR fixes an undefined-behavior bug in csrc/kernels/legacy/internode.cu where a 32-bit dispatch transaction window (window) could be shifted right by 32. When window == 0xffffffffu, num_empty_slots evaluates to 32; the previous window &gt;&gt;= num_empty_slots; shifts a 32-bit value by its full width, which is UB in C++/CUDA and may leave stale completion bits in the window bitmap. The fix window = num_empty_slots == 32 ? 0u : window &gt;&gt; num_empty_slots; explicitly assigns zero for the full-window case (the intended state: no remaining pending bits) and preserves the existing shift behavior for counts < 32. The boundary handling is correct and consistent with how num_empty_slots is computed ((~window) == 0 ? 32 : __ffs(~window) - 1): since window already ORs in the current offset bit before this block, num_empty_slots == 32 correctly corresponds to a fully contiguous/complete window. The remainder of the diff is formatting-only line-wrapping of long expressions (notify_dispatch, combine, etc.), consistent with running ./format.sh, and does not alter behavior. Overall the functional change is minimal, well-scoped, and correct.

Files reviewed: 1
Issues found: 🔵 1 suggestion
Inline comments posted: 1

⚠️ Parse warning: [v4] no_result_file:模型未产出结果文件

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.

[Suggest] Fix undefined uint32 right-shift behavior in the transaction window

2 participants