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.

Sign up for GitHub

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 some syncs from MLMG #4340

Open
wants to merge 11 commits into
base: development
Choose a base branch
from
Open

Remove some syncs from MLMG #4340

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
05d0f54
Remove some syncs from MLMG
AlexanderSinn Feb 19, 2025
2f33b84
improve copyFrom
AlexanderSinn Feb 20, 2025
9be5ad5
reuse TagVector in applyBC
AlexanderSinn Feb 21, 2025
dc9c62d
fix cpu
AlexanderSinn Feb 21, 2025
e0283ea
fix cpu 2
AlexanderSinn Feb 21, 2025
5578ab8
build_arrays memcpy_async
AlexanderSinn Feb 21, 2025
8215c09
fix TagVector
AlexanderSinn Feb 21, 2025
1ed47dd
fix TagVector
AlexanderSinn Feb 21, 2025
36a4b87
add abec_gsrb_v2
AlexanderSinn Feb 21, 2025
4cea7f3
fix
AlexanderSinn Feb 21, 2025
09b7fe8
fix 2
AlexanderSinn Feb 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions Src/Base/AMReX_FBI.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,9 @@ void
fab_to_fab (Vector<Array4CopyTag<T0, T1> > const& copy_tags, int scomp, int dcomp, int ncomp,
F && f)
{
detail::ParallelFor_doit(copy_tags,
TagVector<Array4CopyTag<T0, T1>> tv{copy_tags};

detail::ParallelFor_doit(tv,
[=] AMREX_GPU_DEVICE (
#ifdef AMREX_USE_SYCL
sycl::nd_item<1> const& /*item*/,
Expand Down Expand Up @@ -85,7 +87,9 @@ fab_to_fab (Vector<Array4CopyTag<T0, T1> > const& copy_tags, int scomp, int dcom

amrex::Abort("xxxxx TODO This function still has a bug. Even if we fix the bug, it should still be avoided because it is slow due to the lack of atomic operations for this type.");

detail::ParallelFor_doit(tags,
TagVector<TagType> tv{tags};

detail::ParallelFor_doit(tv,
[=] AMREX_GPU_DEVICE (
#ifdef AMREX_USE_SYCL
sycl::nd_item<1> const& item,
Expand Down
14 changes: 10 additions & 4 deletions Src/Base/AMReX_FabArray.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1623,7 +1623,10 @@ FabArray<FAB>::build_arrays () const
#ifdef AMREX_USE_GPU
m_arrays.dp = (A*)m_dp_arrays;
m_const_arrays.dp = (AC*)m_dp_arrays + n;
Gpu::htod_memcpy(m_dp_arrays, m_hp_arrays, n*2*sizeof(A));
Gpu::htod_memcpy_async(m_dp_arrays, m_hp_arrays, n*2*sizeof(A));
if (!Gpu::inNoSyncRegion()) {
Gpu::streamSynchronize();
}
#endif
}
}
Expand All @@ -1633,9 +1636,12 @@ void
FabArray<FAB>::clear_arrays ()
{
#ifdef AMREX_USE_GPU
The_Pinned_Arena()->free(m_hp_arrays);
The_Arena()->free(m_dp_arrays);
m_dp_arrays = nullptr;
if (m_dp_arrays) {
Gpu::streamSynchronize();
The_Pinned_Arena()->free(m_hp_arrays);
The_Arena()->free(m_dp_arrays);
m_dp_arrays = nullptr;
}
#else
std::free(m_hp_arrays);
#endif
Expand Down
4 changes: 2 additions & 2 deletions Src/Base/AMReX_GpuControl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,8 +6,8 @@ namespace amrex::Gpu {
#if defined(AMREX_USE_GPU)
bool in_launch_region = true;
bool in_graph_region = false;
bool in_single_stream_region = false;
bool in_nosync_region = false;
bool in_single_stream_region = true;
bool in_nosync_region = true;
#endif

}
Loading
Loading