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

Remove attributes from math functions #1695

Merged
merged 2 commits into from
Oct 18, 2024

Conversation

MarcelKoch
Copy link
Member

This PR removes GKO_ATTRIBUTES from constexpr functions in math.hpp. Since we build cuda with --expt-relaxed-constexpr the constexpr is already enough to allow those functions on the device.

Fixes #1693.

@MarcelKoch MarcelKoch added the 1:ST:ready-for-review This PR is ready for review label Oct 15, 2024
@MarcelKoch MarcelKoch added this to the Ginkgo 1.9.0 milestone Oct 15, 2024
@MarcelKoch MarcelKoch requested a review from a team October 15, 2024 07:41
@MarcelKoch MarcelKoch self-assigned this Oct 15, 2024
@ginkgo-bot ginkgo-bot added the mod:core This is related to the core module. label Oct 15, 2024
Copy link
Collaborator

@greole greole left a comment

Choose a reason for hiding this comment

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

LGTM!

Copy link
Member

@upsj upsj left a comment

Choose a reason for hiding this comment

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

LGTM!

Copy link
Member

@yhmtsai yhmtsai left a comment

Choose a reason for hiding this comment

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

It will work with __half, but I think it is somehow by accident or not strict rule from cuda compiler. they might inline it before checking whether it can be constexpr or not.
__half does not have any constexpr constructor.
zero<__half> and one<__half> will not be done in compiled time, and then these functions are not __device__ kernel. Thus, the device kernel should not be able to use these function, right?

@upsj
Copy link
Member

upsj commented Oct 17, 2024

Inlining happens much later than any checks for function attributes, so this should not be an issue at all. If zero and one are constexpr, they can be called from device code. If the function they themselves are calling is not callable from device code, the compiler should complain.

@yhmtsai
Copy link
Member

yhmtsai commented Oct 17, 2024

but even if the function is __device__ __host__ but not constexpr, the zero/one function should behave like normal function in __host__ because of no attribute.
no matter if function is device callable or not, device kernel should not be able to call zero/one in this case

@MarcelKoch
Copy link
Member Author

The one and zero functions are already marked constexpr and we can't remove it, since that would be an interface break. So the only question is if __device__ __host__ is necessary, which I argue it is not due to cuda's compiler option. AFAIK, it doesn't matter if the constexpr function is actually evaluated at compile time.
If these functions can't be constexpr for gko::half, or __half, then this needs to be addressed anyway, but not in this PR.

@yhmtsai
Copy link
Member

yhmtsai commented Oct 17, 2024

I know constexpr does not always need to be evaluated in compile time.
when constexpr is failed, I thought __device__ __host__ will be necessary.
So the original one will work no matter if __half is constexpr construable or not.
it sounds like just stronger inline in some sense.

@upsj
Copy link
Member

upsj commented Oct 17, 2024

To make the terminology a bit clearer: A function is marked constexpr if it has the attribute next to its name - that's what nvcc cares about. A function is constexpr if it is marked constexpr and can be evaluated at compile-time. For us, only the former matters right now.

Copy link
Member

@yhmtsai yhmtsai left a comment

Choose a reason for hiding this comment

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

If the constexpr in the first place does not need to be checked, I am fine with the changes

@MarcelKoch MarcelKoch added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels Oct 17, 2024
@MarcelKoch MarcelKoch force-pushed the remove-attributes-from-math-functions branch from 20f7f32 to 98f29f6 Compare October 17, 2024 18:14
Copy link

sonarcloud bot commented Oct 18, 2024

Copy link

codecov bot commented Oct 18, 2024

Codecov Report

All modified and coverable lines are covered by tests ✅

Project coverage is 89.96%. Comparing base (532566d) to head (98f29f6).
Report is 3 commits behind head on develop.

Additional details and impacted files
@@           Coverage Diff            @@
##           develop    #1695   +/-   ##
========================================
  Coverage    89.96%   89.96%           
========================================
  Files          763      763           
  Lines        62877    62878    +1     
========================================
+ Hits         56570    56571    +1     
  Misses        6307     6307           

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

@MarcelKoch MarcelKoch merged commit db80c7d into develop Oct 18, 2024
11 of 14 checks passed
@MarcelKoch MarcelKoch deleted the remove-attributes-from-math-functions branch October 18, 2024 08:38
@yhmtsai yhmtsai mentioned this pull request Oct 18, 2024
12 tasks
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
1:ST:ready-to-merge This PR is ready to merge. mod:core This is related to the core module.
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Can't build ginkgo 1.8.0 with HIP
5 participants