Test unaligned temp storage#9746
Conversation
| // randomly offset the temporary storage address by one byte | ||
| const int offset = GENERATE(random(0, 1)); | ||
| c2h::device_vector<std::uint8_t> temp_storage(temp_storage_bytes + offset, thrust::no_init); | ||
|
|
||
| error = action(thrust::raw_pointer_cast(temp_storage.data()), temp_storage_bytes, args...); | ||
| error = action(thrust::raw_pointer_cast(temp_storage.data()) + offset, temp_storage_bytes, args...); |
There was a problem hiding this comment.
I don't mind this, but my issue with it is you may have a case where we don't test the regression this PR aims to fix. It might be better to add an explicit test case for scan that has unaligned temp storage
There was a problem hiding this comment.
I deem the probability of ALL calls to cub::DeviceScan across all unit tests across all CI jobs generating ONLY an offset of 0 so low that I believe it's almost certain we are going to cover it ;)
There was a problem hiding this comment.
The reason I didn't add it for scan explicitly is, that we do actually need this test for every CUB device-scope algorithm. And for a second I considered asking claude to add those, until I noticed that every CUB device-scope algorithm goes through the launch helper, so there is a single point where we could test this for every algorithm at essentially no additional compile and runtime cost (there is no new scan instantiation and no additional test, we just perturb the existing tests a bit). I found that elegant.
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
📝 WalkthroughSummary by CodeRabbit
WalkthroughThe host-side test launch helper now randomly offsets temporary storage by zero or one byte before invoking the action, while preserving existing CUDA error checks and synchronization. ChangesMisaligned temporary storage coverage
Assessment against linked issues
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cub/cub/device/dispatch/dispatch_scan.cuh (1)
502-506: 🩺 Stability & Availability | 🔴 Critical | ⚡ Quick wincritical:
DispatchScanstill underreports temp storage here.aligned_temp_storage = ::cuda::align_up(...)can advance by up toalignment - 1bytes, so allocating exactlytemp_storage_bytescan overrun the buffer. Add the same padding used ininvoke_lookahead.
🧹 Nitpick comments (1)
cub/cub/device/dispatch/dispatch_scan.cuh (1)
526-526: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
aligned_temp_storageis never reassigned; declare itconst auto. As per coding guidelines: "All variables that are not modified must be declaredconst".Also applies to: 1131-1131
Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: d4810340-1fc7-45c5-9e1b-d1a8dd25f73b
📒 Files selected for processing (2)
cub/cub/device/dispatch/dispatch_scan.cuhcub/test/catch2_test_launch_helper.h
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
|
Since #9565 also solved this issue, this PR becomes obsolete for
I will limit this PR to the temporary storage test and open a new PR with a targeted fix for 3.4: #9781 |
bernhardmgruber
force-pushed
the
temp_storage_alignment
branch
from
9d54f77 to
1a33c68
Compare
bernhardmgruber
changed the title
🥳 CI Workflow Results🟩 Finished in 2h 36m: Pass: 100%/242 | Total: 11d 13h | Max: 2h 35m | Hits: 16%/915581See results here. |
|
Successfully created backport PR for |
2 participants
Follow-up to: #9742