                                             float* d_partial_sum_vsq, 
                                             float* d_partial_sum_asq)
     {
-    int idx_local = blockIdx.x * blockDim.x + threadIdx.x;
-    int idx_global = idx_local + pdata.local_beg;
-    
-    float P, vsq, asq;
-     if (idx_local < pdata.local_num)
-        {   
-        float4 a = tex1Dfetch(pdata_accel_tex, idx_global);
-        float4 v = tex1Dfetch(pdata_vel_tex, idx_global);
+    // determine which particle this thread works on (MEM TRANSFER: 4 bytes)
+    int group_idx = blockIdx.x * blockDim.x + threadIdx.x;
+    
+    if (group_idx < group_size)
+        {
+        unsigned int idx = d_group_members[group_idx];
+    
+        float P, vsq, asq; 
+        float4 a = tex1Dfetch(pdata_accel_tex, idx);
+        float4 v = tex1Dfetch(pdata_vel_tex, idx);
         P = a.x*v.x + a.y*v.y + a.z*v.z;
         vsq = v.x*v.x + v.y*v.y + v.z*v.z;
         asq = a.x*a.x + a.y*a.y + a.z*a.z;
-        }
-    else
-        {
-        P = 0.0f;
-        asq = 0.0f;
-        vsq = 0.0f;
-        }
-
-    fire_sdata1[threadIdx.x] = P;
-    fire_sdata2[threadIdx.x] = vsq;
-    fire_sdata3[threadIdx.x] = asq;
-    __syncthreads();
-    
-    // reduce the sum in parallel
-    int offs = blockDim.x >> 1;
-    while (offs > 0)
-        {
-        if (threadIdx.x < offs)
-            {
-            fire_sdata1[threadIdx.x] += fire_sdata1[threadIdx.x + offs];
-            fire_sdata2[threadIdx.x] += fire_sdata2[threadIdx.x + offs];
-            fire_sdata3[threadIdx.x] += fire_sdata3[threadIdx.x + offs];
-            }
-        offs >>= 1;
-        __syncthreads();
-        }
-        
-    // write out our partial sum
-    if (threadIdx.x == 0)
-        {
-        d_partial_sum_P[blockIdx.x] = fire_sdata1[0];
-        d_partial_sum_vsq[blockIdx.x] = fire_sdata2[0];
-        d_partial_sum_asq[blockIdx.x] = fire_sdata3[0];
-        }
-    }
-
+
+        fire_sdata1[threadIdx.x] = P;
+        fire_sdata2[threadIdx.x] = vsq;
+        fire_sdata3[threadIdx.x] = asq;
+        __syncthreads();
+    
+        // reduce the sum in parallel
+        int offs = blockDim.x >> 1;
+        while (offs > 0)
+            {
+            if (threadIdx.x < offs)
+                {
+                fire_sdata1[threadIdx.x] += fire_sdata1[threadIdx.x + offs];
+                fire_sdata2[threadIdx.x] += fire_sdata2[threadIdx.x + offs];
+                fire_sdata3[threadIdx.x] += fire_sdata3[threadIdx.x + offs];
+                }
+            offs >>= 1;
+            __syncthreads();
+            }
+        
+        // write out our partial sum
+        if (threadIdx.x == 0)
+            {
+            d_partial_sum_P[blockIdx.x] = fire_sdata1[0];
+            d_partial_sum_vsq[blockIdx.x] = fire_sdata2[0];
+            d_partial_sum_asq[blockIdx.x] = fire_sdata3[0];
+            }
+        }
+    }
 
 //! Kernel function to simultaneously reduce three partial sums at the same time
 /*! \param d_sum Array to hold the sums

                                            unsigned int group_size,
                                            float* d_partial_sum_vsq)
    {
    int idx_local = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_global = idx_local + pdata.local_beg;
    
    float vsq;
     if (idx_local < pdata.local_num)
        {   
        float4 v = tex1Dfetch(pdata_vel_tex, idx_global);
        vsq = v.x*v.x + v.y*v.y + v.z*v.z;
        }
    else
        {
        vsq = 0.0f;
        }

    fire_sdata[threadIdx.x] = vsq;
    __syncthreads();
    
    // reduce the sum in parallel
    int offs = blockDim.x >> 1;
    while (offs > 0)
        {
        if (threadIdx.x < offs)
            fire_sdata[threadIdx.x] += fire_sdata[threadIdx.x + offs];
        offs >>= 1;
        __syncthreads();
        }
        
    // write out our partial sum
    if (threadIdx.x == 0)
        d_partial_sum_vsq[blockIdx.x] = fire_sdata[0];
    }

//! Kernel function to compute the partial sum over the asq term in the FIRE algorithm
/*! \param pdata Particle data to compute asq for
    \param d_group_members Array listing members of the group

                                           float4* d_net_force, 
                                           float* d_partial_sum_pe)
    {
    // determine which particle this thread works on (MEM TRANSFER: 4 bytes)
    int group_idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (group_idx < group_size)
        {
        unsigned int idx = d_group_members[group_idx];
 
        float pe;
        pe = d_net_force[idx].w;
         
               
        fire_sdata[threadIdx.x] = pe;
        __syncthreads();
    
        // reduce the sum in parallel
        int offs = blockDim.x >> 1;
        while (offs > 0)
            {
            if (threadIdx.x < offs)
                fire_sdata[threadIdx.x] += fire_sdata[threadIdx.x + offs];
            offs >>= 1;
            __syncthreads();
            }
        
        // write out our partial sum
        if (threadIdx.x == 0)
            {
            d_partial_sum_pe[blockIdx.x] = fire_sdata[0];
            }
        }
    }
    
//! Kernel function for reducing a partial sum to a full sum (one value)
/*! \param d_sum Placeholder for the sum
    \param d_partial_sum Array containing the parial sum
