     __syncthreads();
 
     // read in the position of our particle. (MEM TRANSFER: 16 bytes)
+    #ifdef SINGLE_PRECISION
     Scalar4 postype = tex1Dfetch(pdata_pos_tex, idx);
+    #else
+    Scalar4 postype = d_pos[idx];
+    #endif
     Scalar3 pos = make_scalar3(postype.x, postype.y, postype.z);
 
     // read in the diameter of our particle if needed
     Scalar diam(0);
     if (evaluator::needsDiameter())
+        {
+        #ifdef SINGLE_PRECISION
         diam = tex1Dfetch(pdata_diam_tex, idx);
+        #else
+        diam = d_diameter[idx];
+        #endif
+        }
     else
         diam += 0; // shutup compiler warning
 
     Scalar q(0);
     if (evaluator::needsCharge())
+        {
+        #ifdef SINGLE_PRECISION
         q = tex1Dfetch(pdata_charge_tex, idx);
+        #else
+        q = d_charge[idx];
+        #endif
+        }
     else
         q += 0; // shutup compiler warning
 

        // get the bonded particle's diameter if needed
        if (evaluator::needsDiameter())
            {
            Scalar neigh_diam = tex1Dfetch(pdata_diam_tex, cur_bond_idx);
            eval.setDiameter(diam, neigh_diam);
            }
        if (evaluator::needsCharge())
            {
            Scalar neigh_q = tex1Dfetch(pdata_charge_tex, cur_bond_idx);
            eval.setCharge(q, neigh_q);
            }


        int cur_bond_type = cur_bond.y;

        // get the bonded particle's position (MEM_TRANSFER: 16 bytes)
        #ifdef SINGLE_PRECISION
        Scalar4 neigh_postypej = tex1Dfetch(pdata_pos_tex, cur_bond_idx);
        #else
        Scalar4 neigh_postypej = d_pos[cur_bond_idx];
        #endif
        Scalar3 neigh_pos= make_scalar3(neigh_postypej.x, neigh_postypej.y, neigh_postypej.z);

        // calculate dr (FLOPS: 3)
        Scalar3 dx = pos - neigh_pos;
