From 1b058db11c14e847cdd4948660aef0cdda7e5eb0 Mon Sep 17 00:00:00 2001 From: Noperi0r Date: Wed, 15 Jul 2026 02:05:37 +0900 Subject: [PATCH 1/4] Fix DeviceAdjacentDifference for device_buffer iterators [skip-docs] --- cub/cub/agent/agent_adjacent_difference.cuh | 2 +- ...ch2_test_device_adjacent_difference_api.cu | 35 +++++++++++++++++++ 2 files changed, 36 insertions(+), 1 deletion(-) diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index e3ae0e7a80f..4f142c7cf39 100644 --- a/cub/cub/agent/agent_adjacent_difference.cuh +++ b/cub/cub/agent/agent_adjacent_difference.cuh @@ -105,7 +105,7 @@ struct AgentDifference OffsetT num_items) : temp_storage(temp_storage.Alias()) , input_it(input_it) - , load_it(LoadIt(input_it)) + , load_it(try_make_cache_modified_iterator(input_it)) , first_tile_previous(first_tile_previous) , result(result) , difference_op(difference_op) diff --git a/cub/test/catch2_test_device_adjacent_difference_api.cu b/cub/test/catch2_test_device_adjacent_difference_api.cu index 4eb9f5c3e17..5fb149bfcd2 100644 --- a/cub/test/catch2_test_device_adjacent_difference_api.cu +++ b/cub/test/catch2_test_device_adjacent_difference_api.cu @@ -5,6 +5,12 @@ #include +#include +#include +#include + +#include + #include // Guard: the legacy memory-size query call with all defaults (no explicit difference_op, @@ -51,4 +57,33 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRight legacy size-query is unambiguo REQUIRE(cudaSuccess == cub::DeviceAdjacentDifference::SubtractRight(nullptr, bytes, d_in, n)); } +C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy accepts cuda::device_buffer input", + "[adjacent_difference][device]") +{ + using type = std::int32_t; + + cuda::stream stream{cuda::devices[0]}; + auto input = cuda::make_device_buffer(stream, cuda::devices[0], {2, 5, 9, 14, 20}); + c2h::device_vector output(input.size()); + auto output_it = thrust::raw_pointer_cast(output.data()); + + 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 temporary_buffer(temporary_bytes); + 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())); + stream.sync(); + + const c2h::host_vector expected{2, 3, 4, 5, 6}; + REQUIRE(output == expected); +} + // todo(giannis): extract examples from the docs to literalinclude extracts here From facc0d52f83c205758ce8bca83d986bf5cfae5ba Mon Sep 17 00:00:00 2001 From: Noperi0r Date: Wed, 15 Jul 2026 23:26:29 +0900 Subject: [PATCH 2/4] Move DeviceAdjacentDifference device_buffer test [skip-docs] --- ...ch2_test_device_adjacent_difference_api.cu | 35 ------------------- ...vice_adjacent_difference_substract_left.cu | 30 ++++++++++++++++ 2 files changed, 30 insertions(+), 35 deletions(-) diff --git a/cub/test/catch2_test_device_adjacent_difference_api.cu b/cub/test/catch2_test_device_adjacent_difference_api.cu index 5fb149bfcd2..4eb9f5c3e17 100644 --- a/cub/test/catch2_test_device_adjacent_difference_api.cu +++ b/cub/test/catch2_test_device_adjacent_difference_api.cu @@ -5,12 +5,6 @@ #include -#include -#include -#include - -#include - #include // Guard: the legacy memory-size query call with all defaults (no explicit difference_op, @@ -57,33 +51,4 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRight legacy size-query is unambiguo REQUIRE(cudaSuccess == cub::DeviceAdjacentDifference::SubtractRight(nullptr, bytes, d_in, n)); } -C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy accepts cuda::device_buffer input", - "[adjacent_difference][device]") -{ - using type = std::int32_t; - - cuda::stream stream{cuda::devices[0]}; - auto input = cuda::make_device_buffer(stream, cuda::devices[0], {2, 5, 9, 14, 20}); - c2h::device_vector output(input.size()); - auto output_it = thrust::raw_pointer_cast(output.data()); - - 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 temporary_buffer(temporary_bytes); - 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())); - stream.sync(); - - const c2h::host_vector expected{2, 3, 4, 5, 6}; - REQUIRE(output == expected); -} - // todo(giannis): extract examples from the docs to literalinclude extracts here diff --git a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu index 4e36b2cb99c..05122dcab9d 100644 --- a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu +++ b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu @@ -5,6 +5,7 @@ #include +#include #include #include #include @@ -258,6 +259,35 @@ C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy works with user provided me test_subtract_left_copy(policy); } } + +C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy accepts cuda::device_buffer input", + "[adjacent_difference][device]") +{ + using type = std::int32_t; + + cuda::stream stream{cuda::devices[0]}; + auto input = cuda::make_device_buffer(stream, cuda::devices[0], {2, 5, 9, 14, 20}); + c2h::device_vector output(input.size(), thrust::no_init); + auto output_it = thrust::raw_pointer_cast(output.data()); + + 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 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())); + stream.sync(); + + const c2h::host_vector expected{2, 3, 4, 5, 6}; + REQUIRE(output == expected); +} #endif // TEST_LAUNCH == 0 C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy works with iterators", "[device][adjacent_difference]", types) From 7aed3a462082b7338a0bf89fe144e9bca4b5c615 Mon Sep 17 00:00:00 2001 From: Noperi0r Date: Thu, 16 Jul 2026 00:29:16 +0900 Subject: [PATCH 3/4] Make AdjacentDifference output pointer binding const [skip-docs] --- .../catch2_test_device_adjacent_difference_substract_left.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu index 05122dcab9d..62149ff85e6 100644 --- a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu +++ b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu @@ -268,7 +268,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy accepts cuda::device_buffer cuda::stream stream{cuda::devices[0]}; auto input = cuda::make_device_buffer(stream, cuda::devices[0], {2, 5, 9, 14, 20}); c2h::device_vector output(input.size(), thrust::no_init); - auto output_it = thrust::raw_pointer_cast(output.data()); + const auto output_it = thrust::raw_pointer_cast(output.data()); void* temporary_storage = nullptr; std::size_t temporary_bytes = 0; From 8c1fa4f3074fec94e152d49702e4f6a898851a6b Mon Sep 17 00:00:00 2001 From: Noperi0r Date: Thu, 16 Jul 2026 01:27:36 +0900 Subject: [PATCH 4/4] Use launch wrapper in device_buffer regression test [skip-docs] --- ...t_device_adjacent_difference_substract_left.cu | 15 +-------------- 1 file changed, 1 insertion(+), 14 deletions(-) diff --git a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu index 62149ff85e6..6087d7f6c1a 100644 --- a/cub/test/catch2_test_device_adjacent_difference_substract_left.cu +++ b/cub/test/catch2_test_device_adjacent_difference_substract_left.cu @@ -270,20 +270,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy accepts cuda::device_buffer c2h::device_vector output(input.size(), thrust::no_init); const auto output_it = thrust::raw_pointer_cast(output.data()); - 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 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())); - stream.sync(); + adjacent_difference_subtract_left_copy(input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get()); const c2h::host_vector expected{2, 3, 4, 5, 6}; REQUIRE(output == expected);