[CUB] Fix DeviceAdjacentDifference for cuda::device_buffer iterators - #9861
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthrough
ChangesAdjacent Difference Iterator Compatibility
Assessment against linked issues
Possibly related PRs
Suggested reviewers: 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: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: de06a381-b09a-4138-9843-f913ccbe3dac
📒 Files selected for processing (1)
cub/test/catch2_test_device_adjacent_difference_substract_left.cu
|
Addressed the review feedback in the latest commits:
The target build, focused regression test, and targeted pre-commit checks pass locally. |
| void* temporary_storage = nullptr; | ||
| std::size_t temporary_bytes = 0; | ||
| REQUIRE( | ||
| cudaSuccess | ||
| == cub::DeviceAdjacentDifference::SubtractLeftCopy( | ||
| temporary_storage, temporary_bytes, input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get())); | ||
|
|
||
| c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init); | ||
| temporary_storage = thrust::raw_pointer_cast(temporary_buffer.data()); | ||
| REQUIRE( | ||
| cudaSuccess | ||
| == cub::DeviceAdjacentDifference::SubtractLeftCopy( | ||
| temporary_storage, temporary_bytes, input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get())); |
There was a problem hiding this comment.
Important: Just use the launch wrapper adjacent_difference_subtract_left instead of handling the temporary storage allocation yourself.
There was a problem hiding this comment.
Updated the regression test to use adjacent_difference_subtract_left_copy, since this test covers SubtractLeftCopy. The launch wrapper now handles the temporary storage internally.
The focused target build, regression test, and targeted pre-commit checks pass locally.
|
|
||
| const c2h::host_vector<type> expected{2, 3, 4, 5, 6}; | ||
| REQUIRE(output == expected); |
There was a problem hiding this comment.
Important: I think we do need to sync when we use a custom stream:
| const c2h::host_vector<type> expected{2, 3, 4, 5, 6}; | |
| REQUIRE(output == expected); | |
| stream.sync(); | |
| const c2h::host_vector<type> expected{2, 3, 4, 5, 6}; | |
| REQUIRE(output == expected); |
There was a problem hiding this comment.
Added stream.sync() immediately after the custom-stream launch and before the host-side output comparison.
The updated test builds and passes locally, along with git diff --check and the targeted pre-commit checks.
|
/ok to test 4ced5e9 |
This comment has been minimized.
This comment has been minimized.
|
/ok to test 14fcdc5 |
🥳 CI Workflow Results🟩 Finished in 8h 42m: Pass: 100%/287 | Total: 4d 04h | Max: 46m 14s | Hits: 55%/397258See results here. |
3 participants
Description
Closes #9859
cub::DeviceAdjacentDifference::SubtractLeftCopyfails to compile withcuda::device_buffer<T>::iteratorbecauseAgentDifferenceselectsLoadItthroughtry_make_cache_modified_iterator_tbut constructs it directly withLoadIt(input_it).This change constructs
load_itthroughtry_make_cache_modified_iterator, matching the existing type-selection path. It also adds a regression test that exercises both the temporary-storage query and execution calls withcuda::device_buffer<int>::iterator.Validation:
lid_0suite: 6/6 passedcompute-sanitizer --tool memcheck: 0 errorsChecklist