                                             unsigned int group_size,
                                             float* d_partial_sum_asq)
     {
-    int idx_local = blockIdx.x * blockDim.x + threadIdx.x;
-    int idx_global = idx_local + pdata.local_beg;
-    
-    float asq;
-     if (idx_local < pdata.local_num)
-        {   
-        float4 a = tex1Dfetch(pdata_accel_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 asq;
+        float4 a = tex1Dfetch(pdata_accel_tex, idx);
         asq = a.x*a.x + a.y*a.y + a.z*a.z;
-        }
-    else
-        {
-        asq = 0.0f;
-        }
-
-    fire_sdata[threadIdx.x] = asq;
-    __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_asq[blockIdx.x] = fire_sdata[0];
-    }
-
+
+
+        fire_sdata[threadIdx.x] = asq;
+        __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_asq[blockIdx.x] = fire_sdata[0];
+        }
+    }
 
 //! Kernel function to simultaneously compute the partial sum over P, vsq and asq for the FIRE algorithm
 /*! \param pdata Particle data to compute P, vsq and asq for

        vel.w = 0.0f;
                
        // write out the results (MEM_TRANSFER: 32 bytes)
        pdata.vel[idx_global] = vel;
        }
    }


                                          unsigned int group_size,    
                                          float* d_partial_sum_P)
    {
    // 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];
     
        float4 a = tex1Dfetch(pdata_accel_tex, idx);
        float4 v = tex1Dfetch(pdata_vel_tex, idx);
        float P = a.x*v.x + a.y*v.y + a.z*v.z;
        
        fire_sdata[threadIdx.x] = P;
        __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_P[blockIdx.x] = fire_sdata[0];
        }
    }
    
//! Kernel function to compute the partial sum over the vsq term in the FIRE algorithm
/*! \param pdata Particle data to compute vsq for
    \param d_group_members Array listing members of the group
