tuse particle velocity divergence as forcing term - 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 76388c20dccb3c13a6e654bf584713427b8fd113
(DIR) parent 6fcad793bd1dae1d962f5ea36b5c73a267d7e40f
(HTM) Author: Anders Damsgaard <anders.damsgaard@geo.au.dk>
Date: Thu, 19 Mar 2015 16:16:49 +0100
use particle velocity divergence as forcing term
Diffstat:
M src/darcy.cuh | 20 ++++++++++----------
M src/device.cu | 20 ++++++++++++++++----
2 files changed, 26 insertions(+), 14 deletions(-)
---
(DIR) diff --git a/src/darcy.cuh b/src/darcy.cuh
t@@ -1208,7 +1208,7 @@ __global__ void firstDarcySolution(
const Float* __restrict__ dev_darcy_k, // in
const Float* __restrict__ dev_darcy_phi, // in
const Float* __restrict__ dev_darcy_dphi, // in
- //const Float* __restrict__ dev_darcy_div_v_p, // in
+ const Float* __restrict__ dev_darcy_div_v_p, // in
const Float3* __restrict__ dev_darcy_grad_k, // in
const Float beta_f, // in
const Float mu, // in
t@@ -1244,7 +1244,7 @@ __global__ void firstDarcySolution(
const Float3 grad_k = dev_darcy_grad_k[cellidx];
const Float phi = dev_darcy_phi[cellidx];
const Float dphi = dev_darcy_dphi[cellidx];
- //const Float div_v_p = dev_darcy_div_v_p[cellidx];
+ const Float div_v_p = dev_darcy_div_v_p[cellidx];
const Float p_xn = dev_darcy_p[d_idx(x-1,y,z)];
const Float p = dev_darcy_p[cellidx];
t@@ -1290,8 +1290,8 @@ __global__ void firstDarcySolution(
Float dp_expl =
+ (ndem*devC_dt)/(beta_f*phi*mu)*(k*laplace_p + dot(grad_k, grad_p))
- //- div_v_p/(beta_f*phi);
- - dphi/(beta_f*phi*(1.0 - phi));
+ - div_v_p/(beta_f*phi);
+ //- dphi/(beta_f*phi*(1.0 - phi));
// Dirichlet BC at dynamic top wall. wall0_iz will be larger than the
// grid if the wall isn't dynamic
t@@ -1316,7 +1316,7 @@ __global__ void firstDarcySolution(
"grad_k = %e, %e, %e\n"
"dp_diff = %e\n"
"dp_forc = %e\n"
- //"div_v_p = %e\n"
+ "div_v_p = %e\n"
"dphi = %e\n"
//"dphi/dt = %e\n"
,
t@@ -1330,7 +1330,7 @@ __global__ void firstDarcySolution(
grad_p.x, grad_p.y, grad_p.z,
grad_k.x, grad_k.y, grad_k.z,
dp_diff, dp_forc,
- //div_v_p//,
+ div_v_p,
dphi//,
//dphi/(ndem*devC_dt)
);
t@@ -1356,7 +1356,7 @@ __global__ void updateDarcySolution(
const Float* __restrict__ dev_darcy_k, // in
const Float* __restrict__ dev_darcy_phi, // in
const Float* __restrict__ dev_darcy_dphi, // in
- //const Float* __restrict__ dev_darcy_div_v_p, // in
+ const Float* __restrict__ dev_darcy_div_v_p, // in
const Float3* __restrict__ dev_darcy_grad_k, // in
const Float beta_f, // in
const Float mu, // in
t@@ -1393,7 +1393,7 @@ __global__ void updateDarcySolution(
const Float3 grad_k = dev_darcy_grad_k[cellidx];
const Float phi = dev_darcy_phi[cellidx];
const Float dphi = dev_darcy_dphi[cellidx];
- //const Float div_v_p = dev_darcy_div_v_p[cellidx];
+ const Float div_v_p = dev_darcy_div_v_p[cellidx];
const Float p_old = dev_darcy_p_old[cellidx];
const Float dp_expl = dev_darcy_dp_expl[cellidx];
t@@ -1443,8 +1443,8 @@ __global__ void updateDarcySolution(
//Float p_new = p_old
Float dp_impl =
+ (ndem*devC_dt)/(beta_f*phi*mu)*(k*laplace_p + dot(grad_k, grad_p))
- //- div_v_p/(beta_f*phi);
- - dphi/(beta_f*phi*(1.0 - phi));
+ - div_v_p/(beta_f*phi);
+ //- dphi/(beta_f*phi*(1.0 - phi));
// Dirichlet BC at dynamic top wall. wall0_iz will be larger than the
// grid if the wall isn't dynamic
(DIR) diff --git a/src/device.cu b/src/device.cu
t@@ -1830,7 +1830,7 @@ __host__ void DEM::startTime()
if (PROFILING == 1)
startTimer(&kernel_tic);
- findDarcyPorosities<<<dimGridFluid, dimBlockFluid>>>(
+ /*findDarcyPorosities<<<dimGridFluid, dimBlockFluid>>>(
dev_cellStart,
dev_cellEnd,
dev_x_sorted,
t@@ -1840,7 +1840,19 @@ __host__ void DEM::startTime()
np,
darcy.c_phi,
dev_darcy_phi,
- dev_darcy_dphi);//,
+ dev_darcy_dphi);*/
+ findDarcyPorositiesLinear<<<dimGridFluid, dimBlockFluid>>>(
+ dev_cellStart,
+ dev_cellEnd,
+ dev_x_sorted,
+ dev_vel_sorted,
+ iter,
+ darcy.ndem,
+ np,
+ darcy.c_phi,
+ dev_darcy_phi,
+ dev_darcy_dphi,
+ dev_darcy_div_v_p);
cudaThreadSynchronize();
if (PROFILING == 1)
stopTimer(&kernel_tic, &kernel_toc, &kernel_elapsed,
t@@ -1984,7 +1996,7 @@ __host__ void DEM::startTime()
dev_darcy_k,
dev_darcy_phi,
dev_darcy_dphi,
- //dev_darcy_div_v_p,
+ dev_darcy_div_v_p,
dev_darcy_grad_k,
darcy.beta_f,
darcy.mu,
t@@ -2012,7 +2024,7 @@ __host__ void DEM::startTime()
dev_darcy_k,
dev_darcy_phi,
dev_darcy_dphi,
- //dev_darcy_div_v_p,
+ dev_darcy_div_v_p,
dev_darcy_grad_k,
darcy.beta_f,
darcy.mu,