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

Fix integer overflow in shim device_sum functions #13943

Merged
merged 6 commits into from
Aug 28, 2023

Conversation

brandon-b-miller
Copy link
Contributor

Closes #13873

@brandon-b-miller brandon-b-miller requested a review from a team as a code owner August 23, 2023 19:05
@github-actions github-actions bot added the Python Affects Python cuDF API. label Aug 23, 2023
@brandon-b-miller brandon-b-miller added bug Something isn't working numba Numba issue non-breaking Non-breaking change labels Aug 23, 2023
int64_t size,
T* sum)
template <typename T, std::enable_if_t<std::is_integral_v<T>, int> = 0>
__device__ int64_t device_sum(cooperative_groups::thread_block const& block,
Copy link
Contributor

@bdice bdice Aug 23, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rather than writing an entirely new SFINAE overload, could we have a single kernel and write something like this in a second template parameter:

typename AccumT = std::conditional<std::is_integral_v<T>, int64_t, T>

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yes! was hoping someone would comment with something like this :)

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

giphy

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should be updated! Looking much cleaner now, thanks.

@wence-
Copy link
Contributor

wence- commented Aug 24, 2023

I think this makes sense. Of course, there are many ways to still get a bad result, but pandas also gets a bad result in those cases. Overflowing int64 is the natural one, but pandas hilariously converts a sum of unsigned items to a signed type, so that's a more natural other one.

Copy link
Contributor

@wence- wence- left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Python test changes

python/cudf/udf_cpp/shim.cu Outdated Show resolved Hide resolved
@brandon-b-miller
Copy link
Contributor Author

/merge

@rapids-bot rapids-bot bot merged commit 724e42a into rapidsai:branch-23.10 Aug 28, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working non-breaking Non-breaking change numba Numba issue Python Affects Python cuDF API.
Projects
None yet
3 participants