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

Convert IntVect to IntVectND #3969

Merged

Conversation

AlexanderSinn
Copy link
Member

@AlexanderSinn AlexanderSinn commented Jun 1, 2024

Summary

As described in #3955, this PR converts IntVect to the n dimensional IntVectND and adds
using IntVect = IntVectND<AMREX_SPACEDIM>;. Additionally, deduction guides, support for structured bindings and the helper functions IntVectCat, IntVectSplit and IntVectResize were added to IntVectND.

I didn't find a way to make IntVect::Zero and IntVect::Unit work with n dimensions, so I removed them (converting them to static constexpr or inline static doesn't work with MSVC). Instead, IntVect::TheZeroVector() and IntVect::TheUnitVector() can be used.

Additional background

Using structured binding support of amrex::GpuTuple, the following code should work (on GPU):

int i, j, k, n;
...
// amrex::IntVectND<4>
amrex::IntVectND iv1 (i, j, k, n);
// amrex::IntVectND<8>
amrex::IntVectND iv2 = amrex::IntVectCat(iv1, iv1);
// ... = amrex::GpuTuple<amrex::IntVectND<3>,amrex::IntVectND<2>,amrex::IntVectND<2>,amrex::IntVectND<1>>
auto [iv3, iv4, iv5, iv6] = amrex::IntVectSplit<3,2,2,1>(iv2); 
// int, int, int = amrex::IntVectND<3>
auto [i2, j2, k2] = iv3;
// int = amrex::IntVectND<1>
auto [n2] = iv6;

assert(i==i2 && j==j2 && k==k2 && n==n2);

Checklist

The proposed changes:

  • fix a bug or incorrect behavior in AMReX
  • add new capabilities to AMReX
  • changes answers in the test suite to more than roundoff level
  • are likely to significantly affect the results of downstream AMReX users
  • include documentation in the code and/or rst files, if appropriate

@AlexanderSinn AlexanderSinn marked this pull request as ready for review June 8, 2024 17:03
@AlexanderSinn AlexanderSinn mentioned this pull request Jun 8, 2024
5 tasks
WeiqunZhang pushed a commit that referenced this pull request Jun 10, 2024
## Summary

This PR adds a `TupleSplit` function for `GpuTuple` similar
to `IntVectSplit` in #3969. It can be used as an inverse function
of `TupleCat` and `TypeMultiplier`.
Example:
```C++
auto tup = amrex::makeTuple(2,4,5,7,2.324,7,8,342.3f,4ull,1ll,-38,"test");
auto [t0,t1,t2,t3] = amrex::TupleSplit<3,3,4,2>(tup);
// t0 = GpuTuple( 2, 4, 5 )
// t1 = GpuTuple( 7, 2.324, 7 )
// t2 = GpuTuple( 8, 342.2999878, 4, 1 )
// t3 = GpuTuple( -38, test )
```

## Additional background

The current implementation does not move the tuple elements and makes a
copy instead.
@WeiqunZhang
Copy link
Member

WeiqunZhang commented Jun 11, 2024

Should we have extern template class IntVectND<AMREX_SPACEDIM>, which might help reduce compile time?

@AlexanderSinn
Copy link
Member Author

Wouldn't that prevent the member functions from getting inlined, which could be bad when IntVect is used in a parallel for?

@WeiqunZhang
Copy link
Member

My understanding is it should not affect run time performance. But I might be wrong.

@WeiqunZhang
Copy link
Member

WeiqunZhang commented Jun 11, 2024 via email

@AlexanderSinn
Copy link
Member Author

I tried to add extern template class IntVectND<AMREX_SPACEDIM>;however, without adding template class IntVectND<AMREX_SPACEDIM>;to the cpp file. I expected to get a linker error however it compiled just fine.

@WeiqunZhang
Copy link
Member

Maybe it's because all those functions have already been inlined.

@WeiqunZhang
Copy link
Member

diff --git a/Src/Base/AMReX_IntVect.H b/Src/Base/AMReX_IntVect.H
index 70a12811a2..7c472742ae 100644
--- a/Src/Base/AMReX_IntVect.H
+++ b/Src/Base/AMReX_IntVect.H
@@ -710,6 +710,9 @@ public:
         return IntVectND<dim>(std::numeric_limits<int>::lowest());
     }
 
+    static const IntVectND<dim> Zero;
+    static const IntVectND<dim> Unit;
+
     AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
     static constexpr std::size_t size () noexcept {
         return static_cast<std::size_t>(dim);
@@ -767,6 +770,13 @@ private:
     int vect[dim] = {};
 };
 
