tVarious wforce changes - 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 14400684ad8dbfc7b9e2e9efa0de245227f025a6
 (DIR) parent 18b731037ae5f967592fe454bd63bf2a679b3253
 (HTM) Author: Anders Damsgaard <adc@geo.au.dk>
       Date:   Wed, 28 Nov 2012 13:41:39 +0100
       
       Various wforce changes
       
       Diffstat:
         M src/contactsearch.cuh               |       9 ++++++---
         M src/device.cu                       |       2 +-
         M src/integration.cuh                 |       3 +++
       
       3 files changed, 10 insertions(+), 4 deletions(-)
       ---
 (DIR) diff --git a/src/contactsearch.cuh b/src/contactsearch.cuh
       t@@ -544,15 +544,16 @@ __global__ void interact(unsigned int* dev_gridParticleIndex, // Input: Unsorted
            //// Interact with walls
            Float delta_w; // Overlap distance
            Float3 w_n;    // Wall surface normal
       -    Float w_force = 0.0f; // Force on wall from particle A
       +    Float w_force = 0.0; // Force on wall from particle A
        
            // Upper wall (idx 0)
            delta_w = w_up_nx.w - (x_a.z + radius_a);
            w_n = MAKE_FLOAT3(0.0f, 0.0f, -1.0f);
            if (delta_w < 0.0f) {
       -      w_force = contactLinear_wall(&F, &T, &es_dot, &ev_dot, &p, idx_a, radius_a,
       +        w_force = contactLinear_wall(&F, &T, &es_dot, &ev_dot, &p, idx_a, radius_a,
                                             dev_vel_sorted, dev_angvel_sorted,
                                           w_n, delta_w, w_up_mvfd.y);
       +        //cuPrintf("particle %d collides with upper wall, wforce = %f\n", idx_a, w_force);
            }
        
            // Lower wall (force on wall not stored)
       t@@ -637,8 +638,10 @@ __global__ void interact(unsigned int* dev_gridParticleIndex, // Input: Unsorted
            dev_es[orig_idx]     += es_dot * devC_dt;
            dev_ev[orig_idx]     += ev_dot * devC_dt;
            dev_p[orig_idx]       = p;
       -    if (devC_nw > 0)
       +    if (devC_nw > 0 && w_force != 0.0) {
              dev_walls_force_pp[orig_idx] = w_force;
       +      //cuPrintf("wforce written\n");
       +    }
          }
        } // End of interact(...)
        
 (DIR) diff --git a/src/device.cu b/src/device.cu
       t@@ -656,9 +656,9 @@ __host__ void DEM::startTime()
              checkForCudaErrors("Post topology: One or more particles moved outside the grid.\nThis could possibly be caused by a numerical instability.\nIs the computational time step too large?", iter);
            }
        
       -    //cudaPrintfInit();
        
            // For each particle: Process collisions and compute resulting forces.
       +    //cudaPrintfInit();
            if (PROFILING == 1)
              startTimer(&kernel_tic);
            interact<<<dimGrid, dimBlock>>>(dev_gridParticleIndex,
 (DIR) diff --git a/src/integration.cuh b/src/integration.cuh
       t@@ -238,6 +238,7 @@ __global__ void integrateWalls(Float4* dev_walls_nx,
            int wmode = dev_walls_wmode[idx];  // Wall BC, 0: fixed, 1: devs, 2: vel
            Float acc;
        
       +
            if (wmode == 0) // Wall fixed: do nothing
              return;
        
       t@@ -271,6 +272,8 @@ __global__ void integrateWalls(Float4* dev_walls_nx,
            // Update linear velocity
            w_mvfd.y += acc * dt;
        
       +    //cuPrintf("\nwall %d, wmode = %d, force = %f, acc = %f\n", idx, wmode, w_mvfd.z, acc);
       +
            // Store data in global memory
            dev_walls_nx[idx]   = w_nx;
            dev_walls_mvfd[idx] = w_mvfd;