Dropped duplicated math function from Thrust by viralbhadeshiya · Pull Request #6188 · NVIDIA/cccl

viralbhadeshiya · 2025-10-10T22:50:41Z

Description

closes #6104

Removed duplicated math function from Thrust
Moved all duplicated math function call from Thrust to cuda/std or cuda

Checklist

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

copy-pr-bot · 2025-10-10T22:50:44Z

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.

bernhardmgruber

Please remove

#include <thrust/detail/integer_math.h>

from files where we no longer need it.

miscco

Thanks a lot for working on this issue. It is already looking great.

I found one incorrect replacement and some potential conversion issues

viralbhadeshiya · 2025-10-11T18:14:51Z

@bernhardmgruber @miscco all changes that you guys suggested is done, lmk if there are more Changes. :)

miscco · 2025-10-13T05:36:34Z

@viralbhadeshiya This looks great 🎉

Thanks a lot for working on this.

I added a few minor changes:

dropped unused is_odd
improved is_negative
dropped two more includes

miscco · 2025-10-13T05:36:48Z

/ok to test 07b811c

miscco · 2025-10-16T16:13:03Z

/ok to test 34fc385

miscco · 2025-10-17T07:31:38Z

@fbusato There is an issue with thrust::pool_options

I suspect there is a differing behavior for is_power_of_two could you please inevestigate why it differs?

fbusato · 2025-10-17T16:45:49Z

::cuda::is_power_of_two correctly evaluates 0 as non-power of two, while thrust::detail::is_power_of_2() misses the check x != 0.

This causes a different behavior in !::cuda::is_power_of_two(alignment), which is initialized to 0 by default.
https://github.com/NVIDIA/cccl/pull/6188/files#diff-d85bf89d70ec2f2a069c4c0ffeccd2722fe2a108a012566503392885900f40a1R115

fbusato · 2025-10-17T16:48:06Z

@viralbhadeshiya my suggestion is just to add the condition alignment != 0 && !::cuda::is_power_of_two(alignment)

viralbhadeshiya · 2025-10-17T16:53:47Z

::cuda::is_power_of_two correctly evaluates 0 as non-power of two, while thrust::detail::is_power_of_2() misses the check x != 0.

This causes a different behavior in !::cuda::is_power_of_two(alignment), which is initialized to 0 by default. https://github.com/NVIDIA/cccl/pull/6188/files#diff-d85bf89d70ec2f2a069c4c0ffeccd2722fe2a108a012566503392885900f40a1R115

@fbusato would this similar changes also required in !::cuda::is_power_of_two(smallest_block_size) & !::cuda::is_power_of_two(largest_block_size)?

fbusato · 2025-10-17T17:30:27Z

yes, but I'm not sure that this edge case is evaluated in the test case. You can also run the test locally to make sure it works

viralbhadeshiya · 2025-10-17T17:41:21Z

yes, but I'm not sure that this edge case is evaluated in the test case. You can also run the test locally to make sure it works

Okay made changes, unfortunately my remote server which had NVIDIA device is facing some issue so could not run tests. but changes are done.

fbusato · 2025-10-17T17:47:18Z

/ok to test 0ab0b72

github-actions · 2025-10-17T22:19:15Z

🥳 CI Workflow Results

🟩 Finished in 4h 29m: Pass: 100%/118 | Total: 3d 10h | Max: 4h 28m | Hits: 90%/166971

See results here.

