tforgot to set ghost nodes on device - 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 de26bf7c571b8579d75dca40e7fe0944e6eb8b23
 (DIR) parent 1658688f8ba0d701c10244801f13fd46821c799c
 (HTM) Author: Anders Damsgaard <anders.damsgaard@geo.au.dk>
       Date:   Wed,  9 Oct 2013 11:22:23 +0200
       
       forgot to set ghost nodes on device
       
       Diffstat:
         M src/darcy.cuh                       |       2 +-
         M src/device.cu                       |      22 ++++++++++++++++++++++
       
       2 files changed, 23 insertions(+), 1 deletion(-)
       ---
 (DIR) diff --git a/src/darcy.cuh b/src/darcy.cuh
       t@@ -60,7 +60,7 @@ void DEM::transferDarcyToGlobalDeviceMemory(int statusmsg)
                    "transferDarcyToGlobalDeviceMemory");
        
            if (verbose == 1 && statusmsg == 1)
       -        std::cout << "  Transfering darcy data to the device:           ";
       +        std::cout << "  Transfering fluid data to the device:           ";
        
            // number of cells
            //unsigned int ncells = d_nx*d_ny*d_nz; // without ghost nodes
 (DIR) diff --git a/src/device.cu b/src/device.cu
       t@@ -702,6 +702,7 @@ __host__ void DEM::startTime()
        
            double t_findPorositiesDev = 0.0;
            double t_findDarcyTransmissivitiesDev = 0.0;
       +    double t_setDarcyGhostNodesDev = 0.0;
            double t_explDarcyStepDev = 0.0;
            double t_findDarcyGradientsDev = 0.0;
            double t_findDarcyVelocitiesDev = 0.0;
       t@@ -930,6 +931,25 @@ __host__ void DEM::startTime()
                    // Perform explicit Darcy time step
                    if (PROFILING == 1)
                        startTimer(&kernel_tic);
       +            setDarcyGhostNodesDev<<<dimGridFluid, dimBlockFluid>>>(
       +                    dev_d_H,
       +                    dev_d_H_new,
       +                    dev_d_V,
       +                    dev_d_dH,
       +                    dev_d_K,
       +                    dev_d_T,
       +                    dev_d_Ss,
       +                    dev_d_W,
       +                    dev_d_phi);
       +            cudaThreadSynchronize();
       +            if (PROFILING == 1)
       +                stopTimer(&kernel_tic, &kernel_toc, &kernel_elapsed,
       +                        &t_setDarcyGhostNodesDev);
       +            checkForCudaErrors("Post setDarcyGhostNodesDev", iter);
       +
       +            // Perform explicit Darcy time step
       +            if (PROFILING == 1)
       +                startTimer(&kernel_tic);
                    explDarcyStepDev<<<dimGridFluid, dimBlockFluid>>>(
                            dev_d_H,
                            dev_d_H_new,
       t@@ -1192,6 +1212,8 @@ __host__ void DEM::startTime()
                    << "  - findDarcyTransmissivitiesDev:\t" <<
                    t_findDarcyTransmissivitiesDev/1000.0 << " s"
                    << "\t(" << 100.0*t_findDarcyTransmissivitiesDev/t_sum << " %)\n"
       +            << "  - setDarcyGhostNodesDev:\t" << t_setDarcyGhostNodesDev/1000.0
       +            << " s" << "\t(" << 100.0*t_setDarcyGhostNodesDev/t_sum << " %)\n"
                    << "  - explDarcyStepDev:\t" << t_explDarcyStepDev/1000.0 << " s"
                    << "\t(" << 100.0*t_explDarcyStepDev/t_sum << " %)\n"
                    << "  - findDarcyGradientsDev:\t" << t_findDarcyGradientsDev/1000.0