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

bernhardmgruber · 2026-03-25T11:33:09Z


  thrust::device_vector<float> expected{1.0f, 1.0f, 2.0f, 4.0f};
  // example-end exclusive-scan-env-stream
+  stream.sync();


Q: Why do we need a sync here? I don't think we need it.

Applies more often in this file.

Initializing expected implicitly syncs, but it is good to be explicit about it.

i added for me to remove this discrepancy here. afaiu we should be using stream.sync() but now show it to the users as it should be left on their judgement for when to use it

#8175 (comment)

Initializing expected implicitly syncs

But the code here does not rely on this, so it's fine if it wouldn't.

but it is good to be explicit about it.

I disagree. No sync is needed here, so seeing one makes readers like me nervous.

I discussed this with @gonidelis and I now understand why we need the sync. We wondered why the code still works without the sync. Where does

Initializing expected implicitly syncs

Is it the memory allocation? I just want to understand why it works without a sync.

We should definitely add the syncs now!

bernhardmgruber · 2026-04-02T21:26:47Z


  thrust::device_vector<float> expected{1.0f, 1.0f, 2.0f, 4.0f};
  // example-end exclusive-scan-env-stream
+  stream.sync();


Initializing expected implicitly syncs

But the code here does not rely on this, so it's fine if it wouldn't.

but it is good to be explicit about it.

I disagree. No sync is needed here, so seeing one makes readers like me nervous.

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

…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

…hrust::equal

bernhardmgruber · 2026-04-08T15:14:42Z


  thrust::device_vector<float> expected{1.0f, 1.0f, 2.0f, 4.0f};
  // example-end exclusive-scan-env-stream
+  stream.sync();


I discussed this with @gonidelis and I now understand why we need the sync. We wondered why the code still works without the sync. Where does

Initializing expected implicitly syncs

Is it the memory allocation? I just want to understand why it works without a sync.

We should definitely add the syncs now!

bernhardmgruber · 2026-04-08T17:23:12Z

There are errors in the docs build:

/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:348: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:654: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:1123: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:1393: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:1782: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:1867: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:2603: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:2732: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:2860: error: explicit link request to 'cuda::std::execution::env' could not be resolved
/home/runner/_work/cccl/cccl/cub/cub/device/device_scan.cuh:2981: error: explicit link request to 'cuda::std::execution::env' could not be resolved

Please address those.

github-actions · 2026-04-08T19:18:13Z

🥳 CI Workflow Results

🟩 Finished in 1h 26m: Pass: 100%/269 | Total: 3d 15h | Max: 1h 07m | Hits: 98%/177053

See results here.

) * Add ebv overloads for InclusiveSum and future ExclusiveScan * Add dynamic smem methods to stream_registry_factory_t and tidy-up tests - 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) * Improve DeviceScan env test coverage and clean up duplicate doc comments - 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)

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

Comment thread cub/cub/device/device_scan.cuh Outdated

gonidelis commented Mar 19, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh

gonidelis commented Mar 19, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env.cu

gonidelis commented Mar 19, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env_api.cu

This comment has been minimized.

Sign in to view

bernhardmgruber reviewed Mar 23, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh

bernhardmgruber reviewed Mar 23, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh

Comment thread cub/test/catch2_test_device_scan_env_api.cu Outdated

Comment thread cub/test/catch2_test_device_scan_env_api.cu Outdated

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

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

Comment thread cub/test/catch2_test_device_scan_env.cu Outdated

pauleonix reviewed Mar 30, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env.cu Outdated

pauleonix reviewed Mar 30, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env.cu Outdated

pauleonix reviewed Mar 30, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env_api.cu Outdated

pauleonix reviewed Mar 30, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_scan_env_api.cu Outdated

gonidelis mentioned this pull request Apr 2, 2026

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

Merged

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

This comment has been minimized.

Sign in to view

bernhardmgruber reviewed Apr 4, 2026

View reviewed changes

pauleonix reviewed Apr 7, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh

pauleonix reviewed Apr 7, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh Outdated

pauleonix reviewed Apr 7, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh

pauleonix reviewed Apr 7, 2026

View reviewed changes

Comment thread cub/cub/device/device_scan.cuh Outdated

gonidelis added 12 commits April 8, 2026 07:29

Add ebv overloads for InclusiveSum and future ExclusiveScan

8527737

reviews

1cb127f

fix/add versionadds

05361f4

docs fix

f59c9aa

Doc fixes

d21d8cf

add stream.sync()

3076cfc

Add corner case num_items values and substitute requires loops with t…

685a6c4

…hrust::equal

remove unused non_guaranteed api test

ba3dba6

Add FutureValue float determinism fail tests for ExclusiveScan

b75fb7b

gonidelis force-pushed the init_sum_env branch from 8a7dc45 to b75fb7b Compare April 8, 2026 14:45

gonidelis requested review from bernhardmgruber and pauleonix April 8, 2026 14:46

bernhardmgruber reviewed Apr 8, 2026

View reviewed changes

Substitute thrust vectors with c2h vectors

8225594

This comment has been minimized.

Sign in to view

bernhardmgruber approved these changes Apr 8, 2026

View reviewed changes

fix docs

1569cbe

gonidelis merged commit 8d77295 into NVIDIA:main Apr 8, 2026
285 of 288 checks passed

github-project-automation Bot moved this from In Review to Done in CCCL Apr 8, 2026

Conversation

gonidelis commented Mar 17, 2026

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber Mar 25, 2026

Choose a reason for hiding this comment

Uh oh!

pauleonix Mar 30, 2026

Choose a reason for hiding this comment

Uh oh!

gonidelis Apr 2, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Apr 2, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Apr 8, 2026

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

bernhardmgruber Apr 2, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

bernhardmgruber Apr 8, 2026

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

bernhardmgruber commented Apr 8, 2026

Uh oh!

github-actions Bot commented Apr 8, 2026

🥳 CI Workflow Results

🟩 Finished in 1h 26m: Pass: 100%/269 | Total: 3d 15h | Max: 1h 07m | Hits: 98%/177053

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants