Skip to content

Lees-Edwards boundary conditions for shear flow #2122

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

Open
wkdarko wants to merge 29 commits into
base: trunk
Choose a base branch
from
Open

Lees-Edwards boundary conditions for shear flow #2122

Show file tree
Hide file tree
Changes from 27 commits
Commits
Show all changes
29 commits
Select commit Hold shift + click to select a range
147b1fe
added variable m_xy_rate to box
wkdarko Apr 10, 2025
ab6982b
wrap pos and vel back into box
wkdarko Apr 15, 2025
6284dd0
Apply suggestions from code review
wkdarko Apr 16, 2025
dc3c4e5
updated velocity wrap
wkdarko Apr 16, 2025
aed343d
Merge branch 'lees-edwards' into trunk-minor
wkdarko Apr 23, 2025
6877fb9
updated release version
wkdarko Apr 23, 2025
7013d49
velocity wrap in 2step CV integrator
wkdarko Apr 24, 2025
bf0a1bb
velocity wrap on GPU 2step CV
Apr 24, 2025
16ffcab
velocity wrap on GPU 2step CV
wkdarko Apr 24, 2025
8b9ad3e
Merge remote-tracking branch 'refs/remotes/origin/lees-edwards' into …
May 1, 2025
70f547b
remove wrap from 2nd step
wkdarko May 1, 2025
1bb8ff7
edit minimage for DPD
wkdarko May 29, 2025
74c778d
wrap/minimage/shift modifications
wkdarko Jun 27, 2025
a31b0cc
Apply suggestions from code review
wkdarko Jul 23, 2025
0ac2ef1
Merge remote-tracking branch 'upstream/trunk-minor' into lees-edwards
wkdarko Jul 23, 2025
d3c6e88
further changes to box methods
wkdarko Jul 31, 2025
ac3166f
revert file hoomd/test/test_pdata.cc to version at 49a93b32
wkdarko Aug 19, 2025
5babb15
Apply suggestions from code review
wkdarko Aug 21, 2025
a4ee6a8
code review suggestions - manual
wkdarko Aug 21, 2025
8aa0752
Merge remote-tracking branch 'origin/lees-edwards' into lees-edwards
wkdarko Aug 21, 2025
d50715c
unit tests on box deformation
wkdarko Aug 21, 2025
23527da
unit tests
wkdarko Aug 22, 2025
ac2f517
comprehensive unit tests on box deformation
wkdarko Sep 11, 2025
f227dc0
Merge remote-tracking branch 'upstream/trunk-minor' into lees-edwards
wkdarko Oct 3, 2025
3dfdc47
fix pre-commit format issues
wkdarko Oct 3, 2025
b80186a
Revert change to tutorial submodule
mphoward Oct 6, 2025
003d90b
modified gpu wrap kernel for ghost communication
wkdarko Oct 8, 2025
b5ccfd5
apply code review suggestions
wkdarko Oct 22, 2025
3ffe596
Merge remote-tracking branch 'upstream/trunk' into lees-edwards
wkdarko Nov 26, 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
325 changes: 325 additions & 0 deletions hoomd/BoxDim.h
View file
Open in desktop

Large diffs are not rendered by default.

5 changes: 4 additions & 1 deletion hoomd/BoxResizeUpdater.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -96,6 +96,9 @@ void BoxResizeUpdater::scaleAndWrapParticles(const BoxDim& cur_box, const BoxDim

// ensure that the particles are still in their
// local boxes by wrapping them if they are not
ArrayHandle<Scalar4> h_vel(m_pdata->getVelocities(),
access_location::host,
access_mode::readwrite);
ArrayHandle<int3> h_image(m_pdata->getImages(), access_location::host, access_mode::readwrite);

const BoxDim& local_box = m_pdata->getBox();
Expand All @@ -104,7 +107,7 @@ void BoxResizeUpdater::scaleAndWrapParticles(const BoxDim& cur_box, const BoxDim
{
// need to update the image if we move particles from one side
// of the box to the other
local_box.wrap(h_pos.data[i], h_image.data[i]);
local_box.wrap(h_pos.data[i], h_vel.data[i], h_image.data[i]);
}
}

Expand Down
5 changes: 5 additions & 0 deletions hoomd/BoxResizeUpdaterGPU.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,10 @@ void BoxResizeUpdaterGPU::scaleAndWrapParticles(const BoxDim& cur_box, const Box
access_location::device,
access_mode::readwrite);

ArrayHandle<Scalar4> d_vel(m_pdata->getVelocities(),
access_location::device,
access_mode::readwrite);

ArrayHandle<int3> d_image(m_pdata->getImages(),
access_location::device,
access_mode::readwrite);
Expand All @@ -62,6 +66,7 @@ void BoxResizeUpdaterGPU::scaleAndWrapParticles(const BoxDim& cur_box, const Box
m_tuner_wrap->begin();
kernel::gpu_box_resize_wrap(m_pdata->getN(),
d_pos.data,
d_vel.data,
d_image.data,
new_box,
m_tuner_wrap->getParam()[0]);
Expand Down
11 changes: 8 additions & 3 deletions hoomd/BoxResizeUpdaterGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,14 +31,17 @@ __global__ void gpu_box_resize_scale_kernel(Scalar4* d_pos,
}
}

__global__ void
gpu_box_resize_wrap_kernel(unsigned int N, Scalar4* d_pos, int3* d_image, const BoxDim new_box)
__global__ void gpu_box_resize_wrap_kernel(unsigned int N,
Scalar4* d_pos,
Scalar4* d_vel,
int3* d_image,
const BoxDim new_box)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < N)
{
new_box.wrap(d_pos[idx], d_image[idx]);
new_box.wrap(d_pos[idx], d_vel[idx], d_image[idx]);
}
}

