tWall viscosities now in constant parameter structure. Dynamic arrays now transfered one-by-one - sphere - GPU-based 3D discrete element method algorithm with optional fluid coupling
(HTM) git clone git://src.adamsgaard.dk/sphere
(DIR) Log
(DIR) Files
(DIR) Refs
(DIR) LICENSE
---
(DIR) commit e13db0e693ca0fcddcdf072a561c99e01fc74f6e
(DIR) parent f07848d9a5dda616d45ce0e8fe5696882370b9fc
(HTM) Author: Anders Damsgaard <adc@geo.au.dk>
Date: Wed, 31 Oct 2012 08:58:19 +0100
Wall viscosities now in constant parameter structure. Dynamic arrays now transfered one-by-one
Diffstat:
M src/datatypes.h | 9 +++++----
M src/device.cu | 70 ++++++++++++++++++++++++-------
M src/file_io.cpp | 14 ++++++++------
3 files changed, 68 insertions(+), 25 deletions(-)
---
(DIR) diff --git a/src/datatypes.h b/src/datatypes.h
t@@ -29,7 +29,7 @@ struct Kinematics {
Float4 *delta_t; // Accumulated shear distance of contacts
};
-// Structure containing individual physical particle parameters
+// Structure containing individual particle energies
struct Energies {
Float *es_dot; // Frictional dissipation rates
Float *es; // Frictional dissipations
t@@ -75,6 +75,10 @@ struct Params {
Float gamma_n; // Normal viscosity
Float gamma_t; // Tangential viscosity
Float gamma_r; // Rotational viscosity
+ Float mu_ws; // Wall static friction coefficient
+ Float mu_wd; // Wall dynamic friction coefficient
+ Float gamma_wn; // Wall normal viscosity
+ Float gamma_wt; // Wall tangential viscosity
Float mu_s; // Static friction coefficient
Float mu_d; // Dynamic friction coefficient
Float mu_r; // Rotational friction coefficient
t@@ -89,9 +93,6 @@ struct Params {
struct Walls {
unsigned int nw; // Number of walls (<= MAXWALLS)
int wmode[MAXWALLS]; // Wall modes
- Float gamma_wn; // Wall normal viscosity
- Float gamma_wt; // Wall tangential viscosity
- Float gamma_wr; // Wall rolling viscosity
Float4* nx; // Wall normal and position
Float4* mvfd; // Wall mass, velocity, force and dev. stress
Float* force; // Resulting forces on walls per particle
(DIR) diff --git a/src/device.cu b/src/device.cu
t@@ -218,7 +218,7 @@ __host__ void DEM::allocateGlobalDeviceMemory(void)
unsigned int memSizeF4 = sizeof(Float4) * np;
if (verbose == 1)
- std::cout << " Allocating device memory: ";
+ std::cout << " Allocating global device memory: ";
// Particle arrays
cudaMalloc((void**)&dev_k->x, memSizeF4);
t@@ -316,15 +316,60 @@ __host__ void DEM::transferToGlobalDeviceMemory()
if (verbose == 1)
std::cout << " Transfering data to the device: ";
- // Copy structure data from host to global device memoryj
- cudaMemcpy(dev_k, &k, sizeof(Kinematics), cudaMemcpyHostToDevice);
- cudaMemcpy(dev_e, &e, sizeof(Energies), cudaMemcpyHostToDevice);
+ // Commonly-used memory sizes
+ unsigned int memSizeF = sizeof(Float) * np;
+ unsigned int memSizeF4 = sizeof(Float4) * np;
+
+ // Copy static-size structure data from host to global device memory
cudaMemcpy(dev_time, &time, sizeof(Time), cudaMemcpyHostToDevice);
- cudaMemcpy(dev_walls, &walls, sizeof(Walls), cudaMemcpyHostToDevice);
- /*cudaMemcpy(dev_k, &k, sizeof(k), cudaMemcpyHostToDevice);
- cudaMemcpy(dev_e, &e, sizeof(e), cudaMemcpyHostToDevice);
- cudaMemcpy(dev_time, &time, sizeof(time), cudaMemcpyHostToDevice);
- cudaMemcpy(dev_walls, &walls, sizeof(walls), cudaMemcpyHostToDevice);*/
+
+ // Kinematic particle values
+ cudaMemcpy( dev_k->x, k.x,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->xysum, k.xysum,
+ sizeof(Float2)*np, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->vel, k.vel,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->acc, k.acc,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->force, k.force,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->angpos, k.angpos,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->angvel, k.angvel,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->angacc, k.angacc,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->torque, k.torque,
+ memSizeF4, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->contacts, k.contacts,
+ sizeof(unsigned int)*np*NC, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->distmod, k.distmod,
+ memSizeF4*NC, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_k->delta_t, k.delta_t,
+ memSizeF4*NC, cudaMemcpyHostToDevice);
+
+ // Individual particle energy values
+ cudaMemcpy( dev_e->es_dot, e.es_dot,
+ memSizeF, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_e->es, e.es,
+ memSizeF, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_e->ev_dot, e.ev_dot,
+ memSizeF, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_e->ev, e.ev,
+ memSizeF, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_e->p, e.p,
+ memSizeF, cudaMemcpyHostToDevice);
+
+ // Wall parameters
+ cudaMemcpy( dev_walls->wmode, walls.wmode,
+ sizeof(int)*walls.nw, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_walls->nx, walls.nx,
+ sizeof(Float4)*walls.nw, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_walls->mvfd, walls.mvfd,
+ sizeof(Float4)*walls.nw, cudaMemcpyHostToDevice);
+ cudaMemcpy( dev_walls->force, walls.force,
+ memSizeF*walls.nw, cudaMemcpyHostToDevice);
checkForCudaErrors("End of transferToGlobalDeviceMemory");
if (verbose == 1)
t@@ -363,12 +408,7 @@ __host__ void DEM::startTime()
// Copy data to constant global device memory
transferToConstantDeviceMemory();
-
- // Particle memory size
- //unsigned int memSizeF = sizeof(Float) * np;
- //unsigned int memSizeF4 = sizeof(Float4) * np;
-
- // Allocate device memory for particle variables,
+ // Allocate device memory for particle variables,
// tied to previously declared pointers in structures
allocateGlobalDeviceMemory();
(DIR) diff --git a/src/file_io.cpp b/src/file_io.cpp
t@@ -193,6 +193,10 @@ void DEM::readbin(const char *target)
ifs.read(as_bytes(params.mu_s), sizeof(params.mu_s));
ifs.read(as_bytes(params.mu_d), sizeof(params.mu_d));
ifs.read(as_bytes(params.mu_r), sizeof(params.mu_r));
+ ifs.read(as_bytes(params.gamma_wn), sizeof(params.gamma_wn));
+ ifs.read(as_bytes(params.gamma_wt), sizeof(params.gamma_wt));
+ ifs.read(as_bytes(params.mu_ws), sizeof(params.mu_s));
+ ifs.read(as_bytes(params.mu_wd), sizeof(params.mu_d));
ifs.read(as_bytes(params.rho), sizeof(params.rho));
ifs.read(as_bytes(params.contactmodel), sizeof(params.contactmodel));
ifs.read(as_bytes(params.kappa), sizeof(params.kappa));
t@@ -228,9 +232,6 @@ void DEM::readbin(const char *target)
ifs.read(as_bytes(walls.mvfd[i].z), sizeof(Float));
ifs.read(as_bytes(walls.mvfd[i].w), sizeof(Float));
}
- ifs.read(as_bytes(walls.gamma_wn), sizeof(walls.gamma_wn));
- ifs.read(as_bytes(walls.gamma_wt), sizeof(walls.gamma_wt));
- ifs.read(as_bytes(walls.gamma_wr), sizeof(walls.gamma_wr));
// Close file if it is still open
if (ifs.is_open())
t@@ -338,6 +339,10 @@ void DEM::writebin(const char *target)
ofs.write(as_bytes(params.mu_s), sizeof(params.mu_s));
ofs.write(as_bytes(params.mu_d), sizeof(params.mu_d));
ofs.write(as_bytes(params.mu_r), sizeof(params.mu_r));
+ ofs.write(as_bytes(params.gamma_wn), sizeof(params.gamma_wn));
+ ofs.write(as_bytes(params.gamma_wt), sizeof(params.gamma_wt));
+ ofs.write(as_bytes(params.mu_ws), sizeof(params.mu_ws));
+ ofs.write(as_bytes(params.mu_wd), sizeof(params.mu_wd));
ofs.write(as_bytes(params.rho), sizeof(params.rho));
ofs.write(as_bytes(params.contactmodel), sizeof(params.contactmodel));
ofs.write(as_bytes(params.kappa), sizeof(params.kappa));
t@@ -359,9 +364,6 @@ void DEM::writebin(const char *target)
ofs.write(as_bytes(walls.mvfd[i].z), sizeof(Float));
ofs.write(as_bytes(walls.mvfd[i].w), sizeof(Float));
}
- ofs.write(as_bytes(walls.gamma_wn), sizeof(walls.gamma_wn));
- ofs.write(as_bytes(walls.gamma_wt), sizeof(walls.gamma_wt));
- ofs.write(as_bytes(walls.gamma_wr), sizeof(walls.gamma_wr));
// Close file if it is still open
if (ofs.is_open())