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

                                           float4* d_net_force, 
                                           float* d_partial_sum_pe)
    {
    int idx_local = blockIdx.x * blockDim.x + threadIdx.x;
    
    float pe;
     if (idx_local < pdata.local_num)
        {        
        pe = d_net_force[idx_local].w;
        }
    else
        {
        pe = 0.0f;
        }

    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

                                            unsigned int group_size,
                                            float* d_partial_sum_vsq)
    {
    // 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 vsq;
   
        float4 v = tex1Dfetch(pdata_vel_tex, idx);
        vsq = v.x*v.x + v.y*v.y + v.z*v.z;

        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
