Add env overloads for InclusiveSum and future ExclusiveScan by gonidelis · Pull Request #8059 · NVIDIA/cccl

gonidelis · 2026-03-17T05:40:48Z

Adds miscellaneous env overloads for InclusiveSum and ExclusiveScan with FutureValue

gonidelis · 2026-03-19T17:16:33Z

cub/cub/device/device_scan.cuh

+  //! - When ``d_in`` and ``d_out`` are equal, the scan is performed in-place. The
+  //!   range ``[d_in, d_in + num_items)`` and ``[d_out, d_out + num_items)``
+  //!   shall not overlap in any other way.
+  //!


add env overload documentation info

cub/cub/device/device_scan.cuh

cub/test/catch2_test_device_scan_env.cu

cub/test/catch2_test_device_scan_env_api.cu

cub/cub/device/device_scan.cuh

bernhardmgruber · 2026-03-23T08:29:33Z

cub/test/catch2_test_device_scan_env_api.cu

+C2H_TEST("cub::DeviceScan::ExclusiveScan accepts not_guaranteed determinism requirements", "[scan][env]")
+{
+  auto op     = cuda::std::plus{};
+  auto input  = thrust::device_vector<int>{0, 1, 2, 3};
+  auto output = thrust::device_vector<int>(4);
+  auto init   = 0;
+
+  auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
+
+  auto error = cub::DeviceScan::ExclusiveScan(input.begin(), output.begin(), op, init, input.size(), env);
+  if (error != cudaSuccess)
+  {
+    std::cerr << "cub::DeviceScan::ExclusiveScan failed with status: " << error << std::endl;
+  }
+
+  thrust::device_vector<int> expected{0, 0, 1, 3};
+
+  REQUIRE(error == cudaSuccess);
+  REQUIRE(output == expected);
+}


Important: This example seems unused by any documentation. Please add it.

@gonidelis you resolved this comment but there are still no // example-begin etc. comments in this example. Therefore, it's still not used anywhere. Please address this.

there might be api test cases where we don't insert in the docs in order to avoid bloating. i added the guards but did not insert it in the rst docs with a literalinclude as a non_guaranteed example already exists

we can also remove that

cub/test/catch2_test_device_scan_env_api.cu

- dispatch_scan's warpspeed path calls max_dynamic_smem_size_for and set_max_dynamic_smem_size_for on the launcher factory - Previously untriggered because env tests used constant_iterator (non-contiguous), which skips the warpspeed path - Now needed for InclusiveSum env tests with device_vector (contiguous iterators)

…ents - Add default environment test for ExclusiveScan with FutureValue - Add not_guaranteed determinism test for ExclusiveScan - Remove duplicate in-place precondition from env overload descriptions (already in Preconditions section) - Remove @devicestorage from env overloads (no temp storage parameter)

cub/cub/device/device_scan.cuh

bernhardmgruber · 2026-03-25T11:32:21Z

cub/test/catch2_test_device_scan_env_api.cu

+C2H_TEST("cub::DeviceScan::ExclusiveScan accepts not_guaranteed determinism requirements", "[scan][env]")
+{
+  auto op     = cuda::std::plus{};
+  auto input  = thrust::device_vector<int>{0, 1, 2, 3};
+  auto output = thrust::device_vector<int>(4);
+  auto init   = 0;
+
+  auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed);
+
+  auto error = cub::DeviceScan::ExclusiveScan(input.begin(), output.begin(), op, init, input.size(), env);
+  if (error != cudaSuccess)
+  {
+    std::cerr << "cub::DeviceScan::ExclusiveScan failed with status: " << error << std::endl;
+  }
+
+  thrust::device_vector<int> expected{0, 0, 1, 3};
+
+  REQUIRE(error == cudaSuccess);
+  REQUIRE(output == expected);
+}


@gonidelis you resolved this comment but there are still no // example-begin etc. comments in this example. Therefore, it's still not used anywhere. Please address this.

cub/test/catch2_test_device_scan_env_api.cu

cub/test/catch2_test_device_scan_env.cu

pauleonix · 2026-03-30T17:29:42Z

cub/test/catch2_test_device_scan_env.cu

