         unsigned int idx = d_group_members[group_idx];
 
         // read in the net force and compute the acceleration
-        float4 accel = tex1Dfetch(net_force_tex, idx);
-        float mass = tex1Dfetch(pdata_mass_tex, idx);
+        float4 accel = net_force[idx];
+        float mass = pdata.mass[idx];
         accel.x /= mass;
         accel.y /= mass;
         accel.z /= mass;
 
         // fetch velocities
-        float4 vel = tex1Dfetch(pdata_vel_tex, idx);
+        float4 vel = pdata.vel[idx];
 
         // propagate velocities from t+1/2*deltaT to t+deltaT
         vel.x +=  (1.0f/2.0f)*deltaT*accel.x;

    dim3 grid( (group_size / block_size) + 1, 1, 1);
    dim3 threads(block_size, 1, 1);

    // bind the texture
    cudaError_t error = cudaBindTexture(0, pdata_vel_tex, pdata.vel, sizeof(float4) * pdata.N);
    if (error != cudaSuccess)
        return error;

    error = cudaBindTexture(0, pdata_mass_tex, pdata.mass, sizeof(float) * pdata.N);
    if (error != cudaSuccess)
        return error;

    error = cudaBindTexture(0, net_force_tex, d_net_force, sizeof(float4) * pdata.N);
    if (error != cudaSuccess)
        return error;

    // run the kernel
    gpu_nph_step_two_kernel<<< grid, threads >>>(pdata, d_group_members, group_size, deltaT);

    return cudaSuccess;
    }

    unsigned int block_size = 256;
    dim3 grid( (group_size / block_size) + 1, 1, 1);
    dim3 threads(block_size, 1, 1);

    // run the kernel
    gpu_nph_step_one_kernel<<< grid, threads >>>(pdata,
