fix(cache): route null decode pages to dummy slot - #1167
Conversation
Keep fused target and DFlash decode writes consistent with the general cache-location contract so null or reclaimed pages cannot scatter writes across page zero.
|
I have read the DCO document and I hereby sign off on the commits contained in this pull request. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: a70724a610
ℹ️ About Codex in GitHub
Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".
| page_ids = tl.load(page_ptrs, mask=mask, other=0) | ||
| cache_locs = page_ids * page_size + offsets_in_page | ||
| cache_locs = tl.where(overflow, 0, cache_locs) | ||
| cache_locs = tl.where(overflow | (page_ids <= 0), 0, cache_locs) |
There was a problem hiding this comment.
When block_granularity > draft_kernel_page_size, DraftPageStaging.publish() expands a null or -1 scheduler page into kernel page IDs 0..page_ratio-1 (page_table.py:80-86); for example, the existing ratio-2 contract explicitly produces [0, 1]. Consequently, DFlash positions in the second subpage see page_ids == 1, bypass this condition, and continue scattering writes across the shared physical null page instead of the fixed dummy slot. Pass the null-page kernel span/page ratio into this kernel (or otherwise identify all expanded null IDs), and add coverage where the staged table uses a ratio greater than one.
Useful? React with 👍 / 👎.
|
Superseded by #1169 with an identical tree and a commit-level DCO sign-off. Closing this branch rather than rewriting its published history. |
1 participant
Summary
This is an independently landable correctness fix found while investigating #1132. At exact 131K it only changed average accepted length from 1.0000 to 1.0365, so it is not presented as the full acceptance-collapse fix.
Test plan
pytest -q test/runtime/test_decode_cache_loc.py test/runtime/test_sliding_cache_loc.py(11 passedon gfx950)isort --check-onlyandblack --checkon the changed Python files