+
+template <>
+inline constexpr IntVectND<AMREX_SPACEDIM> IntVectND<AMREX_SPACEDIM>::Zero{0};
+
+template <>
+inline constexpr IntVectND<AMREX_SPACEDIM> IntVectND<AMREX_SPACEDIM>::Unit{1};
+
 // Template deduction guide for IntVectND
 template<std::size_t dim>
 AMREX_GPU_HOST_DEVICE // __device__ for HIP

This seems to work. So we could get IntVect::Zero and IntVect::Unit back for backward compatibility.

Unfortunately template <int dim> inline constexpr IntVectND<dim> IntVectND<dim>::Zero{0} might only work for gcc. https://stackoverflow.com/questions/72409091/define-a-static-constexpr-member-of-same-type-of-a-template-class

@WeiqunZhang
Copy link
Member

Unfortunately, nvcc does not work.

@AlexanderSinn
Copy link
Member Author

Maybe instead of being IntVectND, Zero and Unit could have a different type that is implicitly convertible to IntVectND. This would allow them to be static constexpr.

@WeiqunZhang
Copy link
Member

That may not work in cases like

IntVect::Zero + sign(a_side) * 2 * BASISV(idir);

This is copied from a real code.

Maybe we can try to figure out how to make nvcc happy. The error message is

/home/runner/work/amrex/amrex/Src/Base/AMReX_IntVect.H:775:34: error: conflicting declaration ‘amrex::IntVectND<3> amrex::IntVectND<3>::Zero’
  775 | inline constexpr IntVectND<AMREX_SPACEDIM> IntVectND<AMREX_SPACEDIM>::Zero{0};
      |                                  ^~~~~~~~~~~~~
/home/runner/work/amrex/amrex/Src/Base/AMReX_IntVect.H:713:24: note: previous declaration as ‘const amrex::IntVectND<3> amrex::IntVectND<3>::Zero’
  713 |     static const IntVectND<dim> Zero;
      |                        ^~~~

@WeiqunZhang
Copy link
Member

Maybe it wants constexpr const?

@AlexanderSinn
Copy link
Member Author

It seems to work if both are just const with no constexpr.

@WeiqunZhang
Copy link
Member

WeiqunZhang commented Jun 12, 2024

Great. Maybe it will even work for all <int dim>.

@AlexanderSinn
Copy link
Member Author

constexpr const actually works somehow

@WeiqunZhang
Copy link
Member

@ax3l This PR will break pyamrex. How do you want to handle this? Should we wait till pyamrex has a PR ready?

@ax3l
Copy link
Member

ax3l commented Jun 12, 2024

Thanks for the ping. That is all we need to track.

Please go ahead here to merge and we update pyAMReX shortly after.

@WeiqunZhang WeiqunZhang enabled auto-merge (squash) June 12, 2024 20:00
@WeiqunZhang WeiqunZhang merged commit dad25a4 into AMReX-Codes:development Jun 12, 2024
69 of 70 checks passed
@ax3l
Copy link
Member

ax3l commented Jun 18, 2024

pyAMReX update coming via AMReX-Codes/pyamrex#332 :)

WeiqunZhang pushed a commit that referenced this pull request Jul 3, 2024
## Summary

Similar to #3969 but for IndexType.

## Additional background

A maximum of 31 dimensions are supported so that (1u << dim) can fit
into an unsigned int.
@AlexanderSinn AlexanderSinn mentioned this pull request Jul 3, 2024
5 tasks
WeiqunZhang pushed a commit that referenced this pull request Aug 2, 2024
## Summary

Similar to #3969 and #3988 but for Box.

## Additional background

It should be checked that the changes to BoxIndexer do not affect the
compiled GPU code.
In my testing, it gives the same performance as development. 

Example usage:
```C++
amrex::BoxND b1{amrex::IntVectND{1,2,3}, amrex::IntVectND{4,5,6}, amrex::IntVectND{1,0,1}};
// ((1,2,3) (4,5,6) (1,0,1))
auto b2 = amrex::BoxCat(b1, b1, b1);
// ((1,2,3,1,2,3,1,2,3) (4,5,6,4,5,6,4,5,6) (1,0,1,1,0,1,1,0,1))
auto [b3, b4, b5, b6, b7] = amrex::BoxSplit<1, 4, 2, 1, 1>(b2);
// ((1) (4) (1))((2,3,1,2) (5,6,4,5) (0,1,1,0))((3,1) (6,4) (1,1))((2) (5) (0))((3) (6) (1))
auto b8 = amrex::BoxResize<2>(b4);
// ((2,3) (5,6) (0,1))
auto b9 = amrex::BoxResize<5>(b8);
// ((2,3,0,0,0) (5,6,0,0,0) (0,1,0,0,0))
```
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

3 participants