+{
+  using num_items_t = int;
+
+  num_items_t num_items = 10;


Use (small) random value to include corner-cases.

this would require to make the logic for the expected vectors more complex, do you think this change is necessary? are you referring to zero sized inputs type of corner cases?

Do use GENERATE() to generate size of input but elements are still all 1 using constant_iter

cub/test/catch2_test_device_scan_env.cu

cub/test/catch2_test_device_scan_env_api.cu

…ents - Use non-identity init values - Replace per-element REQUIRE(d_out[i]) with bulk thrust::equal or device_vector comparison where possible - Add example-begin/end markers for exclusive-scan-env-not-guaranteed - Fix mismatched init/expected in inclusive-scan-init env test

github-actions · 2026-04-02T21:06:48Z

🥳 CI Workflow Results

🟩 Finished in 1h 56m: Pass: 100%/249 | Total: 5d 03h | Max: 1h 38m | Hits: 93%/160113

See results here.

gonidelis · 2026-04-03T15:50:51Z

cub/test/catch2_test_device_scan_env.cu

  REQUIRE(d_block_size[0] == target_block_size);
 }

+TEST_CASE("Device scan exclusive scan with FutureValue works with default environment", "[scan][device]")


Add failed tests for float run_to_run here:
/home/ggonidelis/cccl/cub/test/test_device_exclusive_scan_determinism_fail.cu

bernhardmgruber · 2026-04-02T21:24:33Z

cub/test/catch2_test_device_scan_env_api.cu

+
+  thrust::device_vector<float> expected{1.0f, 3.0f, 6.0f, 10.0f};
+  // example-end inclusive-sum-env-stream
+  stream.sync();


Important: we should not need any sync here. Also not in any other tests in this file. Why is it needed?

cub/test/catch2_test_device_scan_env_api.cu

gonidelis requested a review from a team as a code owner March 17, 2026 05:40

gonidelis requested a review from pauleonix March 17, 2026 05:40

github-project-automation bot added this to CCCL Mar 17, 2026

github-project-automation bot moved this to Todo in CCCL Mar 17, 2026

gonidelis changed the title ~~Add ebv overloads for InclusiveSum and future ExclusiveScan~~ Add env overloads for InclusiveSum and future ExclusiveScan Mar 17, 2026

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Mar 17, 2026

gonidelis mentioned this pull request Mar 17, 2026

Add env-based API for cub::DeviceScan #7548

Closed

9 tasks

This comment has been minimized.

Sign in to view

gonidelis requested a review from srinivasyadav18 March 19, 2026 02:33

gonidelis commented Mar 19, 2026

View reviewed changes

cub/cub/device/device_scan.cuh Show resolved Hide resolved

gonidelis commented Mar 19, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env.cu Show resolved Hide resolved

gonidelis commented Mar 19, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env_api.cu Show resolved Hide resolved

This comment has been minimized.

Sign in to view

bernhardmgruber reviewed Mar 23, 2026

View reviewed changes

cub/cub/device/device_scan.cuh Show resolved Hide resolved

bernhardmgruber reviewed Mar 23, 2026

View reviewed changes

gonidelis enabled auto-merge (squash) March 25, 2026 10:34

gonidelis added 4 commits March 25, 2026 03:38

Add ebv overloads for InclusiveSum and future ExclusiveScan

4f2e167

reviews

1328c7c

gonidelis force-pushed the init_sum_env branch from 06a483a to 1328c7c Compare March 25, 2026 10:41

gonidelis requested a review from bernhardmgruber March 25, 2026 10:41

bernhardmgruber reviewed Mar 25, 2026

View reviewed changes

This comment has been minimized.

Sign in to view

pauleonix reviewed Mar 30, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env.cu Outdated Show resolved Hide resolved

pauleonix reviewed Mar 30, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env.cu Outdated Show resolved Hide resolved

pauleonix reviewed Mar 30, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env_api.cu Outdated Show resolved Hide resolved

pauleonix reviewed Mar 30, 2026

View reviewed changes

cub/test/catch2_test_device_scan_env_api.cu Outdated Show resolved Hide resolved

gonidelis added 2 commits March 30, 2026 17:22

fix/add versionadds

5c1267c

docs fix

43c02b5

gonidelis mentioned this pull request Apr 2, 2026

Improve DeviceScan env test coverage and clean up stale code #8280

Open

gonidelis requested review from bernhardmgruber and pauleonix April 2, 2026 19:10

gonidelis commented Apr 3, 2026

View reviewed changes

bernhardmgruber reviewed Apr 4, 2026

View reviewed changes

Conversation

gonidelis commented Mar 17, 2026

Uh oh!

This comment has been minimized.

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

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!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gonidelis Apr 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions bot commented Apr 2, 2026

🥳 CI Workflow Results

🟩 Finished in 1h 56m: Pass: 100%/249 | Total: 5d 03h | Max: 1h 38m | Hits: 93%/160113

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

gonidelis Apr 2, 2026 •

edited

Loading