…tion from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Partial: upstream deletes thrust/detail/integer_math.h and switches execute_with_allocator.h / mr/{pool,disjoint_pool,pool_options}.h to ::cuda::ceil_div / ceil_ilog2 / is_power_of_two. Kept integer_math.h and the thrust::detail helpers in generic headers for compatibility (those ::cuda:: cmath utilities are unavailable without libhipcxx; behaviour is identical). system/cuda: dropped the now-redundant integer_math.h include from scan.h/sort.h (is_negative still reaches scan.h via dispatch.h); cuda dispatch.h keeps including integer_math.h instead of defining its own is_negative, which would collide with integer_math.h's overloads. HIP: no change (hip dispatch.h still uses thrust::detail::is_negative from integer_math.h). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…function from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Completes the earlier partial port. libhipcxx is now mandatory (libcxx.h requires >= 3.4), so the compatibility reason for keeping thrust::detail integer helpers no longer applies: - delete thrust/detail/integer_math.h. - execute_with_allocator.h: divide_ri -> _THRUST_LIBCXX::ceil_div. - mr/{pool,disjoint_pool}.h: log2_ri -> _THRUST_LIBCXX::ceil_ilog2, is_power_of_2 -> _THRUST_LIBCXX::is_power_of_two. - mr/pool_options.h: validate() uses is_power_of_two with the upstream explicit != 0 checks. - system/cuda/detail/dispatch.h: take upstream (defines its own thrust::detail::is_negative). - system/hip/detail/dispatch.h: same is_negative, spelled with _THRUST_STD/THRUST_HOST_DEVICE; hip scan.h/sort.h drop the integer_math.h include. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…tion from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Partial: upstream deletes thrust/detail/integer_math.h and switches execute_with_allocator.h / mr/{pool,disjoint_pool,pool_options}.h to ::cuda::ceil_div / ceil_ilog2 / is_power_of_two. Kept integer_math.h and the thrust::detail helpers in generic headers for compatibility (those ::cuda:: cmath utilities are unavailable without libhipcxx; behaviour is identical). system/cuda: dropped the now-redundant integer_math.h include from scan.h/sort.h (is_negative still reaches scan.h via dispatch.h); cuda dispatch.h keeps including integer_math.h instead of defining its own is_negative, which would collide with integer_math.h's overloads. HIP: no change (hip dispatch.h still uses thrust::detail::is_negative from integer_math.h). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…function from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Completes the earlier partial port. libhipcxx is now mandatory (libcxx.h requires >= 3.4), so the compatibility reason for keeping thrust::detail integer helpers no longer applies: - delete thrust/detail/integer_math.h. - execute_with_allocator.h: divide_ri -> _THRUST_LIBCXX::ceil_div. - mr/{pool,disjoint_pool}.h: log2_ri -> _THRUST_LIBCXX::ceil_ilog2, is_power_of_2 -> _THRUST_LIBCXX::is_power_of_two. - mr/pool_options.h: validate() uses is_power_of_two with the upstream explicit != 0 checks. - system/cuda/detail/dispatch.h: take upstream (defines its own thrust::detail::is_negative). - system/hip/detail/dispatch.h: same is_negative, spelled with _THRUST_STD/THRUST_HOST_DEVICE; hip scan.h/sort.h drop the integer_math.h include. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…tion from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Partial: upstream deletes thrust/detail/integer_math.h and switches execute_with_allocator.h / mr/{pool,disjoint_pool,pool_options}.h to ::cuda::ceil_div / ceil_ilog2 / is_power_of_two. Kept integer_math.h and the thrust::detail helpers in generic headers for compatibility (those ::cuda:: cmath utilities are unavailable without libhipcxx; behaviour is identical). system/cuda: dropped the now-redundant integer_math.h include from scan.h/sort.h (is_negative still reaches scan.h via dispatch.h); cuda dispatch.h keeps including integer_math.h instead of defining its own is_negative, which would collide with integer_math.h's overloads. HIP: no change (hip dispatch.h still uses thrust::detail::is_negative from integer_math.h). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…function from Thrust (NVIDIA/cccl#6188) Upstream: NVIDIA/cccl@5303f65 Completes the earlier partial port. libhipcxx is now mandatory (libcxx.h requires >= 3.4), so the compatibility reason for keeping thrust::detail integer helpers no longer applies: - delete thrust/detail/integer_math.h. - execute_with_allocator.h: divide_ri -> _THRUST_LIBCXX::ceil_div. - mr/{pool,disjoint_pool}.h: log2_ri -> _THRUST_LIBCXX::ceil_ilog2, is_power_of_2 -> _THRUST_LIBCXX::is_power_of_two. - mr/pool_options.h: validate() uses is_power_of_two with the upstream explicit != 0 checks. - system/cuda/detail/dispatch.h: take upstream (defines its own thrust::detail::is_negative). - system/hip/detail/dispatch.h: same is_negative, spelled with _THRUST_STD/THRUST_HOST_DEVICE; hip scan.h/sort.h drop the integer_math.h include. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

viralbhadeshiya added 8 commits October 10, 2025 10:33

