tadd pressure gradient array - 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 8ef27c9281201850da0bcc2095604dec0d743aac
 (DIR) parent f746002709f2c7c3ddad8b68ce7a745125abcadc
 (HTM) Author: Anders Damsgaard <anders.damsgaard@geo.au.dk>
       Date:   Fri, 20 Mar 2015 10:08:26 +0100
       
       add pressure gradient array
       
       Diffstat:
         M src/darcy.cuh                       |       2 ++
         M src/device.cu                       |      16 ++++++++++++++++
         M src/sphere.h                        |       1 +
       
       3 files changed, 19 insertions(+), 0 deletions(-)
       ---
 (DIR) diff --git a/src/darcy.cuh b/src/darcy.cuh
       t@@ -37,6 +37,7 @@ void DEM::initDarcyMemDev(void)
            cudaMalloc((void**)&dev_darcy_k, memSizeF);        // hydraulic permeability
            cudaMalloc((void**)&dev_darcy_grad_k, memSizeF*3); // grad(permeability)
            cudaMalloc((void**)&dev_darcy_div_v_p, memSizeF);  // divergence(v_p)
       +    cudaMalloc((void**)&dev_darcy_grad_p, memSizeF*3); // grad(pressure)
            //cudaMalloc((void**)&dev_darcy_v_p_x, memSizeFace); // v_p.x
            //cudaMalloc((void**)&dev_darcy_v_p_y, memSizeFace); // v_p.y
            //cudaMalloc((void**)&dev_darcy_v_p_z, memSizeFace); // v_p.z
       t@@ -65,6 +66,7 @@ void DEM::freeDarcyMemDev()
            //cudaFree(dev_darcy_v_p_x);
            //cudaFree(dev_darcy_v_p_y);
            //cudaFree(dev_darcy_v_p_z);
       +    cudaFree(dev_darcy_grad_p);
        }
        
        // Transfer to device
 (DIR) diff --git a/src/device.cu b/src/device.cu
       t@@ -1811,6 +1811,22 @@ __host__ void DEM::startTime()
        
                            if (PROFILING == 1)
                                startTimer(&kernel_tic);
       +                    findDarcyPressureGradient<<<dimGridFluid, dimBlockFluid>>>(
       +                            dev_darcy_p,
       +                            dev_darcy_grad_p);
       +                    cudaThreadSynchronize();
       +                    checkForCudaErrorsIter("After findDarcyPressureGradient",
       +                            iter);
       +
       +                    setDarcyGhostNodes<Float3><<<dimGridFluid, dimBlockFluid>>>(
       +                            dev_darcy_grad_p, darcy.bc_bot, darcy.bc_top);
       +                    cudaThreadSynchronize();
       +                    if (PROFILING == 1)
       +                        stopTimer(&kernel_tic, &kernel_toc, &kernel_elapsed,
       +                                &t_setDarcyGhostNodes);
       +                    checkForCudaErrorsIter("Post setDarcyGhostNodes("
       +                            "dev_darcy_grad_p)", iter);
       +
                            /*findDarcyPressureForce<<<dimGrid, dimBlock>>>(
                                    dev_x,
                                    dev_darcy_p,
 (DIR) diff --git a/src/sphere.h b/src/sphere.h
       t@@ -312,6 +312,7 @@ class DEM {
                Float4* dev_darcy_f_p;       // Pressure gradient force on particles
                Float*  dev_darcy_k;         // Cell hydraulic permeability
                Float3* dev_darcy_grad_k;    // Spatial gradient of permeability
       +        Float3* dev_darcy_grad_p;    // Spatial gradient of fluid pressure
        
                // Darcy functions
                void initDarcyMem();