[CUB] Fix DeviceAdjacentDifference for cuda::device_buffer iterators by Noperi0r · Pull Request #9861 · NVIDIA/cccl

Noperi0r · 2026-07-14T17:20:57Z

Description

cub::DeviceAdjacentDifference::SubtractLeftCopy fails to compile with cuda::device_buffer<T>::iterator because AgentDifference selects LoadIt through try_make_cache_modified_iterator_t but constructs it directly with LoadIt(input_it).

This change constructs load_it through try_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 with cuda::device_buffer<int>::iterator.

Validation:

Focused regression test: 1/1 passed
Applicable AdjacentDifference lid_0 suite: 6/6 passed
Standalone ClangCUDA 21.1.8 and NVCC 12.9.86 reproducer/control matrix: 6/6 compile-and-run combinations passed
compute-sanitizer --tool memcheck: 0 errors
Targeted pre-commit hooks: passed
Raw-pointer AdjacentDifference SASS: byte-identical before and after

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2026-07-14T17:21:00Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

coderabbitai · 2026-07-14T17:24:37Z

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: cecdb9a4-96db-4540-beef-d8674bc5cbd4

📥 Commits

Reviewing files that changed from the base of the PR and between 7aed3a4 and 8c1fa4f.

📒 Files selected for processing (1)

cub/test/catch2_test_device_adjacent_difference_substract_left.cu

📝 Walkthrough

Summary by CodeRabbit

Bug Fixes
- Improved adjacent-difference behavior when using cached input iterators by changing how the cached-modified input iterator is created, ensuring correct data loading for adjacent-difference operations.
Tests
- Added a missing CUDA device-buffer header include.
- Added a new device-buffer-based test covering subtract-left copy, including CUDA stream execution and verification of expected output values.

Walkthrough

AgentDifference now uses the cache-modified iterator helper for adjacent-difference loads. A Catch2 test covers cuda::device_buffer input for SubtractLeftCopy and validates the resulting subtraction output.

Changes

Adjacent Difference Iterator Compatibility

Layer / File(s)	Summary
Cache-modified iterator construction `cub/cub/agent/agent_adjacent_difference.cuh`	`AgentDifference` creates `load_it` through `try_make_cache_modified_iterator<Policy::LOAD_MODIFIER>`.
Device-buffer API regression coverage `cub/test/catch2_test_device_adjacent_difference_substract_left.cu`	The test includes `cuda/buffer` and validates `SubtractLeftCopy` with a `cuda::device_buffer` input and expected subtraction results.

Assessment against linked issues

Objective	Addressed	Explanation
Accept `cuda::device_buffer<T>::iterator` in `DeviceAdjacentDifference::SubtractLeftCopy`, including temporary-storage query and execution calls [`#9859`]	✅

Suggested reviewers: pauleonix

_{Comment @coderabbitai help to get the list of available commands.}

bernhardmgruber · 2026-07-14T22:08:27Z

Important: Please move the new test into the file catch2_test_device_adjacent_difference_substract_left.cu. The _api.cu files cover tested API examples for the documentation.

Moved the regression test to
catch2_test_device_adjacent_difference_substract_left.cu
and restored catch2_test_device_adjacent_difference_api.cu to its previous scope.

The relocated focused test, target build, and targeted pre-commit checks pass locally.

bernhardmgruber · 2026-07-14T22:08:40Z

+
+  cuda::stream stream{cuda::devices[0]};
+  auto input = cuda::make_device_buffer<type>(stream, cuda::devices[0], {2, 5, 9, 14, 20});
+  c2h::device_vector<type> output(input.size());


Suggested change

c2h::device_vector<type> output(input.size());

c2h::device_vector<type> output(input.size(), thrust::no_init);

Applied thrust::no_init to the output allocation in the relocated test.

bernhardmgruber · 2026-07-14T22:09:05Z

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


Suggested change

c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes);

c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init);

Applied thrust::no_init to the temporary-storage allocation in the relocated test.

coderabbitai

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

📥 Commits

Reviewing files that changed from the base of the PR and between 1b058db and facc0d5.

📒 Files selected for processing (1)

cub/test/catch2_test_device_adjacent_difference_substract_left.cu

Noperi0r · 2026-07-15T15:55:52Z

Addressed the review feedback in the latest commits:

Moved the cuda::device_buffer regression test from
catch2_test_device_adjacent_difference_api.cu to
catch2_test_device_adjacent_difference_substract_left.cu.
Restored the API test file to its previous scope.
Applied thrust::no_init to both the output and temporary-storage allocations.
Made the output_it raw-pointer binding const following the CodeRabbit coding-guideline suggestion. This only prevents reassignment of the local pointer variable and does not make the output buffer const.

The target build, focused regression test, and targeted pre-commit checks pass locally.

bernhardmgruber · 2026-07-15T16:02:02Z

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


Important: Just use the launch wrapper adjacent_difference_subtract_left instead of handling the temporary storage allocation yourself.

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.

bernhardmgruber · 2026-07-15T17:13:28Z

+
+  const c2h::host_vector<type> expected{2, 3, 4, 5, 6};
+  REQUIRE(output == expected);


Important: I think we do not to sync when we use a custom stream:

Suggested change

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

Fix DeviceAdjacentDifference for device_buffer iterators [skip-docs]

1b058db

Noperi0r requested a review from a team as a code owner July 14, 2026 17:20

Noperi0r requested a review from pauleonix July 14, 2026 17:20

github-project-automation Bot added this to CCCL Jul 14, 2026

github-project-automation Bot moved this to Todo in CCCL Jul 14, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 14, 2026

bernhardmgruber reviewed Jul 14, 2026

View reviewed changes

Move DeviceAdjacentDifference device_buffer test [skip-docs]

facc0d5

coderabbitai Bot reviewed Jul 15, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_adjacent_difference_substract_left.cu Outdated

Make AdjacentDifference output pointer binding const [skip-docs]

7aed3a4

bernhardmgruber reviewed Jul 15, 2026

View reviewed changes

Use launch wrapper in device_buffer regression test [skip-docs]

8c1fa4f

bernhardmgruber reviewed Jul 15, 2026

View reviewed changes

	c2h::device_vector<type> output(input.size());
	c2h::device_vector<type> output(input.size(), thrust::no_init);

	c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes);
	c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init);


		const c2h::host_vector<type> expected{2, 3, 4, 5, 6};
		REQUIRE(output == expected);

Uh oh!

Conversation

Noperi0r commented Jul 14, 2026

Description

Checklist

Uh oh!

copy-pr-bot Bot commented Jul 14, 2026

Uh oh!

coderabbitai Bot commented Jul 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Assessment against linked issues

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Noperi0r commented Jul 15, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

coderabbitai Bot commented Jul 14, 2026 •

edited

Loading