fixed duplication for log2_ri

1a37095

changed cub for log2_ri deplication

87b98f1

fix ceil_ilog2 import

96a8c4d

fix sytax error

8b1eadd

Drop all duplicated math function from Thrust

8eec94e

fix syntax issue for divide_rz

70938ba

remove duplicate declaration

1a883b7

fix pair error on thrust

0b4597d

viralbhadeshiya requested review from a team as code owners October 10, 2025 22:50

viralbhadeshiya requested a review from NaderAlAwar October 10, 2025 22:50

github-project-automation Bot added this to CCCL Oct 10, 2025

viralbhadeshiya requested a review from elstehle October 10, 2025 22:50

github-project-automation Bot moved this to Todo in CCCL Oct 10, 2025

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 10, 2025

viralbhadeshiya added 2 commits October 10, 2025 15:56

Merge remote-tracking branch 'remote/main' into FEA/6104

237f908

changes for pre-commit hook

5a43614

bernhardmgruber approved these changes Oct 11, 2025

View reviewed changes

Comment thread thrust/thrust/detail/integer_math.h Resolved

PR review changes

01a27de

miscco requested changes Oct 11, 2025

View reviewed changes

Comment thread cub/cub/device/dispatch/dispatch_merge_sort.cuh Resolved

Comment thread cub/cub/device/dispatch/dispatch_merge_sort.cuh Resolved

Comment thread thrust/thrust/detail/execute_with_allocator.h Resolved

Comment thread thrust/thrust/mr/disjoint_pool.h Resolved

github-project-automation Bot moved this from In Review to In Progress in CCCL Oct 11, 2025

fixes for PR review suggestion 2

f0d7b38

viralbhadeshiya requested a review from miscco October 11, 2025 18:14

miscco added 3 commits October 13, 2025 07:12

Drop unused is_odd

1c3060f

Improve is_negative implementation and drop unused headers

701af01

Drop unused includes

07b811c

This comment has been minimized.

Sign in to view

viralbhadeshiya and others added 2 commits October 15, 2025 14:35

Merge branch 'main' into FEA/6104

8a80170

Merge branch 'main' into FEA/6104

34fc385

This comment has been minimized.

Sign in to view

miscco assigned fbusato Oct 17, 2025

fbusato assigned viralbhadeshiya and unassigned fbusato Oct 17, 2025

viralbhadeshiya added 2 commits October 17, 2025 09:55

bug: fix different behavior for ::cuda::is_power_of_two

d147cd1

Merge branch 'main' into FEA/6104

e117d67

viralbhadeshiya added 2 commits October 17, 2025 10:38

fix for ::cuda::is_power_of_two

16d4048

Merge branch 'FEA/6104' of github.com:viralbhadeshiya/cccl into FEA/6104

0ab0b72

bernhardmgruber merged commit 5303f65 into NVIDIA:main Oct 20, 2025
128 checks passed

github-project-automation Bot moved this from In Review to Done in CCCL Oct 20, 2025

Conversation

viralbhadeshiya commented Oct 10, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

copy-pr-bot Bot commented Oct 10, 2025

Uh oh!

bernhardmgruber left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

miscco left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

viralbhadeshiya commented Oct 11, 2025

Uh oh!

miscco commented Oct 13, 2025

Uh oh!

miscco commented Oct 13, 2025

Uh oh!

This comment has been minimized.

miscco commented Oct 16, 2025

Uh oh!

This comment has been minimized.

miscco commented Oct 17, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

fbusato commented Oct 17, 2025

Uh oh!

fbusato commented Oct 17, 2025

Uh oh!

viralbhadeshiya commented Oct 17, 2025

Uh oh!

fbusato commented Oct 17, 2025

Uh oh!

viralbhadeshiya commented Oct 17, 2025

Uh oh!

fbusato commented Oct 17, 2025

Uh oh!

github-actions Bot commented Oct 17, 2025

🥳 CI Workflow Results

🟩 Finished in 4h 29m: Pass: 100%/118 | Total: 3d 10h | Max: 4h 28m | Hits: 90%/166971

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

viralbhadeshiya commented Oct 10, 2025 •

edited

Loading

miscco commented Oct 17, 2025 •

edited

Loading