ginkgo-project / ginkgo

Numerical linear algebra software package
https://ginkgo-project.github.io/
BSD 3-Clause "New" or "Revised" License
414 stars 88 forks source link

Remove attributes from math functions #1695

Closed MarcelKoch closed 1 month ago

MarcelKoch commented 1 month ago

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.

upsj commented 1 month ago

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 commented 1 month ago

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 commented 1 month ago

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 commented 1 month ago

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 commented 1 month ago

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.

sonarcloud[bot] commented 1 month ago

Quality Gate Passed Quality Gate passed

Issues
16 New issues
0 Accepted issues

Measures
0 Security Hotspots
80.0% Coverage on New Code
0.0% Duplication on New Code

See analysis details on SonarCloud

codecov[bot] commented 1 month ago

Codecov Report

All modified and coverable lines are covered by tests :white_check_mark:

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

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

:umbrella: View full report in Codecov by Sentry.
:loudspeaker: Have feedback on the report? Share it here.