Expand Down Expand Up @@ -74,6 +77,7 @@ hipError_t gpu_box_resize_scale(Scalar4* d_pos,

hipError_t gpu_box_resize_wrap(const unsigned int N,
Scalar4* d_pos,
Scalar4* d_vel,
int3* d_image,
const BoxDim& new_box,
unsigned int block_size)
Expand All @@ -94,6 +98,7 @@ hipError_t gpu_box_resize_wrap(const unsigned int N,
0,
N,
d_pos,
d_vel,
d_image,
new_box);

Expand Down
1 change: 1 addition & 0 deletions hoomd/BoxResizeUpdaterGPU.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@ hipError_t gpu_box_resize_scale(Scalar4* d_pos,

hipError_t gpu_box_resize_wrap(const unsigned int N,
Scalar4* d_pos,
Scalar4* d_vel,
int3* d_image,
const BoxDim& new_box,
unsigned int block_size);
Expand Down
15 changes: 12 additions & 3 deletions hoomd/Communicator.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1719,9 +1719,10 @@ void Communicator::migrateParticles()
{
detail::pdata_element& p = m_recvbuf[idx];
Scalar4& postype = p.pos;
Scalar4& vel = p.vel;
int3& image = p.image;

shifted_box.wrap(postype, image);
shifted_box.wrap(postype, vel, image);
}

// remove particles that were sent and fill particle data with received particles
Expand Down Expand Up @@ -2394,6 +2395,9 @@ void Communicator::exchangeGhosts()
ArrayHandle<Scalar4> h_pos(m_pdata->getPositions(),
access_location::host,
access_mode::readwrite);
ArrayHandle<Scalar4> h_vel(m_pdata->getVelocities(),
access_location::host,
access_mode::readwrite);
ArrayHandle<int3> h_image(m_pdata->getImages(),
access_location::host,
access_mode::readwrite);
Expand All @@ -2403,10 +2407,11 @@ void Communicator::exchangeGhosts()
for (unsigned int idx = start_idx; idx < start_idx + m_num_recv_ghosts[dir]; idx++)
{
Scalar4& pos = h_pos.data[idx];
Scalar4& vel = h_vel.data[idx];

// wrap particles received across a global boundary
int3& img = h_image.data[idx];
shifted_box.wrap(pos, img);
shifted_box.wrap(pos, vel, img);
}
}

Expand Down Expand Up @@ -2966,15 +2971,19 @@ void Communicator::beginUpdateGhosts(uint64_t timestep)
ArrayHandle<Scalar4> h_pos(m_pdata->getPositions(),
access_location::host,
access_mode::readwrite);
ArrayHandle<Scalar4> h_vel(m_pdata->getVelocities(),
access_location::host,
access_mode::readwrite);

const BoxDim shifted_box = getShiftedBox();
for (unsigned int idx = start_idx; idx < start_idx + m_num_recv_ghosts[dir]; idx++)
{
Scalar4& pos = h_pos.data[idx];
Scalar4& vel = h_vel.data[idx];

// wrap particles received across a global boundary
int3 img = make_int3(0, 0, 0);
shifted_box.wrap(pos, img);
shifted_box.wrap(pos, vel, img);
}
}

