Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
2 changes: 1 addition & 1 deletion environment.yml
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@ channels:
dependencies:
- python <3.13
- cmake >=3.24
- cuda-version <=13.1
- cuda-version <= 13.1
- gxx >9,<15
- cuda-libraries-dev
- cuda-nvcc
Expand Down
4 changes: 1 addition & 3 deletions solvers/DPStokes/extra/uammd_interface.h
Original file line number Diff line number Diff line change
Expand Up @@ -20,16 +20,14 @@ std::string getPrecision();

struct PyParameters {
// The number of cells in each direction
// If -1, they will be autocomputed from the tolerance if possible (DP cannot
// do it, FCM can)
int nx = -1;
int ny = -1;
int nz = -1;
real viscosity;
real Lx;
real Ly;
real zmin, zmax;
// Tolerance will be ignored in DP mode, TP will use only tolerance and nxy/nz
// Tolerance will be ignored in DP mode
real tolerance = 1e-5;
real delta = 1e-3; // RFD step size
real w, w_d;
Expand Down
56 changes: 9 additions & 47 deletions solvers/DPStokes/extra/uammd_wrapper.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/* Raul P. Pelaez 2021. Doubly Periodic Stokes UAMMD wrapper
Allows to call the DPStokes or TP FCM modules from via a simple contained
Allows to call the DPStokes module from via a simple contained
class to compute the product between the mobility tensor and a list forces
and torques acting on a group of positions.

Expand All @@ -9,14 +9,10 @@
*/
#include <uammd.cuh>
// Doubly Periodic FCM implementation (currently without noise)
#include <Integrator/BDHI/DoublyPeriodic/DPStokesSlab.cuh>
// Triply Periodic FCM implementation
#include "uammd_interface.h"
#include <Integrator/BDHI/BDHI_FCM.cuh>
#include <Integrator/BDHI/DoublyPeriodic/DPStokesSlab.cuh>
// Some convenient aliases
namespace uammd_dpstokes {
using FCM_BM = uammd::BDHI::FCM_ns::Kernels::BarnettMagland;
using FCM = uammd::BDHI::FCM_impl<FCM_BM, FCM_BM>;
using DPStokesSlab = uammd::DPStokesSlab_ns::DPStokes;
using uammd::System;
using uammd::DPStokesSlab_ns::WallMode;
Expand Down Expand Up @@ -45,24 +41,6 @@ struct Real3ToReal4SubstractOriginZ {
}
};

auto createFCMParameters(PyParameters pypar) {
FCM::Parameters par;
par.temperature =
0; // FCM can compute fluctuations, but they are turned off here
par.viscosity = pypar.viscosity;
par.tolerance = pypar.tolerance;
par.box = uammd::Box({pypar.Lx, pypar.Ly, pypar.zmax - pypar.zmin});
par.cells = {pypar.nx, pypar.ny, pypar.nz};
par.kernel =
std::make_shared<FCM_BM>(pypar.w, pypar.alpha,
pypar.beta.x, // TODO beta parameter may need to
// be adjusted for non-square?
pypar.Lx / pypar.nx);
par.kernelTorque = std::make_shared<FCM_BM>(
pypar.w_d, pypar.alpha_d, pypar.beta_d.x, pypar.Lx / pypar.nx);
return par;
}

WallMode stringToWallMode(std::string str) {
if (str.compare("nowall") == 0) {
return WallMode::none;
Expand Down Expand Up @@ -100,24 +78,14 @@ private:
const uammd::real4 *d_force,
const uammd::real4 *d_torques,
int numberParticles, cudaStream_t st) {
if (fcm) {
return fcm->computeHydrodynamicDisplacements(
(uammd::real4 *)(d_pos), (uammd::real4 *)(d_force),
(uammd::real4 *)(d_torques), numberParticles, 0.0, 0.0, st);
} else if (dpstokes) {
return dpstokes->Mdot(reinterpret_cast<const uammd::real4 *>(d_pos),
reinterpret_cast<const uammd::real4 *>(d_force),
reinterpret_cast<const uammd::real4 *>(d_torques),
numberParticles, st);
} else {
throw std::runtime_error("DPStokesUAMMD: No DPStokes or FCM module "
"initialized. This should not had happened");
}
return dpstokes->Mdot(reinterpret_cast<const uammd::real4 *>(d_pos),
reinterpret_cast<const uammd::real4 *>(d_force),
reinterpret_cast<const uammd::real4 *>(d_torques),
numberParticles, st);
}

public:
std::shared_ptr<DPStokesSlab> dpstokes;
std::shared_ptr<FCM> fcm;
cudaStream_t st;
thrust::device_vector<uammd::real3> tmp3;
thrust::device_vector<uammd::real4> force4;
Expand All @@ -127,15 +95,9 @@ public:
real zOrigin;

DPStokesUAMMD(PyParameters pypar) {
if (pypar.mode.compare("periodic") == 0) {
auto par = createFCMParameters(pypar);
this->fcm = std::make_shared<FCM>(par);
zOrigin = 0;
} else {
auto par = createDPStokesParameters(pypar);
this->dpstokes = std::make_shared<DPStokesSlab>(par);
zOrigin = pypar.zmin + par.H * 0.5;
}
auto par = createDPStokesParameters(pypar);
this->dpstokes = std::make_shared<DPStokesSlab>(par);
zOrigin = pypar.zmin + par.H * 0.5;
CudaSafeCall(cudaStreamCreate(&st));
}

Expand Down
Loading