Skip to content

Test unaligned temp storage #9746

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bernhardmgruber merged 2 commits into from
Jul 10, 2026
+4 −2
Merged

Test unaligned temp storage #9746

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d623b3f
Randomly perturb temporary storage alignment
bernhardmgruber Jul 8, 2026
1a33c68
take 1
bernhardmgruber Jul 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 4 additions & 2 deletions cub/test/catch2_test_launch_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -188,9 +188,11 @@ void launch(ActionT action, Args... args)

REQUIRE(temp_storage_bytes > 0); // required by API contract

c2h::device_vector<std::uint8_t> temp_storage(temp_storage_bytes, thrust::no_init);
// randomly offset the temporary storage address by one byte
const int offset = GENERATE(take(1, 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...);
Comment on lines +191 to +195

@bernhardmgruber bernhardmgruber Jul 8, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

🍿

@NaderAlAwar NaderAlAwar Jul 9, 2026

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.

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

@bernhardmgruber bernhardmgruber Jul 9, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

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 ;)

@bernhardmgruber bernhardmgruber Jul 9, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

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.

Comment thread
coderabbitai[bot] marked this conversation as resolved.
REQUIRE(cudaSuccess == cudaPeekAtLastError());
REQUIRE(cudaSuccess == cudaDeviceSynchronize());
REQUIRE(cudaSuccess == error);
Expand Down
Loading