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

Use LvArray::math functions #2927

Merged
merged 5 commits into from
Jan 22, 2024
Merged

Use LvArray::math functions #2927

merged 5 commits into from
Jan 22, 2024

Conversation

untereiner
Copy link
Contributor

closes #2926

@untereiner untereiner self-assigned this Jan 8, 2024
@untereiner untereiner added ci: run CUDA builds Allows to triggers (costly) CUDA jobs ci: run integrated tests Allows to run the integrated tests in GEOS CI flag: no rebaseline Does not require rebaseline flag: ready for review labels Jan 8, 2024
real32 const localIncrementx = density[e] * (velocityVp[e] * abs( nx ) + velocityVs[e] * sqrt( pow( ny, 2 ) + pow( nz, 2 ) ) ) * aux;
real32 const localIncrementy = density[e] * (velocityVp[e] * abs( ny ) + velocityVs[e] * sqrt( pow( nx, 2 ) + pow( nz, 2 ) ) ) * aux;
real32 const localIncrementz = density[e] * (velocityVp[e] * abs( nz ) + velocityVs[e] * sqrt( pow( nx, 2 ) + pow( ny, 2 ) ) ) * aux;
real32 const localIncrementx = density[e] * (velocityVp[e] * std::abs( nx ) + velocityVs[e] * sqrt( pow( ny, 2 ) + pow( nz, 2 ) ) ) * aux;
Copy link
Contributor

Choose a reason for hiding this comment

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

Don't you think you should be using LvArray::math::abs instead

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Why? Someone might change one day the define real32 ?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

What about sqrt and pow? I don’t see any using namespace Lvarray::math. That would be koenig’s lookup aware by the way

Copy link
Contributor

Choose a reason for hiding this comment

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

Why? Someone might change one day the define real32 ?

Because the LvArray::math functions are supposed to be GPU proof. Which makes sense in a __device__ function.

Copy link
Contributor

Choose a reason for hiding this comment

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

std::abs will not compile on device. Sadly, LvArray::math::abs does not provide correct overload for floating point values either (it dispatches to ::abs(int) on device which causes truncation).

You have 2 options:

  • use fabsf (no namespace)
  • or implement correct overloads in LvArray/src/math.hpp (dispatch to ::fabsf for float and ::fabs for double) and use them.

sqrt, pow, fabs, etc. are functions in the global namespace inherited from the C standard. On device compilation trajectory, device compilers substitute their own versions, so these are always safe to use. The value of LvArray::math versions is in generic code, since they can usually dispatch appropriately depending on argument type (e.g. to fabsf vs fabs vs abs).

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@TotoGaz @klevzoff is last commit acceptable ? If a pow function is added to lvarray it will be used first.

Copy link
Contributor

@klevzoff klevzoff Jan 17, 2024

Choose a reason for hiding this comment

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

It's all good by me as long as it compiles on all platforms :)

Note that adding using namespace LvArray::math does not necessarily mean the LvArray version will be called, it just adds another set of templates/overloads in the set and lets overload resolution pick the winner. For example, given

double pow(double base, double exp);

in the global namespace coming from C stdlib, and

template<typename T>
T pow(T base, T exp);

in LvArray::math, the first one will be picked (since a non-template perfect match wins over a template). The end result is the same though. The details about when C stdlib functions are added or not to the global namespace are notoriously murky and even less clear with CUDA (we seem to get more than is documented). The whole situation is a mess :) The one thing we should always avoid is using std:: version in kernel code, which this PR correctly does.

Copy link
Contributor

Choose a reason for hiding this comment

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

I don't know if this PR (or anywhere else for that matter) has using namespace LvArray::math, but for exactly that reason it should be avoided. Using LvArray::math::pow is not a elegant solution but at least there is zero ambiguity.

Copy link
Contributor

Choose a reason for hiding this comment

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

OK, so let's wait for @untereiner to explicitly use LvArray::math.

Maybe using LvArray::math::pow;, using LvArray::math::abs;... can help at having the formula a bit short?

In the long term, maybe could we have a geos::math namespace that would allow us to forward the LvArray::math namespace, so its usage will be more natural?

Copy link
Contributor Author

@untereiner untereiner Jan 18, 2024

Choose a reason for hiding this comment

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

Thank you @klevzoff. I thought the opposite: if you bring functions from a namespace into scope then they would be prioritized. Maybe the C versions should be put in a namespace and should not hang around and pollute the global one.

@CusiniM
Copy link
Collaborator

CusiniM commented Jan 10, 2024

This does not seem to be ready to be merged. Neither @TotoGaz nor @rrsettgast have approved it. And it looks like @TotoGaz had some comments.

@tbeltzun tbeltzun requested review from tbeltzun and removed request for tbeltzun January 17, 2024 11:23
Copy link
Contributor

@TotoGaz TotoGaz left a comment

Choose a reason for hiding this comment

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

This LGTM. I'm waiting for @klevzoff 's expertise though!

@untereiner
Copy link
Contributor Author

as a side note, if this solution is acceptable it should be added to another places.

Copy link

codecov bot commented Jan 18, 2024

Codecov Report

Attention: 6 lines in your changes are missing coverage. Please review.

Comparison is base (3ccaaf9) 51.33% compared to head (402286d) 51.42%.
Report is 1 commits behind head on develop.

❗ Current head 402286d differs from pull request most recent head 2daac51. Consider uploading reports for the commit 2daac51 to get more accurate results

Files Patch % Lines
...agation/ElasticFirstOrderWaveEquationSEMKernel.hpp 0.00% 3 Missing ⚠️
...s/wavePropagation/ElasticWaveEquationSEMKernel.hpp 0.00% 3 Missing ⚠️
Additional details and impacted files
@@             Coverage Diff             @@
##           develop    #2927      +/-   ##
===========================================
+ Coverage    51.33%   51.42%   +0.08%     
===========================================
  Files          968      966       -2     
  Lines        86819    86672     -147     
===========================================
+ Hits         44570    44571       +1     
+ Misses       42249    42101     -148     

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

@TotoGaz TotoGaz changed the title add std namespace Use LvArray::math functions Jan 22, 2024
@rrsettgast rrsettgast merged commit e322e0a into develop Jan 22, 2024
@rrsettgast rrsettgast deleted the bugfix/2926 branch January 22, 2024 23:24
ouassimkh pushed a commit that referenced this pull request Feb 16, 2024
* use lvarray::math instead of std::
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
ci: run CUDA builds Allows to triggers (costly) CUDA jobs ci: run integrated tests Allows to run the integrated tests in GEOS CI flag: no rebaseline Does not require rebaseline
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Compilation issue
7 participants