Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Restrict CUB device histograms to NUM_ACTIVE_CHANNELS <= NUM_CHANNELS #1792

Closed
bernhardmgruber opened this issue May 31, 2024 · 1 comment · Fixed by #1796
Closed

Restrict CUB device histograms to NUM_ACTIVE_CHANNELS <= NUM_CHANNELS #1792

bernhardmgruber opened this issue May 31, 2024 · 1 comment · Fixed by #1796
Assignees
Labels
bug: functional cub For all items related to CUB

Comments

@bernhardmgruber
Copy link
Contributor

bernhardmgruber commented May 31, 2024

The cub::DeviceHistogram::MultiHistogram*** APIs can be compiled and run with more active channels than channels in the input. E.g. having 4 active channels on a 3 channel sample. This leads to interesting results and unit test failures.

This rather seems to be by accident than a feature, so I propose disabling it via a static_assert.

@bernhardmgruber bernhardmgruber added the feature request New feature or request. label May 31, 2024
@bernhardmgruber bernhardmgruber added bug: functional and removed feature request New feature or request. labels May 31, 2024
@bernhardmgruber bernhardmgruber changed the title [FEA]: Restrict CUB device histograms to NUM_ACTIVE_CHANNELS <= NUM_CHANNELS Restrict CUB device histograms to NUM_ACTIVE_CHANNELS <= NUM_CHANNELS May 31, 2024
@gevtushenko
Copy link
Collaborator

@bernhardmgruber I believe this is an oversight, thank you for spotting this!

The documentation states:

Of the NUM_CHANNELS specified, the function will only compute histograms for the first NUM_ACTIVE_CHANNELS

Which implies to me that the number of active channels is smaller than the total number of channels. On the implementation side, active channels exceeding num channels leads to OOB access, so catching this use case as a compile-time error is preferrable.

Apart from adding a static assert, maybe we could clarify this aspect in documentation as well?

@bernhardmgruber bernhardmgruber self-assigned this May 31, 2024
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue May 31, 2024
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue May 31, 2024
@bernhardmgruber bernhardmgruber added the cub For all items related to CUB label Jun 1, 2024
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jun 4, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug: functional cub For all items related to CUB
Projects
Archived in project
Development

Successfully merging a pull request may close this issue.

2 participants