Expand Down
51 changes: 30 additions & 21 deletions hoomd/CommunicatorGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -260,7 +260,7 @@ __global__ void gpu_wrap_particles_kernel(const unsigned int n_recv,
return;

detail::pdata_element p = d_recv[idx];
box.wrap(p.pos, p.image);
box.wrap(p.pos, p.vel, p.image);
d_recv[idx] = p;
}

Expand Down Expand Up @@ -789,8 +789,10 @@ gpu_pack_kernel(unsigned int n_out, const uint2* d_ghost_idx_adj, const T* in, T
__global__ void gpu_pack_wrap_kernel(unsigned int n_out,
const uint2* d_ghost_idx_adj,
const Scalar4* d_postype,
const Scalar4* d_vel,
const int3* d_img,
Scalar4* out_pos,
Scalar4* out_vel,
int3* out_img,
Index3D di,
uint3 my_pos,
Expand Down Expand Up @@ -860,14 +862,24 @@ __global__ void gpu_pack_wrap_kernel(unsigned int n_out,
}

box.setPeriodic(periodic);
Scalar4 postype = d_postype[idx];
int3 img = make_int3(0, 0, 0);
if (d_img)
if (out_img && d_img)
{
img = d_img[idx];
Scalar4 postype = d_postype[idx];
box.wrap(postype, img, wrap);
}
if (out_vel && d_vel)
{
Scalar4 vel = d_vel[idx];
box.wrap(postype, vel, img, wrap);
out_vel[buf_idx] = vel;
}
else
{
box.wrap(postype, img, wrap);
}

out_pos[buf_idx] = postype;

if (out_img)
{
out_img[buf_idx] = img;
Expand Down Expand Up @@ -928,10 +940,19 @@ void gpu_exchange_ghosts_pack(unsigned int n_out,
d_tag,
d_tag_sendbuf);
}
if (send_pos)
/*
* combined packing pathway for positions , velocities and images
* call gpu_pack_wrap_kernel whenever any of send_pos, send_vel and send_image is true
* for buffers that are not being sent, pass a nullptr to the respective output
*/
if (send_pos || send_image || send_vel)
{
assert(d_pos);
assert(d_pos_sendbuf);
if (send_vel)
{
assert(d_vel);
}
if (send_image)
{
assert(d_img);
Expand All @@ -944,27 +965,15 @@ void gpu_exchange_ghosts_pack(unsigned int n_out,
n_out,
d_ghost_idx_adj,
d_pos,
d_img,
send_vel ? d_vel : 0,
send_image ? d_img : 0,
d_pos_sendbuf,
send_vel ? d_vel_sendbuf : 0,
send_image ? d_img_sendbuf : 0,
di,
my_pos,
box);
}
if (send_vel)
{
assert(d_vel);
assert(d_vel_sendbuf);
hipLaunchKernelGGL(gpu_pack_kernel,
dim3(n_blocks),
dim3(block_size),
0,
0,
n_out,
d_ghost_idx_adj,
d_vel,
d_vel_sendbuf);
}
if (send_charge)
{
assert(d_charge);
Expand Down
6 changes: 5 additions & 1 deletion hoomd/GSDDumpWriter.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1113,6 +1113,9 @@ void GSDDumpWriter::populateLocalFrame(GSDDumpWriter::GSDFrame& frame, uint64_t
ArrayHandle<Scalar4> h_postype(m_pdata->getPositions(),
access_location::host,
access_mode::read);
ArrayHandle<Scalar4> h_veltype(m_pdata->getVelocities(),
access_location::host,
access_mode::read);
ArrayHandle<int3> h_image(m_pdata->getImages(), access_location::host, access_mode::read);

if (m_dynamic[gsd_flag::particles_position] || m_nframes == 0)
Expand All @@ -1132,6 +1135,7 @@ void GSDDumpWriter::populateLocalFrame(GSDDumpWriter::GSDFrame& frame, uint64_t
{
vec3<Scalar> position
= vec3<Scalar>(h_postype.data[index]) - vec3<Scalar>(m_pdata->getOrigin());
vec3<Scalar> velocity = vec3<Scalar>(h_veltype.data[index]);
unsigned int type = __scalar_as_int(h_postype.data[index].w);
int3 image = make_int3(0, 0, 0);

Expand All @@ -1140,7 +1144,7 @@ void GSDDumpWriter::populateLocalFrame(GSDDumpWriter::GSDFrame& frame, uint64_t
image = h_image.data[index];
}

frame.global_box.wrap(position, image);
frame.global_box.wrap(position, velocity, image);

if (m_dynamic[gsd_flag::particles_position] || m_nframes == 0)
{
Expand Down
Loading
Loading